# NVFP4 at the bit level > Every other post says '4-bit' and moves on. Here's what NVFP4 actually is, down to the four bits and the two scales, and the two Blackwell instructions (cvt and mma) that make it fly. Source: https://rfriedmann.de/blog/nvfp4-at-the-bit-level/ Published: 2026-06-15 · Track: log · Level: Expert 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
[diagram omitted — see the page for the chart]
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
[diagram omitted — see the page for the chart]
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](/blog/optimizing-kernels-consumer-blackwell/) 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](/blog/gguf-decode-beats-llama-cpp/). 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](/blog/what-the-5090-lacks-vs-datacenter/), alongside memory bandwidth, the decode cache, and the attention kernels.