Low precision: FP8, FP4, and microscaling

Every other technique in this series moves a workload along the roofline. Low precision moves the roofline itself. On an H100, going from BF16 to FP8 halves the bytes a decode step must read and doubles the peak FLOP/s — 989.5 to 1979 TFLOP/s, dense. On a B200, FP4 does it again: 9 PFLOPS dense (18 with sparsity) against 4.5 for FP8.

The catch is that the numbers still have to be right. A BF16 value has 8 exponent bits and covers the same range as FP32; an FP8 E4M3 value has 4, and tops out at 448. One activation outlier can make an entire tensor unrepresentable. Everything below is a way of shrinking the group of numbers that has to share a single scale.

All of these are the same idea — a sign bit, some exponent bits, some mantissa bits — with the budget split differently. Exponent bits buy range; mantissa bits buy precision.

An FP8 E4M3 tensor can only hold values up to 448, so before you cast anything you multiply by a scale chosen so the largest element lands near the top of the range. The GEMM then runs in FP8 and exposes an FP32 accumulator/output interface; on Hopper, however, the tensor core retains fewer bits internally during FP8 accumulation, as §8.4 shows. The result is divided by the product of the two input scales.

// Quantize two FP32 values to E4M3 and back, with a per-tensor scale.
// s = amax / 448 comes out of an amax reduction; here we pass s and 1/s.
// sm_89+ (PTX ISA 8.1). .rn = round to nearest even; .satfinite = clamp
// to +-448 on overflow instead of producing inf -- E4M3 has none.
__device__ float2 qdq_e4m3(float2 v, float s, float s_inv)
{
    unsigned short q;                    // two packed e4m3 values
    // First source lands in the HIGH half, so the operands are swapped
    // to keep v.x in the low byte.
    asm("cvt.rn.satfinite.e4m3x2.f32 %0, %2, %1;"
        : "=h"(q) : "f"(v.x*s_inv), "f"(v.y*s_inv));
    unsigned int p;
    asm("cvt.rn.f16x2.e4m3x2 %0, %1;" : "=r"(p) : "h"(q));   // exact in f16
    float2 r = __half22float2(*reinterpret_cast<__half2*>(&p));
    r.x *= s;  r.y *= s;                 // the dequantization multiply
    return r;
}

# The same round trip in Triton: the compiler emits the same cvt
# instructions, and casts to fp8 saturate to +-448 — unlike eager
# PyTorch, whose float8_e4m3fn cast produces NaN on overflow.
xq = (x * s_inv).to(tl.float8e4nv)     # quantize: E4M3, clamped
xd = xq.to(tl.float32) * s             # dequantize, never leaves registers

That works well until the tensor contains an outlier. LLM activations reliably do: a handful of channels carry values one or two orders of magnitude larger than the rest. Scale for the outlier and many ordinary values may underflow or collapse into the bottom few representable values.

The central scaling question is therefore: how few numbers can be made to share one scale?

NVIDIA's taxonomy of the options, in increasing order of granularity:

• Per-tensor, delayed — the scale comes from a rolling history of amax. Smooths transient outliers, but a large one dominates the history for as long as the window lasts.
• Per-tensor, current — the scale comes from this pass's amax. Adapts instantly, no historical stability.
• Generic block scaling — user-chosen 1×128 or 128×128 blocks with FP32 scales. This is what DeepSeek-V3 uses.
• Microscaling — 32 elements per scale, in hardware. NVIDIA's MXFP8 pre-training recipes report less than 0.5% difference in validation perplexity against BF16 across a full pre-training run, on 843M and 8B dense models.

The OCP Microscaling (MX) specification — written jointly by AMD, Arm, Intel, Meta, Microsoft, NVIDIA and Qualcomm — fixes the block structure so hardware can implement it: 32 consecutive elements plus one shared scale in UE8M0 — eight exponent bits, no mantissa, so the scale is a pure power of two and applying it is an exponent add.

NVIDIA's Blackwell-native NVFP4 tightens both knobs: blocks of 16 instead of 32, and a two-level scale — an E4M3 microscale per block, which unlike a power of two can express fractions, plus a single FP32 scale per tensor that keeps those microscales inside E4M3's own range.

