Numerics: FP32 to INT4, and why bf16 won

§0.4 ended with an uncomfortable conclusion: a decode step reads every weight once and does almost nothing with it, so the step time is bytes moved ÷ HBM bandwidth and nothing else. The only lever with real leverage is bytes per parameter. Halve it, halve the floor.

So you halve it. And sometimes the model is fine, and sometimes it emits fluent nonsense, and sometimes it emits nan on token 400 of a 2000-token generation and only on long prompts. The failure is silent by construction: a format that cannot represent your number does not raise — it saturates, or rounds to zero, or rounds to inf and poisons a softmax.

vLLM's config layer has a hard-coded list of models that must not be run in fp16 at all:

# model_type -> reason
_FLOAT16_NOT_SUPPORTED_MODELS = {
    "gemma2": "Numerical instability. Please use bfloat16 or float32 instead.",
    "gemma3": "Numerical instability. Please use bfloat16 or float32 instead.",
    "gemma3_text": "Numerical instability. Please use bfloat16 or float32 instead.",
    "glm4": "Numerical instability. Please use bfloat16 or float32 instead.",
}


def _is_valid_dtype(model_type: str, dtype: torch.dtype):
    if model_type in _FLOAT16_NOT_SUPPORTED_MODELS and dtype == torch.float16:  # noqa: E501, SIM103
        return False

    return True


def _check_valid_dtype(model_type: str, dtype: torch.dtype):
    if model_type in _FLOAT16_NOT_SUPPORTED_MODELS and dtype == torch.float16:
        reason = _FLOAT16_NOT_SUPPORTED_MODELS[model_type]
        raise ValueError(
            f"The model type {model_type!r} does not support float16. Reason: {reason}"
        )

fp16 and bf16 are both 16 bits. They cost identical bandwidth. One of them is on a denylist. That difference is entirely a question of where the bit boundary sits, and it generalises to every format below 16 bits too.

A floating-point number is sign × 2^exponent × 1.mantissa. The exponent bits buy you dynamic range — how far apart the largest and smallest representable magnitudes are. The mantissa bits buy you precision — how finely you can resolve values within one power of two. These are independent budgets, and a format is just a choice of how to split a fixed number of bits between them.

Transformers are asymmetric consumers of these two budgets. They need range, because activations in a deep residual stream span many orders of magnitude and a handful of outlier channels run far above the rest. They need surprisingly little precision, because every value is about to be summed with thousands of others in a dot product and the summation error dominates the per-element rounding error. bf16 won because it spends its 16 bits exactly that way: FP32's entire exponent, and whatever mantissa is left over.

The dashed line is the whole argument. FP32, TF32 and BF16 all put the mantissa boundary at the same place; they have identical exponent widths but slightly different maximum finite values because their mantissas differ. FP16 and FP8-E5M2 stop the exponent early to buy mantissa, and pay for it at the top end.

Let $e$ be the number of exponent bits and $m$ the number of mantissa bits. IEEE-754 stores the exponent biased by $B = 2^{e-1} - 1$, reserves the all-ones exponent field for infinity/NaN, and reserves the all-zeros field for zero and subnormals. So the largest representable finite magnitude is

and the smallest positive normal is $2^{1-B}$. The gap between $1.0$ and the next representable value — machine epsilon — is $2^{-m}$, independent of $e$. That is the clean statement of the two budgets: $e$ alone sets the range, $m$ alone sets the relative precision.

Two of the formats break the IEEE rules deliberately. FP8-E4M3 as standardised by OCP (float8_e4m3fn: finite, naN) has no infinities, so it reclaims the all-ones exponent field and only burns the single mantissa pattern 111 for NaN. That buys it one extra binade. FP4-E2M1 has neither infinity nor NaN. Running the formula with $e{=}4, m{=}3$ and the extended-range rule gives $(1 + 6/8)\cdot 2^{15-7} = 1.75 \cdot 256 = 448$.

You do not have to trust my arithmetic, because vLLM encodes the layouts as data and derives the same numbers at import time:

