NVFP4 at the bit level - Raphael Friedmann

The rest of these posts lean on “4-bit” and “NVFP4” as if they were obvious. They are not, quite. This is the format spelled out: the four bits, the two scales that rescue them, and the two hardware instructions that make the whole thing fast on Blackwell, the cvt that dequantizes and the mma that does the GEMM.

Four bits: E2M1

NVFP4’s element is E2M1: one sign bit, two exponent bits, one mantissa bit. That is it. Those four bits can encode exactly eight magnitudes.

E2M1: the entire vocabulary of a 4-bit weight

Four bits per number: one sign, two exponent, one mantissa. These eight magnitudes (plus the mirror for negatives) are all a single NVFP4 weight can be. The dynamic range comes from the scales, not the value.

Round-to-nearest into this set is not a loop or a lookup, it is a branchless cascade of comparisons against the midpoints. Eight magnitudes, mirrored by the sign into sixteen values, is the entire vocabulary every weight gets squeezed into.

Two scales rescue four bits

Sixteen values obviously cannot represent a layer full of weights with wildly different magnitudes. The dynamic range does not come from the value, it comes from scales, and NVFP4 keeps two of them.

Two scales rescue four bits

A single 4-bit value can't span a whole layer's range, so NVFP4 keeps two scales: a fine FP8 one for every 16 values, and a coarse FP32 one for the tensor. MXFP4 instead uses a power-of-two scale per 32 values, which is what the tensor cores want natively.

A fine FP8 (E4M3) micro-scale covers every group of 16 values, and a single coarse FP32 tensor scale covers the whole tensor. Reconstructing a weight is just e2m1_value * micro_scale * tensor_scale. The per-tensor FP32 scale exists so the per-block scales fit into E4M3’s range, and it is folded into the scales rather than applied as a separate per-element multiply at GEMM time. The two-level scheme is what lets 4-bit hold an outlier-heavy weight distribution without flattening it.

MXFP4 is the close cousin: same E2M1 nibbles, but the per-group scale is a power-of-two (UE8M0) over 32 values, and there is no separate tensor scale. Both formats are equally native to the Blackwell tensor cores: the single MMA opcode kind mxf4nvf4 consumes MXFP4 and NVFP4 alike, distinguished only by the scale operand type (UE8M0 vs UE4M3) and the scale_vec granularity (per-32 vs per-16). The hardware does not merely tolerate NVFP4, it reads it directly, which matters for the next part.

The hardware that makes it fly

A format is only as good as the silicon that reads it, and this is where consumer Blackwell earns its keep.

On the decode path, a single PTX instruction, cvt.rn.f16x2.e2m1x2, takes one byte (two packed FP4 values) and produces two FP16 numbers in one shot, replacing roughly a dozen software operations. The dequant that would dominate a naive kernel is reduced to a single cvt instruction rather than eliminated.

On the prefill path, the matrix multiply runs on the tensor-core instruction mma.sync.aligned.m16n8k64.row.col.kind::mxf4nvf4.block_scale.scale_vec::4X.f32.e2m1.e2m1.f32.ue4m3: 64 FP4 elements per step, with the per-16 block scale (the scale_vec::4X and ue4m3 operands) applied inside the hardware. That one instruction is the foundation the entire NVFP4 GEMM is built on.

Packing and provenance

Two FP4 values share a byte (low nibble even index, high nibble odd). The weights usually arrive already in this form, exported by NVIDIA’s Model Optimizer or llm-compressor, and imp registers them as-is with no re-quantization. Or imp makes them itself, converting GGUF weights to NVFP4 at load time to feed the decode cache.

So when the other posts say a model “runs in 4-bit”, this is the machinery: eight magnitudes, two scales, and a one-instruction dequant that only exists because the 5090 has real FP4 tensor cores. It is a primary reason the 5090 is fast at NVFP4 inference, alongside memory bandwidth, the decode cache, and the attention kernels.