Number formats: BF16, FP8, FP4

Code version: DeepInfer @ 9bd9cf2 · 2026-10-21

The idea#

Why fewer bits?#

Every weight, activation and KV cache entry is a number in device memory. The dtype of a tensor sets how many bits each number uses. Fewer bits help in three ways:

• Memory. A model in 8-bit numbers needs half the memory of the same model in 16-bit numbers.
• Bandwidth. Decode is memory-bound (see GPUs, kernels, and speed). Half the bytes means about half the time.
• Math. Tensor Cores do more FP8 or FP4 operations per second than BF16 operations.

The cost is accuracy. A small format has fewer possible values, so most numbers must be rounded. Engineers make this rounding small with quantization methods, mainly with scales.

Sign, exponent and mantissa#

A floating-point number has three bit fields:

value = (−1)^sign × 1.mantissa × 2^(exponent − bias)

sign      1 bit: + or −
exponent  e bits: the power of two (sets the range)
mantissa  m bits: the digits after the binary point (sets the precision)

More exponent bits give a larger range: very large and very small numbers. More mantissa bits give finer steps between neighbour values. When the exponent field is 0, the number is subnormal: 0.mantissa × 2^(1 − bias). Subnormals fill the gap between 0 and the smallest normal number.

The formats on this page#

BF16 has the same exponent as FP32, so it has the same range with less precision. This is why models train and run in BF16 without overflow. FP16 is more precise, but its range ends at 65,504.

FP8 has two common layouts. E4M3 has more precision; E5M2 has more range. Inference engines usually use E4M3 for weights and activations. FP4 E2M1 has only 16 codes. Its positive values are 0, 0.5, 1, 1.5, 2, 3, 4 and 6.

A worked conversion: 0.3#

0.3 = 1.2 × 2⁻². So the exponent is −2 and the mantissa is 0.2. Each format rounds 0.2 to its mantissa bits:

The exponent field of E4M3 is −2 + 7 = 5 = 0101. The exponent field of BF16 is −2 + 127 = 125 = 01111101.

Scales: keep small formats accurate#

E2M1 cannot store 0.3 well on its own. But if we first divide by a scale factor, the value can land near a code. Divide 0.3 by 0.0625: the result is 4.8, and the nearest code is 4. Then 4 × 0.0625 = 0.25, an error of 17 % instead of 67 %.

The scale must fit the numbers that share it. With one scale for a whole tensor, one large value forces a large scale on all other values. So modern formats use block scaling: a small block of values (for example 16 or 32) shares one scale.

• MXFP4: 32 E2M1 values + 1 UE8M0 scale = 17 bytes, 4.25 bits per value.
• MXFP8: 32 FP8 values + 1 UE8M0 scale = 33 bytes, 8.25 bits per value.
• FP4 with an E4M3 scale per 16 values: 9 bytes per 16 values, 4.5 bits per value.

A UE8M0 scale holds only a power of two: 2^(code − 127). Multiplying by a power of two changes only the exponent, so it adds no rounding error of its own. An E4M3 scale can be any FP8 value, so it fits the block better, but it costs a few more bits per value.

Try it. The interactive number-format converter shows one number in BF16, FP8 E4M3 and FP4 E2M1. It also encodes a block of up to 32 values with one shared UE8M0 scale. Add an outlier and see how many small values become 0.

An everyday comparison#

Think of a price list with the header "all prices in thousands of euros". Each row needs only a short number, such as 4 or 1.5, and the header gives the size. A block scale is that header. A list with very different prices (5 euros and 5 million euros) needs several headers, one per small group of rows.

The comparison stops here: a price list can write any short number, but an FP4 row can only write one of 16 codes. And a UE8M0 header can only be a power of two, such as "in units of 4" or "in units of 8".

The picture#

Read panel 1 from top to bottom: the strips get shorter, and the mantissa (amber) shrinks first. Panel 2 shows the price: the error grows as the mantissa shrinks.

In DeepInfer#

This is a basics page, so this part only gives pointers. The page Low-precision formats in DeepInfer explains the full code.

1. The checkpoint must be FP8 with FP4 experts. The config loader accepts only one quantization setup, and turns it into the flags fp8_attn, fp8_mlp, mxfp8 and mxfp4:

    quant = hf["quantization_config"]
    for key, expected in {
        "quant_method": "fp8",
        "expert_dtype": "fp4",
        "activation_scheme": "dynamic",
        "weight_block_size": [32, 32],
        "scale_fmt": "ue8m0",
    }.items():
        if quant[key] != expected:
            raise ValueError(f"Unsupported HF {key}: {quant[key]!r}; expected {expected!r}")
    fp8 = quant["quant_method"] == "fp8"
    mxfp4 = quant["expert_dtype"] == "fp4"