class scalar_types:
    int4 = ScalarType.int_(4, None)
    uint4 = ScalarType.uint(4, None)
    int8 = ScalarType.int_(8, None)
    uint8 = ScalarType.uint(8, None)
    float8_e4m3fn = ScalarType.float_(4, 3, True, NanRepr.EXTD_RANGE_MAX_MIN)
    float8_e5m2 = ScalarType.float_IEEE754(5, 2)
    float8_e8m0fnu = ScalarType(8, 0, False, 0, True, NanRepr.EXTD_RANGE_MAX_MIN)
    float16_e8m7 = ScalarType.float_IEEE754(8, 7)
    float16_e5m10 = ScalarType.float_IEEE754(5, 10)

Note the naming: bf16 is literally spelled float16_e8m7 and aliased to bfloat16 at the bottom of the class. The two 16-bit types differ in this file by three characters. ScalarType._floating_point_max() (vllm/scalar_type.py:L71-L105) reconstructs the max by packing the biased exponent and mantissa into an IEEE double and bit-casting — the same derivation, executed. Copy the file somewhere and run it; it has no vLLM imports:

float8_e4m3fn   max= 448.0
float8_e5m2     max= 57344.0
float16_e8m7    max= 3.3895313892515355e+38     # bf16
float16_e5m10   max= 65504.0                    # fp16
int4            max= 7          min= -8
int8            max= 127        min= -128

*TF32 occupies a 32-bit register; only 19 bits are numerically meaningful. It is a tensor-core input format, never a storage format, which is why it does not appear in scalar_types.

Why fp16 breaks and bf16 does not

bf16 shares FP32's exponent width, not every finite endpoint: its maximum is approximately $3.3895\times10^{38}$, slightly below FP32's $3.4028\times10^{38}$. For a finite result rounding to a normal bf16 value, round-to-nearest relative error is bounded by $2^{-8}$; that bound excludes underflow, overflow, and nonfinite inputs. Its broad range usually removes the need for fp16-style loss scaling, but does not eliminate numerical validation.

fp16 tops out at 65504, i.e. $2^{16}$. Transformer residual streams do not stay under $2^{16}$: Dettmers et al. document emergent outlier features whose magnitudes are up to 20× the rest of the hidden state and which appear in essentially every sequence past roughly 6.7B parameters (LLM.int8(), arXiv:2208.07339). Gemma-2 compounds this by design: it multiplies embeddings by $\sqrt{d}$, which raises their magnitude. Its attention soft-cap bounds logits instead of increasing their magnitude. Once one hidden unit hits inf, the next RMSNorm divides by an infinite variance and the entire row becomes nan — which is exactly what that four-word reason string in vllm/config/model.py is protecting you from.

The tie back to the roofline

Llama-3-8B has 8.03×10⁹ parameters, but a decode step streams only 7.50×10⁹ of them: the 128,256×4096 input embedding is a gather of one row, not a stream, and §0.4 counts it out for exactly this reason. On an H100 SXM (3.35 TB/s HBM3), the weight-streaming floor per decode step, from the memory-bound branch of that roofline, is $t = P_{\text{stream}} \cdot b / \mathrm{BW}$ with $b$ = bytes per parameter:

Derived arithmetic, single-request decode, weights only — KV traffic, activations and kernel launch overhead are excluded, so treat these as lower bounds on step time, not predictions. Note the int4 figure includes the group scales: at group size 128 in fp16 you store an extra $16/128 = 0.125$ bits per weight, so $b = 0.5 + 0.016 = 0.516$, not $0.5$. The bf16 row is the same 4.48 ms floor §0.4 derives, which is the point: this chapter changes $b$ and nothing else. Whether the model still works at 4 bits is the subject of §4.4; whether it could go that fast is settled right here.

This is the point people get wrong most often, so it gets its own section. The dtype you store a tensor in and the dtype the hardware sums products in are different axes.

