MLX nvfp4 uses signed E4M3 scales instead of unsigned UE4M3, causing 137x less dynamic range than NVIDIA Blackwell and inferior to Marlin

[Nottlespike]

Problem

MLX's nvfp4 implementation uses signed E4M3 (1 sign bit, 4-bit exponent with bias 7, 3-bit mantissa) for block scaling factors, but NVIDIA Blackwell GPUs use unsigned E4M3 (UE4M3) with no sign bit and no exponent bias, as specified in the PTX documentation.

Dynamic Range Comparison

• MLX nvfp4 scales: Max value 448 ((1+7/8) × 2^7)
• NVIDIA UE4M3 scales: Max value ~61,440 ((1+7/8) × 2^15)
• Difference: NVIDIA has 137x more dynamic range for scale factors

This means MLX's nvfp4 will saturate/clip on large activations and weights that NVIDIA's implementation handles without issue. Combined with the fact that Marlin's MR-GPTQ quantization for FP4 achieves superior accuracy (recovering 96.1% of FP16 accuracy), MLX's current approach is both format-incompatible and suboptimal.

Code References

Metal backend scale dequantization (mlx/backend/metal/kernels/fp_quantized.h: 30-38):

template <typename T, int group_size>
static inline T dequantize_scale(uint8_t s) {
  if constexpr (group_size == 16) {
    // Use nv scale - BUT THIS IS SIGNED E4M3, NOT UE4M3! 
    return T(*(thread fp8_e4m3*)(&s));  
  } else {
    return T(*(thread fp8_e8m0*)(&s));
  }
}

Metal quantize kernel (mlx/backend/metal/kernels/fp_quantized.h:1763):

using ScaleType = metal::conditional_t<use_mx_scale, fp8_e8m0, fp8_e4m3>;
auto s = ScaleType(scale);  // Wrong:  should be UE4M3

FP8 E4M3 struct (mlx/backend/metal/kernels/fp8. h:1-48) implements signed E4M3, not unsigned.

References

• NVIDIA PTX UE4M3 Specification
• Dynamic Range Analysis
• Marlin MR-GPTQ Paper - Shows MR-GPTQ for FP4 recovers 96.1% FP16 accuracy
• PR Metal/CPU nvfp4 and mxfp8 #2946 that introduced nvfp4: Metal/CPU nvfp4 and mxfp8 #2946

Recommendation

1. Implement UE4M3 format for nvfp4 scale factors (unsigned, 4-bit exponent without bias, 3-bit mantissa)
2. Consider adding MR-GPTQ support as Marlin's approach substantially improves FP4 quantization quality
3. Update both Metal and CUDA backends to match NVIDIA's Blackwell specification

This will:

• Achieve full dynamic range parity with NVIDIA hardware (137x improvement)
• Enable quantization of models with large magnitude weights/activations
• Improve compatibility with NVIDIA-quantized models
• Open path to superior MR-GPTQ quantization methods

[awni]

@nastya236 wdyt about this?

NVIDIA Blackwell GPUs use unsigned E4M3 (UE4M3) with no sign bit and no exponent bias

I don't see any evidence that there is no exponent bias and it's not clear to me that that is desirable. A larger range is nice but it's also nice to have granularity. The cublaslt docs just mention that it is unsigned which makes sense (the scales will always be unsigned in MLX as well).

[Nottlespike]

@nastya236 wdyt about this?

NVIDIA Blackwell GPUs use unsigned E4M3 (UE4M3) with no sign bit and no exponent bias

I don't see any evidence that there is no exponent bias and it's not clear to me that that is desirable. A larger range is nice but it's also nice to have granularity. The cublaslt docs just mention that it is unsigned which makes sense (the scales will always be unsigned in MLX as well).

I think there may be some confusion between a few different concepts here, so it might help to step through the specification.

You're right that E4M3 has an exponent bias (bias of 7, specifically), and you're also right that scales being unsigned makes sense since they represent magnitudes. Those points aren't in dispute.

The concern I was raising is about a different aspect of the NVFP4 format: the two-level scaling architecture. According to NVIDIA's documentation, NVFP4 uses both the per-block E4M3 scale and an additional per-tensor FP32 scalar:

"This strategy applies a fine-grained E4M3 scaling factor to each 16-value micro-block... while also leveraging a second-level FP32 scalar applied per tensor. Together, these two levels of scaling enable more accurate value representation and significantly reduce quantization error."

Source: https://developer.nvidia.com/blog/nvidia-blackwell-tensor-core-fp4-performance-and-quantization-techniques/

The FP32 tensor-level scale normalizes the entire tensor so that all block maxima fit within the E4M3 representable range. The formula given is s_g ≈ max_b(amax(b)) / (6 × 448), which maps the largest block maximum to roughly the E4M3 maximum (see arxiv.org/html/2509.25149v1, Section 2.1 for details).

Without this tensor-level normalization, the effective dynamic range is limited to around 2688x (the product of FP4 E2M1 max and E4M3 max). For tensors where cross-block magnitude variance stays within that range, single-level scaling works fine. But tensors with larger variance can run into issues where some blocks saturate while others underflow. Chmiel et al. touch on related dynamics in "FP4 All the Way" (arxiv.org/abs/2508.00612), where they show that scale formats with insufficient dynamic range can cause problems (Figure 1).

It might be worth looking at the Marlin implementation as a reference point, since it targets the same format for CUDA and should have the two-level scaling in place.

It's possible I'm missing something about the MLX implementation that addresses this differently. Happy to dig into the code if that would help clarify things.

[awni]

Yup we are aware of the fp32 scale of scales. We may add it in the future.

[Nottlespike]

Yup we are aware of the fp32 scale of scales. We may add it in the future.

The broader concern here is ecosystem interoperability. Once MLX users start quantizing models and uploading them to Hugging Face tagged as nvfp4, there's no easy way to distinguish them from spec-compliant NVFP4 checkpoints produced by TensorRT Model Optimizer or other tools.

Someone on Blackwell hardware downloads what they think is an NVFP4 checkpoint, loads it expecting two-level scaling, and gets subtly wrong results because the tensor-level scale is missing or set to 1.0. Or the reverse: someone tries to load a proper NVFP4 checkpoint into MLX and the dequantization is off because MLX ignores the FP32 scale. Either way, users end up debugging phantom quality issues with no obvious cause.

This isn't hypothetical. The GGUF ecosystem went through similar growing pains with quantization format naming, and it took considerable effort to sort out which quants were actually compatible with which backends. NVFP4 is early enough that it would be good to avoid repeating that.

If full two-level support is on the roadmap, perhaps the current implementation could use a provisional name like fp4_block_scaled or similar, then migrate to nvfp4 once the FP32 tensor scale is in place. That keeps the format usable now without creating a pool of checkpoints that will be silently incompatible with the broader NVFP4 ecosystem later.

[Artoria2e5]

If a rename is to be done (I think it should be!), it should be done as early as possible to avoid creating more wrongly-named things.

Also, while we’re already making new formats (accidentally), might as well make a version with bf16 block scale instead of e4m3… I even got names for these: bsfp4_e4m3, bsfp4_bf16 :)

(Honestly I’m surprised at nvidia even using fp32 on the outside scale. Surely a e8m0 would be good enough given the inner scale already has some mantissa?)

[nastya236]

Tensor scale for nvfp4 was added in #3022, so closing for now.