On Blackwell the grouping and the scale application happen inside the tensor core — you supply a scale-factor tensor in its required tiled layout alongside the data, and tcgen05.mma.block_scale does the rest. On Hopper there is no hardware block scaling, which is why DeepGEMM's README notes that SM90 wants FP32 scale factors while SM100 wants packed UE8M0.

NVIDIA reports accuracy degradation of 1% or less across seven evaluations when DeepSeek-R1-0528 is requantized from FP8 to NVFP4.

Here is a failure mode that is easy to miss, and DeepSeek-V3 documented it precisely: the Hopper FP8 tensor core does not accumulate in full FP32 internally.

Their fix is a nice illustration of using two units at once: every 128 elements of K — four wgmma instructions — the partial result is copied out to FP32 registers and accumulated on the CUDA cores at full precision. In this schedule, the dequantization multiply can overlap work on CUDA cores with tensor-core MMA. And because two wgmma operations are typically in flight, one warpgroup can be doing the promotion while the other is doing its MMA.

8.4.1 The rest of DeepSeek-V3's recipe

DeepSeek's other choices are worth recording because they contradict the usual advice:

• E4M3 everywhere, including gradients — not the conventional E4M3-forward / E5M2-backward split. Their argument: fine-grained grouping already shares exponent bits among small groups, so the extra exponent bit is not needed.
• Different granularity for activations and weights: 1×128 tiles for activations (per token, per 128 channels), 128×128 blocks for weights.
• These are not interchangeable. Appendix B.2 reports that 128×128 block quantization of activation gradients destabilises training — the backward Dgrad computation is specifically sensitive. Use tiles there.

Measured outcome: relative training-loss error below 0.25% at 16B/1.33T tokens and 230B/0.9T tokens, and DeepGEMM reaching up to 1550 FP8 TFLOP/s on an H800.

There are two fundamentally different reasons to quantize, and they help in opposite regimes.

Weight-only (W4A16 — GPTQ, AWQ) stores weights in 4 bits and dequantizes them in-kernel back to FP16 before the matmul. The tensor cores still run at FP16 rate. That sounds pointless until you remember Part 3: at small batch the decode step can be bandwidth-bound, so 4-bit weights mean up to 4× less weight traffic. Its dequantization cost can be hidden while arithmetic resources and tensor-core throughput have headroom; it is not literally free in every kernel.

W8A8 (FP8 or INT8) quantizes activations too, so the tensor cores genuinely run at 2× the FP16 rate. That is what helps prefill and large-batch decode. Its cost is activation outliers — exactly the problem §8.2 is about.

AWQ's central observation is worth stating on its own: "protecting only 1% salient weights can greatly reduce quantization error" — and the salient channels are identified by activation magnitude, not weight magnitude. Since it needs no backpropagation or reconstruction, it avoids the reconstruction-overfitting failure mode the paper contrasts with GPTQ.

1. Prefer FP32 accumulation for accuracy-sensitive reductions wherever the hardware path provides it, regardless of the storage type. At FP8 on Hopper, remember that an FP32 accumulator interface does not guarantee full internal FP32 accumulation; periodic promotion is one measured remedy.
2. Quantize the KV cache before you quantize the weights, if you serve long contexts. Part 7 showed KV traffic overtaking weight traffic; FP8 KV halves the dominant term.
3. Pick the granularity from the tensor, not from a default. Weights are well-behaved and tolerate coarse blocks; activations have outliers and want 1×128 tiles or finer; activation gradients want tiles specifically.
4. Fuse quantization into the producing kernel's epilogue. Computing a per-block amax and scaling costs a reduction over data that is already in registers. As a separate pass it costs two full trips to HBM.
5. Measure end-task accuracy, not RMSE. The FlashAttention-3 ablation in Part 7 is the cautionary example: two techniques were credited for the FP8 accuracy result, and the ablation shows one of them was doing essentially all the work.

One part left: how to tell whether any of this actually helped.