A tensor core does not compute an fp8 dot product. It multiplies fp8 operands and accumulates the partial sums into an FP32 register file. The same is true for bf16 inputs, fp16 inputs and int8 inputs (which accumulate into int32). You can read this directly in vLLM's Triton block-scaled FP8 GEMM: the accumulator is declared fp32, the per-block dequant scales are folded in at fp32, and the narrowing cast happens exactly once, at the end, on the way to global memory.

    accumulator = tl.zeros((BLOCK_SIZE_M, BLOCK_SIZE_N), dtype=tl.float32)
    for k in range(0, tl.cdiv(K, BLOCK_SIZE_K)):
        a = tl.load(a_ptrs, mask=offs_k[None, :] < K - k * BLOCK_SIZE_K, other=0.0)
        b = tl.load(b_ptrs, mask=offs_k[:, None] < K - k * BLOCK_SIZE_K, other=0.0)

        k_start = k * BLOCK_SIZE_K
        offs_ks = k_start // group_k
        a_s = tl.load(As_ptrs + offs_ks * stride_As_k)
        b_s = tl.load(Bs_ptrs + offs_ks * stride_Bs_k)

        accumulator += tl.dot(a, b) * a_s[:, None] * b_s[None, :]
        a_ptrs += BLOCK_SIZE_K * stride_ak
        b_ptrs += BLOCK_SIZE_K * stride_bk

    if C.dtype.element_ty == tl.bfloat16:
        c = accumulator.to(tl.bfloat16)
    elif C.dtype.element_ty == tl.float16:
        c = accumulator.to(tl.float16)
    else:
        c = accumulator.to(tl.float32)

For a concrete reduction, serial bf16 addition of 4096 ones stalls at 256 under round-to-nearest-even: at 256 the spacing is 2, so adding 1 ties and rounds back to 256. FP32 represents this particular integer sum exactly. There is no universal loss of 12 bits merely because a dot product has 4096 terms. General error depends on conditioning, order, cancellation, and accumulator precision. For unit roundoff $u$ and $(n-1)u<1$, a standard serial bound is $|\widehat s-s|\le\gamma_{n-1}\sum_i|x_i|$, where $\gamma_k=ku/(1-ku)$. Pairwise reduction shortens dependency depth; it does not repair already-rounded inputs.

The same rule explains why the un-quantisable parts of a transformer are un-quantisable. SGLang's reference RMSNorm upcasts unconditionally before computing the variance, because a mean of squares over $d$ elements is a reduction:

        if not x.is_contiguous():
            x = x.contiguous()
        orig_dtype = self.override_orig_dtype or x.dtype
        x = x.to(torch.float32)
        if residual is not None:
            x = x + residual.to(torch.float32)
            if post_residual_addition is not None:
                x = x + post_residual_addition.to(torch.float32)
            if self.fp32_residual:
                residual = x.clone()

And vLLM's sampler upcasts the logits before any processor touches them, because softmax over a 128k-entry vocabulary is a reduction whose result is then exponentiated. Both samplers do it — the V2 model runner's copies into a fresh fp32 tensor at vllm/v1/worker/gpu/sample/sampler.py:L209-L210, the V1 one casts in place:

        # Use float32 for the logits.
        logits = logits.to(torch.float32)

Floats carry their own scale in the exponent field. Integers do not, so an integer format is only half a format — the other half is an affine map back to the reals:

with $s$ the scale (a float) and $z$ the zero-point (an integer). $z = 0$ is symmetric quantisation; $z \neq 0$ is asymmetric, and buys you the ability to represent a range that is not centred on zero at the cost of an extra subtraction in the inner loop. That subtraction is literally one line in vLLM's dequantisation kernel:

    # Load the scales.
    scales = tl.load(scales_ptr + scale_offsets, scale_masks, 0.0)
    scales = tl.broadcast_to(scales, (BLOCK_SIZE_Y, BLOCK_SIZE_X * 8))

    # Dequantize.
    iweights = (iweights - zeros) * scales
    iweights = iweights.to(result_ptr.type.element_ty)