2. The KV cache has its own formats. The main cache uses FP4 with an E4M3 scale per 16 values (288 bytes per entry) by default, or MXFP8 (528 bytes) (deepinfer/model_executor/layers/tilelang_ops/cache_common.py:109-128, deepinfer/hugging_face/config.py:303). See The KV cache.
3. Scales are packed per platform. FP8 GEMMs take UE8M0 scales packed into torch.int32 on NVIDIA and torch.int16 on Ascend (deepinfer/platform.py:63, deepinfer/platform.py:138). See Hardware platforms: NVIDIA and Ascend.

Deep dive#

The sections below give the full detail. Open them one at a time, or use Open all deep dives at the top of the page.

def _ue8m0_to_fp32(scale: torch.Tensor) -> torch.Tensor:
    """F8_E8M0 (exponent only) → fp32: value = 2 ** (bits - 127), lossless restore."""
    assert scale.dtype == torch.float8_e8m0fnu, f"expected F8_E8M0, got {scale.dtype}"
    exponents = scale.view(torch.uint8).to(torch.int32) - 127
    return torch.ldexp(torch.ones_like(exponents, dtype=torch.float32), exponents)

weight  W  : float8_e4m3fn  [N, K]
scale   sf : float32        [ceil(N/32), ceil(K/32)]   (one scale per 32 × 32 block)

    def apply_weight_fp8(
        self,
        x: torch.Tensor | tuple[torch.Tensor, torch.Tensor],
        bias: torch.Tensor | None = None,
        out_acc: torch.Tensor | None = None,
    ) -> torch.Tensor:
        assert bias is None

        if isinstance(x, torch.Tensor):
            assert x.dim() == 2, f"{x.shape=}, {x.stride()=}"
            x = tile_kernels.quant.per_token_cast(x, **get_fp8_cast_params())

        # FP8 GEMM.
        weight = (self.weight, self.weight.sf)
        out_shape = x[0].shape[:-1] + (self.weight.shape[0],)
        if out_acc is None:
            out = self._allocate_output(out_shape, torch.bfloat16, x[0].device)
        else:
            out = out_acc
            assert out.shape == out_shape
            assert out.dtype == torch.bfloat16
            assert out.device == x[0].device

        assert self._has_aligned_output_stride(out), f"unaligned output: {out.shape=}, {out.stride()=}, {out.dtype=}"
        assert x[1].dtype == weight[1].dtype == PLATFORM.fp8_sf_type, f"{x[1].dtype=}, {weight[1].dtype=}"
        deep_gemm.fp8_fp4_gemm_nt(x, weight, out, out_acc, recipe=(1, 1, get_fp8_casting_granularity()))
        return out

x       BF16    [1536, 4096]      12 MiB
x_fp8   E4M3    [1536, 4096]       6 MiB
scales  UE8M0   [1536, 4096/32] = [1536, 128] codes
        packed: [1536, 32] int32 (NVIDIA, 4 codes each) or [1536, 64] int16 (Ascend, 2 codes each)

What can go wrong#

• A checkpoint in another format. The config loader raises ValueError("Unsupported HF ...") if the quantization method, expert dtype, activation scheme, block size or scale format differ (deepinfer/hugging_face/config.py:254-262). A model dtype other than BF16 also raises (deepinfer/hugging_face/config.py:58-59).
• FP4 experts from a plain FP8 checkpoint. The dummy FP4 path asserts that the FP8 bytes are FP4-compatible: "Only support FP4-QAT checkpoint" (deepinfer/model_executor/models/v4_modules/mlp.py:60). QAT means that training already used FP4.
• Wrong scale dtype. An FP8 GEMM asserts that both scale tensors have the platform type, int32 or int16 (deepinfer/model_executor/layers/linear.py:164).
• Overflow and outliers. E4M3 ends at 448 and E2M1 at 6. One very large value in a block makes the scale large, and the small values of the same block round to 0. Block scaling limits this damage to one block.

Why it matters. The format of a tensor sets its memory, its bandwidth and the speed of the math on it. In DeepInfer, FP4 experts need about a quarter of the BF16 memory. FP8 activations halve the GEMM input traffic. The FP4 KV cache lets more sequences run at the same time.