Both integer and floating low-bit formats need a scale policy. Symmetric INT8 can set $z=0$ and compute $s=\max|x|/127$ dynamically, just as FP8 can derive a scale from a live amax. Asymmetric integer quantization additionally learns or computes a zero point. Either family may use static calibration or dynamic scales; the choice depends on outliers, granularity, kernels, and overhead, not simply integer versus float. Define an all-zero-group policy such as $s=1$, validate nonfinite values, and distinguish clipping from rounding error.

See the PyTorch numerical accuracy notes for reduced-precision reduction controls. CPU reference checks do not certify a GPU kernel or its accumulation mode.

Granularity: one scale for how many numbers?

A single scale for a whole 4096×4096 weight matrix is set by its single largest element, so one outlier costs every other element its precision. Narrowing the sharing group is the fix. vLLM names the axes explicitly:

    def is_per_tensor(self) -> bool:
        return self.row == -1 and self.col == -1

    def is_per_token(self) -> bool:
        return self.row == 1 and self.col == -1

    def is_per_channel(self) -> bool:
        return self.row == -1 and self.col == 1

    def is_per_group(self) -> bool:
        return self.row == 1 and self.col >= 1


GroupShape.PER_TENSOR = GroupShape(-1, -1)
GroupShape.PER_TOKEN = GroupShape(1, -1)
GroupShape.PER_CHANNEL = GroupShape(-1, 1)

Per-tensor is one scale, free to store, worst quality. Per-channel (one scale per output column of the weight) and per-token (one scale per row of the activation) cost one float per row or column and compose cleanly with a GEMM, because both scales factor out of the inner product. Per-group narrows further — 128 or 64 contiguous elements along $K$ — and no longer factors out, which is why grouped kernels must fold the scale inside the K-loop, exactly as the Triton kernel above does with a_s and b_s.

MXFP4 and the block-scaled family

Push per-group to its logical extreme and you get the OCP Microscaling formats: a fixed block of 32 elements shares one scale, and the scale itself is compressed to a bare power of two — an 8-bit exponent with no sign and no mantissa, float8_e8m0fnu in the scalar_types listing above. vLLM hard-codes the block size:

OCP_MX_BLOCK_SIZE = 32

OCP_MX_DTYPES = {
    "mxfp4",
    "mxfp6_e3m2",
    "mxfp6_e2m3",
    "mxfp8_e4m3",
    "mxfp8_e5m2",
    "mxint8",
}

MXFP4 therefore costs $4 + 8/32 = 4.25$ bits per weight. The element format is FP4-E2M1, which has exactly eight representable magnitudes — SGLang writes them out as a literal:

class MXFP4QuantizeUtil:
    E2M1_max = 6.0

    E2M1_values = [0, 0.5, 1, 1.5, 2, 3, 4, 6]
    E2M1_bounds = torch.tensor([0.25, 0.75, 1.25, 1.75, 2.5, 3.5, 5])

Sixteen codes total including sign. The quantiser is then obvious: take the amax of each block of 32, divide by 6.0, round the result up to a power of two, and store the exponent biased by 127.

        original_shape = input.shape
        original_dtype = input.dtype
        input = input.view(-1, block_size)
        # get scales
        input_amax = input.abs().max(dim=-1, keepdim=True).values
        descale = input_amax / cls.E2M1_max
        min_value = torch.tensor(-127.0, device=descale.device)
        e8m0_scale = torch.ceil(torch.maximum(torch.log2(descale), min_value))

        input = (input / torch.exp2(e8m0_scale)).view(original_shape)
        input_q = cast_fp4(input)
        input_q = fuse_uint4_to_uint8(input_q)
        e8m0_scale = (e8m0_scale + 127).to(torch.uint8)
        return cls(original_shape, original_dtype, input_q), e8m0_scale

The ceil is deliberate: rounding the scale up guarantees no element of the block clips, at the cost of up to one binade of headroom. NVFP4 makes the opposite tradeoff — block of 16, scale stored in FP8-E4M3 rather than E8M0, so the scale itself has mantissa bits and can land between powers of two. Both appear in the registries below.

The most honest description of the state of the art is the list of things each engine can load. vLLM keeps it as a Literal so the type checker enforces it:

QuantizationMethods = Literal[
    "awq",
    "auto_awq",
    "fp8",
    "fbgemm_fp8",
    "fp_quant",
    "modelopt",
    "modelopt_fp4",
    "modelopt_mxfp8",
    "modelopt_mixed",
    "auto_gptq",
    "gptq",
    "gptq_marlin",
    "awq_marlin",
    "humming",
    "compressed-tensors",
    "experts_int8",
    "quark",
    "moe_wna16",
    "torchao",
    "inc",
    "mxfp4",
    "gpt_oss_mxfp4",
    "deepseek_v4_fp8",
    "online",
    # Below are online quant shorthand names (see vllm.config.quantization).
    # Listed here as strings to avoid a circular import; kept in sync with
    # _ONLINE_SHORTHANDS by the assertion in get_quantization_config().
    "fp8_per_tensor",
    "fp8_per_block",
    "fp8_per_channel",
    "int8_per_channel_weight_only",
    "nvfp4_per_token",
    "mxfp8",
]

SGLang keeps the equivalent as a dict from name to config class, with platform-conditional additions layered on afterwards:

BASE_QUANTIZATION_METHODS: Dict[str, Type[QuantizationConfig]] = {
    "fp8": Fp8Config,
    "mxfp8": Fp8Config,
    "blockwise_int8": BlockInt8Config,
    "modelopt": ModelOptFp8Config,  # Auto-detect, defaults to FP8
    "modelopt_fp8": ModelOptFp8Config,
    "modelopt_fp4": ModelOptFp4Config,
    "nvfp4_online": NvFp4OnlineConfig,
    "modelopt_mixed": ModelOptMixedPrecisionConfig,
    "w8a8_int8": W8A8Int8Config,
    "w8a8_fp8": W8A8Fp8Config,
    "awq": AWQConfig,
    "awq_marlin": AWQMarlinConfig,
    "bitsandbytes": BitsAndBytesConfig,
    "gguf": GGUFConfig,
    "gptq": GPTQConfig,
    "gptq_marlin": GPTQMarlinConfig,
    "moe_wna16": MoeWNA16Config,
    "compressed-tensors": CompressedTensorsConfig,
    "w4afp8": W4AFp8Config,
    "petit_nvfp4": PetitNvFp4Config,
    "quark": QuarkConfig,
    "quark_mxfp4": QuarkConfig,
    "auto-round": AutoRoundConfig,
    "auto-round-int8": W8A8Int8Config,
    "modelslim": ModelSlimConfig,
    "quark_int4fp8_moe": QuarkInt4Fp8Config,
    "humming": HummingConfig,
    "mxfp_w4a8": Mxfp4W4A8Config,
}

Where they differ, and why. The overlap is nearly total on the checkpoint formats that matter — fp8, GPTQ/AWQ (+ Marlin kernels), compressed-tensors, ModelOpt, Quark, MXFP4. The divergence is in how the tail is organised. vLLM has moved toward a single compressed-tensors entry point plus "online" shorthands (fp8_per_tensor, fp8_per_block, fp8_per_channel, nvfp4_per_token) that quantise an unquantised checkpoint at load time; the granularity is part of the method name. SGLang instead spells out separate config classes per scheme (w8a8_int8, w8a8_fp8, w4afp8, blockwise_int8, mxfp_w4a8) and registers backend-specific entries conditionally per platform. Same formats, different factoring: vLLM pushes the variation into config strings, SGLang into classes. Neither list is a quality ranking — a name here means "we can load and run it", nothing about accuracy.

The KV cache is its own dtype axis

Weights and activations are chosen by --dtype; the KV cache has a separate flag because it has separate constraints. It is written once and read thousands of times, it is never a GEMM operand in the usual sense, and it dominates memory at long context. SGLang's resolver is the cleanest statement of what is legal:

    elif server_args_kv_cache_dtype == "fp8_e5m2":
        if _is_hip:  # Using natively supported format
            kv_cache_dtype = fp8_dtype
        else:
            kv_cache_dtype = torch.float8_e5m2
    elif server_args_kv_cache_dtype == "fp8_e4m3":
        if _is_hip:  # Using natively supported format
            kv_cache_dtype = fp8_dtype
        else:
            kv_cache_dtype = torch.float8_e4m3fn
    elif server_args_kv_cache_dtype == "mxfp8":
        kv_cache_dtype = torch.float8_e4m3fn
    elif server_args_kv_cache_dtype in ("bf16", "bfloat16"):
        kv_cache_dtype = torch.bfloat16
    elif server_args_kv_cache_dtype == "fp4_e2m1":
        raise ValueError(
            "--kv-cache-dtype=fp4_e2m1 is deprecated. "
            "Use --kv-cache-dtype=fp4_mx_block16."
        )
    elif server_args_kv_cache_dtype in ("nvfp4", "fp4_mx_block16"):
        if hasattr(torch, "float4_e2m1fn_x2"):
            kv_cache_dtype = torch.float4_e2m1fn_x2
            logger.warning(
                "%s KV Cache might lead to an accuracy drop!",
                server_args_kv_cache_dtype.upper(),
            )
        else:
            raise ValueError(
                f"--kv-cache-dtype={server_args_kv_cache_dtype} requires "
                "torch.float4_e2m1fn_x2 support. Please use PyTorch 2.8.0+ "
                "with CUDA 12.8+."
            )

Note what is absent: there is no fp16. The four float options are bf16, fp8-e4m3, fp8-e5m2 and packed fp4, and only fp4 gets a warning. vLLM's CacheDType literal (vllm/config/cache.py:L19-L37) is longer and does include "float16", plus a set of per-token-head integer modes that SGLang does not expose, enumerated in the V1 attention interface:

class KVQuantMode(IntEnum):
    """KV cache quantization mode.

    Used by attention backends and kernels to dispatch quantization logic
    without string matching on ``kv_cache_dtype``.
    """

    NONE = 0
    FP8_PER_TENSOR = 1  # per-tensor scales (current fp8 path)
    INT8_PER_TOKEN_HEAD = 2  # per-token-head dynamic scales for int8
    FP8_PER_TOKEN_HEAD = 3  # per-token-head dynamic scales for fp8
    INT4_PER_TOKEN_HEAD = 4  # packed 2×int4/byte, RHT + asymmetric zp
    NVFP4 = 5  # packed fp4 data + fp8 block scales

The comment on mode 4 is a compressed statement of everything in this chapter: to survive at four bits, a KV entry needs a random Hadamard transform to spread outliers across channels, an asymmetric zero-point, and per-token-per-head scale granularity. Three separate mitigations, all for the same problem — four bits of integer cannot represent an outlier-heavy distribution.

Below is where each dtype is legal in a transformer block. The rule that generates it: an operation whose output feeds a long reduction, an exponential, or a division must run wider than its inputs.

Follow a single row of the hidden state, shape [1, 4096] in bf16, into an fp8 linear layer in vLLM.

1. QuantFP8.forward_native() — vllm/model_executor/layers/quantization/input_quant_fp8.py:L199-L218. With GroupShape.PER_TOKEN, it takes the amax over the last axis, promotes it to fp32, divides by _FP8_MAX (448.0 for e4m3) to get the scale, clamps the scale away from zero, then divides the row by the scale in fp32 and clamps into range before the narrowing cast:
   vllm/model_executor/layers/quantization/input_quant_fp8.py:L199-L218 vLLM

           if scale is None:
               if self.group_shape == GroupShape.PER_TOKEN:
                   x_max, _ = x.abs().max(dim=-1)
                   x_max = x_max.unsqueeze(-1).to(torch.float32)
                   if scale_ub is not None:
                       x_max = x_max.clamp(max=scale_ub)
               else:
                   x_max = x.abs().max().unsqueeze(-1).to(torch.float32)
   
               scale = (x_max / _FP8_MAX).clamp(min=_FP8_MIN_SCALING_FACTOR)
           else:
               scale = prep_scale_for_group_broadcast(scale, x, self.group_shape)
   
           # Even for dynamic per-token scales,
           # reciprocal performs slightly better than division
           out = (
               x.to(torch.float32)
               * group_broadcast(scale.to(torch.float32), x.shape[-2:]).reciprocal()
           )
           out = out.clamp(_FP8_MIN, _FP8_MAX).to(_FP8_DTYPE)

   Concretely: if the row's amax is 22.4, the scale is $22.4/448 = 0.05$, and every element is multiplied by 20 before the cast. The largest element lands exactly on 448 — full use of the format's range. This is the entire trick. Without it, an element of magnitude 0.0005 would sit below half of fp8-e4m3's smallest subnormal ($2^{-9} = 0.00195$) and round to zero; 0.001, just above that half-way point, survives — as the single value 0.00195, a 95% relative error. Unscaled, the format's dynamic range simply does not overlap the data.
2. The CUDA path does the same thing. csrc/quantization/w8a8/fp8/common.cuh:L57-L77 — the scalar helper every fp8 quant kernel calls. Scale application, saturating clamp, hardware convert, in that order, with the arithmetic in float:
   csrc/quantization/w8a8/fp8/common.cuh:L57-L77 vLLM

   template <bool is_scale_inverted, typename fp8_type>
   __device__ __forceinline__ fp8_type scaled_fp8_conversion(float const val,
                                                             float const scale) {
     float x = 0.0f;
     if constexpr (is_scale_inverted) {
       x = val * scale;
     } else {
       x = val / scale;
     }
   
     float r =
         fmaxf(-quant_type_max_v<fp8_type>, fminf(x, quant_type_max_v<fp8_type>));
   #ifndef USE_ROCM
     // Use hardware cvt instruction for fp8 on nvidia
     // Currently only support fp8_type = c10::Float8_e4m3fn
     return fp8::vec_conversion<fp8_type, float>(r);
   #else
     // Use hardware cvt instruction for fp8 on rocm
     return fp8::cvt_c10<fp8_type>(r);
   #endif
   }

   The fmaxf(-max, fminf(x, max)) is a saturating clamp, not a wrap. An activation that exceeds the range does not become inf; it silently becomes 448. That is the design choice that makes fp8 forgiving and also makes its failures invisible.
3. The GEMM. Quantised row (fp8) × quantised weight (fp8), fp32 accumulator, both scales folded in inside the K-loop, single narrowing cast to bf16 on store — vllm/model_executor/layers/quantization/utils/fp8_utils.py:L793-L812, quoted in §0.5.4. Two dtypes crossed the wire; the maths happened in a third.
4. Out. Result is bf16, ready for the residual add. Net effect on the roofline: the weight tensor for this layer moved 1 byte per element instead of 2. The activation row is negligible at batch 1 — it is 4096 elements against 4096×4096 weights.

The KV cache scale that defaulted to 1.0

Run a model with --kv-cache-dtype fp8_e4m3 against a checkpoint that carries no k_scale/v_scale, and vLLM does not refuse. It uses 1.0:

            elif layer.k_scale < 0.0 and layer.v_scale < 0.0:
                # If no scales were loaded (both scales are invalid negative
                # values), use the default value of 1.0
                k_scale = 1.0
                v_scale = 1.0
# ...
            if k_scale == 1.0 and v_scale == 1.0 and "e5m2" not in layer.kv_cache_dtype:
                logger.warning_once(
                    "Using KV cache scaling factor 1.0 for fp8_e4m3. "
                    "If this is unintended, verify that k/v_scale "
                    "scaling factors are properly set in the checkpoint."
                )

With scale 1.0, an e4m3 KV cache clamps everything above 448 and rounds magnitudes below half the smallest subnormal, $2^{-10}\approx0.000977$, to zero under round-to-nearest-even (the exact tie also rounds to zero). K and V values typically live well inside that window, so the model does not crash and often does not obviously degrade — it just quietly loses the tails. The warning is warning_once, so it scrolls past on startup and never reappears. This is the canonical silent-corruption bug of the area: the format is legal, the flag is accepted, and the one signal you get is a single log line.

Note the guard "e5m2" not in layer.kv_cache_dtype. E5M2 has fp16's exponent field, so it reaches 57344 and unity scale has more range headroom but is not automatically accurate — you trade the extra range for having only two mantissa bits. That is the whole e4m3-vs-e5m2 decision in one conditional: check the scale and accuracy for either format; e5m2 trades precision for greater range.

Mixing fp16 and bf16 checkpoints

vLLM will let you do it, with a log line and no error (vllm/config/model.py:L2347-L2351: logger.warning("Casting %s to %s.", ...)). Casting bf16 → fp16 is the dangerous direction: any weight or buffer above 65504 becomes inf at load time, before you have run a single token.

Symptom-to-cause table

Everything below runs on CPU. No GPU is needed to explore the formats themselves.

V=~/Documents/other_git_repos/vllm
cp "$V"/vllm/scalar_type.py /tmp/st.py
python3 -c "
import sys; sys.path.insert(0,'/tmp')
from st import scalar_types as s
for n in ['float8_e4m3fn','float8_e5m2','float16_e8m7','float16_e5m10','int4','int8']:
    t = getattr(s, n); print(f'{n:16s} bits={t.size_bits:3d} max={t.max()} min={t.min()}')
"

# What does a bf16 round-trip actually cost?
python3 -c "
import torch
x = torch.randn(4096, dtype=torch.float32)
for dt in (torch.bfloat16, torch.float16, torch.float8_e4m3fn, torch.float8_e5m2):
    y = x.to(dt).float()
    print(f'{str(dt):24s} max_rel_err={( (y-x).abs()/x.abs() ).max():.4f}')
"

# Reproduce the overflow: bf16 survives, fp16 does not.
python3 -c "
import torch
big = torch.tensor([70000.0])
print('bf16 ->', big.to(torch.bfloat16).item())
print('fp16 ->', big.to(torch.float16).item())
"

On a GPU, the flag to flip is the KV cache dtype — it is the one dtype knob that changes memory without touching the weights:

vllm serve meta-llama/Llama-3.1-8B-Instruct --kv-cache-dtype fp8_e4m3
python3 -m sglang.launch_server --model meta-llama/Llama-3.1-8B-Instruct --kv-cache-dtype fp8_e4m3

# then watch startup logs for the unity-scale warning quoted in §0.5.9

• OCP Microscaling Formats (MX) Specification v1.0 — the normative definition of MXFP4/MXFP6/MXFP8, block size 32 and the E8M0 shared scale. Cited directly in vllm/scalar_type.py:L338-L339.
• FP8 Formats for Deep Learning (Micikevicius et al., 2022) — the E4M3/E5M2 proposal, including the rationale for E4M3's non-IEEE extended range.
• LLM.int8(): 8-bit Matrix Multiplication for Transformers at Scale (Dettmers et al., 2022) — the emergent-outlier-feature result that explains why per-tensor INT8 activation quantisation fails past ~6.7B parameters.
• A Study of BFLOAT16 for Deep Learning Training (Kalamkar et al., 2019) — the case for trading mantissa for exponent.
• vLLM PR #4532 — FP8 KV cache with per-tensor scaling factors, the change that introduced the k_scale/ v_scale loading path whose unity-scale fallback is quoted in §0.5.9.
• Next: §4.1 for how these formats become quantisation schemes, and §4.3 for the kernels (Marlin, Machete, CUTLASS FP8) that make narrow formats fast rather than merely small.