KV cache quantization

Someone reads that FP8 KV cache “halves KV memory bandwidth”, flips --kv-cache-dtype fp8 on a Llama-3-8B endpoint, measures TPOT at batch 1, and finds it unchanged. Then they file a bug.

The bug is in the expectation. §0.2 already did this arithmetic: for Llama-3-8B at $s = 2304$ and batch 1, the cached decode step moves $1.606\times10^{10}$ bytes of weights and $3.02\times10^{8}$ bytes of KV. Halving the KV term moves total step bytes from $1.636\times10^{10}$ to $1.621\times10^{10}$ — a 0.9% reduction. On a 3.35 TB/s H100 that is roughly 45 µs off a 4.9 ms step. Whether that small difference is measurable depends on experimental noise and sample size.

So what is it for? First, capacity: the same pool holds twice as many tokens, hence twice the concurrency at a given context length, and concurrency is what converts into throughput. Second, past a crossover point, bandwidth: once batch times context gets large enough, KV traffic overtakes weight traffic and halving it does show up in TPOT. Deriving that crossover is the most useful thing this chapter does; everything after it is about how far down you can push the KV term before the model notices.

Per decode step, a GPU moves two piles of bytes. The weight pile is fixed: every parameter, once, no matter how many sequences are in flight. The KV pile is linear in the total number of resident tokens $N = \sum_i s_i$ — batch size times mean context. Batching amortises the first and does nothing to the second. Quantizing the KV cache scales only the second. So the entire question is: which pile is bigger at your operating point?

Read it as a diagnosis tool: find your operating point on the x-axis; left of the circle for your model and dtype the step is weight-bound and KV quantization has a smaller potential bandwidth benefit, right of it the KV line is the tall pile. Note the structural surprise — the weight line is the same for both models. Llama-3-70B at TP=8 moves 17.7 GB of weights per rank per step, Llama-3-8B on one GPU 16.1 GB. TP divides the weight pile by 8 and, because $h_{kv} = 8$ for both, divides the KV pile by 8 too, so the crossover in tokens is invariant to TP for any $\text{TP} \le h_{kv}$. Reason about it once per model and forget the parallelism.

Symbols, all from FORMULAS: $L$ layers, $h_{kv}$ KV heads, $d_h$ head dimension, $b$ bytes per stored KV element, $b_w$ bytes per weight, $P$ parameters, $N = \sum_i s_i$ resident tokens. Per decode step the machine reads every weight once and every cached K and V element once:

Set them equal and solve for the crossover $N^{*}$, where the two piles are the same size:

Llama-3-8B ($L=32$, $h_{kv}=8$, $d_h=128$, $P = 8.03\times10^{9}$). Bytes per token at $b=2$: $2 \cdot 32 \cdot 8 \cdot 128 \cdot 2 = 131{,}072$ = 128 KiB. Weights at $b_w = 2$: $1.606\times10^{10}$.

Llama-3-70B ($L=80$, $h_{kv}=8$, $d_h=128$, $P = 70.6\times10^{9}$). Bytes per token: $2 \cdot 80 \cdot 8 \cdot 128 \cdot 2 = 327{,}680$ = 320 KiB. Weights: $1.412\times10^{11}$.

The 70B crossover sits 3.5× further out, and that is not an accident: under GQA with a fixed $h_{kv} = 8$, scaling the model up grows $P$ by 8.8× but grows per-token KV by only 2.5× (only $L$ moved). Bigger GQA models are more weight-dominated, not less, so KV quantization is a weaker latency lever for them — and an equally strong capacity lever.

The bottom-left cell is the punchline. A single Llama-3-8B request at 128k context is already past the crossover at batch 1. $131{,}072 \times 131{,}072 = 1.72\times10^{10}$ bytes of KV against $1.606\times10^{10}$ of weights; fp8 KV then takes total step bytes from $3.33\times10^{10}$ to $2.46\times10^{10}$, a 26% reduction that is visible in TPOT. “KV quantization is not a latency lever” is true at 2k and false at 128k — the same arithmetic at different $N$. Quantizing moves the crossover right without removing it: fp8 doubles $N^{*}$, and a 4-bit rotated scheme at 3.8× pushes it to $4.7\times10^{5}$ tokens for the 8B model. Since the weight pile never shrinks, halving KV saves about 0.9% of total bytes at this operating point; eliminating KV entirely would save about 1.8%, before other overheads.

The capacity win, in resident tokens

Take the budget §2.1 already derived for an 80 GB H100 SXM: 79.65 GiB of HBM at vLLM's default gpu_memory_utilization of 0.92 (vllm/config/cache.py:L80) is 73.28 GiB requested, less the weights, less a 6.0 GiB allowance for activations, CUDA graphs, NCCL buffers and the CUDA context — §2.6 does that budget properly. That is a 52.32 GiB pool for the 8B model and 50.85 GiB per rank for 70B at TP=8. Only the bytes-per-token column changes below.

Concurrency amortises the weight read, and that is what turns into throughput. Doubling admitted sequences from 52 to 104 roughly doubles weight-GEMM arithmetic intensity — the move from the memory-bound floor toward the ridge that §0.4 plots. That is how fp8 KV increases tokens per second, and it has nothing to do with the KV bytes it saved.

Each cached entry is normally quantized once rather than repeatedly requantized in place. Nevertheless its error enters later attention outputs and hidden states, influencing new K/V and logits even if sampled token IDs have not yet changed. A changed token can further amplify sequence divergence. Evaluate both continuous numerical drift and task-level generation differences; the absence of repeated storage rounding does not remove numerical feedback.

The harder half is the keys, structurally: $K$ feeds a softmax, where a logit error is exponentiated, while $V$ feeds a convex combination $\sum_i p_i v_i$, where errors partly cancel. Every scheme below treats keys more carefully — vLLM's presets are named turboquant_k8v4 and turboquant_k3v4_nc, key bits before value bits. On top of that sits the finding that made rotated quantization necessary: KIVI (arXiv:2402.02750) and KVQuant (arXiv:2401.18079) both report persistent per-channel outliers in the key cache — a handful of the $d_h$ coordinates consistently far larger than the rest across all tokens — and none in the value cache. KVQuant attributes much of the difficulty to RoPE, which mixes channel pairs by a position-dependent angle and smears the clean pre-RoPE structure; production engines cache post-RoPE keys, because that is what the kernel consumes.

A quantized KV entry is reconstructed as $x \approx s\,(q - z)$ — the affine form §0.5 owns. The design question is how many distinct $(s, z)$ pairs, over what group of values. Every extra scale buys accuracy and costs bytes and a load.

The per-tensor and per-KV-head rows carry no inline metadata: their scales are model weights loaded from the checkpoint, four per layer — $q$, $k$, $v$ and the softmax-output scale — each initialised to an invalid sentinel so an uncalibrated checkpoint is detectable:

    def create_weights(self, layer: torch.nn.Module):
        """
        Create "weight" (aka q_scale, k_scale and v_scale)
        for an attention layer.
        """
        # Initialize the Q and KV cache scales to -1.0, an invalid value.
        # If the q and k/v_scales appear in the checkpoint, it will be
        # overwritten when loading weights.
        layer.q_scale = KVCacheScaleParameter()
        layer.k_scale = KVCacheScaleParameter()
        layer.v_scale = KVCacheScaleParameter()
        # Initialize P = softmax(QK^T) scales
        layer.prob_scale = KVCacheScaleParameter()

That path is strictly per-tensor and says so: kv_cache.py:L124-L127 raises "Only support per-tensor scaling factor for fp8 KV cache" if the loaded scale is not a scalar. Per-head scales arrive through a different loader — compressed-tensors' ATTN_HEAD strategy, which allocates n_scales-wide parameters, shards them TP-aware, and reduces $q$ scales from $h$ to $h_{kv}$ by taking the max within each GQA group (compressed_tensors.py:L1129-L1167). They work only with FlashAttention and require llm-compressor calibration (docs/features/quantization/quantized_kvcache.md).

Per-token scales need no calibration — the writer has the whole vector in registers — but the scale must then live in the cache. vLLM's Triton backend widens the head dimension by exactly one fp32 and hides the scale in the pad:

    @classmethod
    def customize_spec(cls, spec: "AttentionSpec") -> "AttentionSpec":
        """Per-token-head modes pack inline fp32 scales after each head's
        data, so the content is (data + one scale) per K/V side."""
        mode = spec.kv_quant_mode
        if spec.state_content_bytes is not None or not mode.is_per_token_head:
            return spec
        hs_k, hs_v = spec.head_size, spec.head_size_v
        if mode == KVQuantMode.INT4_PER_TOKEN_HEAD:
            hs_k, hs_v = hs_k // 2, hs_v // 2
        scale_bytes = get_dtype_size(torch.float32)
        content = (hs_k + hs_v) * get_dtype_size(spec.dtype) + 2 * scale_bytes
        return replace(spec, state_content_bytes=content)

state_content_bytes flows straight into page_size_bytes = num_heads * storage_block_size * state_content_size_bytes (vllm/v1/kv_cache_interface.py:L244-L260) — that is the whole mechanism by which a KV dtype change changes how many blocks the engine allocates. The scale is recovered by _ensure_scale_caches (triton_attn.py:L440-L500) as a strided fp32 view over the pad bytes: a zero-copy reinterpretation of storage the block pool already owns, so §2.2's allocator never has to know scales exist.

Per-block is the fourth point on the axis, and NVFP4's: one fp8 e4m3 scale per 16 consecutive values along the head dimension, over a per-layer fp32 global scale. SGLang names the block size explicitly (SCALE_BLOCK_SIZE = 16, fp4_kv_cache_quant_method.py:L340); vLLM computes the packed width as head_size // 2 + head_size // 16, 64 data bytes plus 8 scale bytes per side at $d_h = 128$ (vllm/utils/torch_utils.py:L531-L534). Finest granularity of the four, highest metadata share in the table.

vLLM's legal KV dtypes are a single Literal. As of a556f3f it has seventeen members, and reading it is the fastest way to see the whole design space at once:

CacheDType = Literal[
    "auto",
    "float16",
    "bfloat16",
    "fp8",
    "fp8_e4m3",
    "fp8_e5m2",
    "fp8_inc",
    "fp8_ds_mla",
    "turboquant_k8v4",
    "turboquant_4bit_nc",
    "turboquant_k3v4_nc",
    "turboquant_3bit_nc",
    "int4_per_token_head",
    "int8_per_token_head",
    "fp8_per_token_head",
    "nvfp4",
    "nvfp4_4over6",
]

Four families. Plain floats (auto, float16, bfloat16) store the model dtype. The fp8 family adds fp8_inc for Intel Gaudi and fp8_ds_mla for DeepSeek's MLA latent (§7.2). The per-token-head family is the dynamic-scale path just described. The NVFP4 pair is packed 4-bit float with block scales, and its second member is the exotic one: nvfp4_4over6 “uses the NVFP4 layout and selects between max/6 and max/4 scales per 16 values by minimizing squared reconstruction error” (cache.py:L95-L97) — a per-block search over two scale conventions at write time, landed in PR #45187.

Storage is blunt: nearly everything narrower than fp16 is torch.uint8, because a byte array is the only thing every allocator, copy engine and offload path handles uniformly (vllm/utils/torch_utils.py:L33-L54). The semantic classification lives in an enum whose comments are the best one-line summary of each scheme in either repo:

    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
    # Hadamard-rotated Lloyd-Max quant, packed K+V per slot.
    TURBOQUANT_K8V4 = 6
    TURBOQUANT_4BIT_NC = 7
    TURBOQUANT_K3V4_NC = 8
    TURBOQUANT_3BIT_NC = 9

SGLang's surface is deliberately smaller: configure_kv_cache_dtype accepts auto, bf16, fp8_e5m2, fp8_e4m3, mxfp8, nvfp4 and fp4_mx_block16, and raises on anything else (python/sglang/srt/mem_cache/kv_cache_dtype.py:L37-L81) — no integer per-token-head modes and no rotated schemes at this SHA. One thing it does that vLLM does not: auto reads quant_config.kv_cache_quant_algo off the checkpoint and silently enables fp8 if the checkpoint says so (kv_cache_dtype.py:L37-L45). Its storage rule matches vLLM's, one line in the pool constructor:

        if dtype in (torch.float8_e5m2, torch.float8_e4m3fn, torch.float8_e4m3fnuz):
            # NOTE: Store as torch.uint8 because Tensor.index_put is not implemented for torch.float8_e5m2
            self.store_dtype = torch.uint8
        else:
            self.store_dtype = dtype

This detail decides whether fp8 KV is a pure memory trick or also a compute trick. The cache holds a byte; between the load and the softmax that byte must become a number the attention math can use — or the math must be rewritten so it never does. There are four answers in these two repos.

Strategy A — dequantize the tile in registers. The general case. vLLM's unified Triton kernel has a helper whose docstring is a complete taxonomy of the modes it handles:

@triton.jit
def _cast_kv_tile(data, Q, tensor_scale, KV_QUANT_MODE: tl.constexpr):
    """Cast a loaded KV tile to Q's dtype, dequantizing if needed.

    Modes handled inside the core kernel:

    - ``KV_QUANT_MODE == 0`` (NONE) and ``2`` (INT8 per-token-head) and
      ``3`` (FP8 per-token-head): plain cast.  Per-token-head modes apply
      their scales separately on S/P inside the loop.
    - ``KV_QUANT_MODE == 1`` (FP8 per-tensor): dequantize using the
      tensor-wide scale, unless Q is also FP8 and the caller folds the scales
      into the attention score and output accumulator.
    """
    if KV_QUANT_MODE == 1:
        if Q.dtype.is_fp8():
            return data.to(Q.dtype)
        return (data.to(tl.float32) * tl.load(tensor_scale)).to(Q.dtype)
    return data.to(Q.dtype)

Three instructions on data already in registers. HBM traffic is halved; the matmul still runs at bf16 rates. This is the mode that produces the 0.9% at batch 1.

Strategy B — never dequantize; fold the scales into the score. With an fp8 query, $QK^\top$ runs on fp8 tensor cores and both scales collapse into one scalar on the pre-softmax score. Same kernel, twelve lines up:

    if USE_FP8_Q_DESCALE:
        score_scale = scale * tl.load(q_scale) * tl.load(k_scale)

FlashAttention takes the same route through a different door: it passes k_descale = layer._k_scale.expand(descale_shape), shape (num_seqs, num_kv_heads), into flash_attn_varlen_func (flash_attn.py:L1016-L1023), having first viewed the uint8 cache as a real fp8 tensor:

        if is_quantized_kv_cache(self.kv_cache_dtype):
            # queries are quantized in the attention layer
            key_cache = key_cache.view(current_platform.fp8_dtype())
            value_cache = value_cache.view(current_platform.fp8_dtype())

This is the only path where fp8 KV buys compute as well as bytes, and vLLM's docs say so outright: “When using the Flash Attention 3 backend with FP8 KV cache, attention operations are also performed in the quantized (FP8) domain” (docs/features/quantization/quantized_kvcache.md). SGLang does the same at flashattention_backend.py:L1288-L1293.

Strategy C — apply the scales to the score matrix, never to K or V. For per-token-head modes the scale varies along the KV-tile axis, so folding it into the score matrix costs the same as multiplying it into the tile and fuses with the softmax scale you were applying anyway:

        if USE_PER_TOKEN_HEAD_SCALES:
            # Per-token-head quant: fuse softmax_scale with per-head k_scale
            # to avoid a separate BLOCK_M × TILE_SIZE multiply on S.
            S += tl.dot(Q, K) * (score_scale * k_token_head_scales[None, :])
        else:
            S += score_scale * tl.dot(Q, K)

Strategy D — dequantize the whole buffer into a scratch workspace first. SGLang's FP4 path formalises this as an access kind, declared per recipe and per phase:

KV_CACHE_ATTENTION_ACCESS_REGISTRY: dict[str, tuple[KVCacheAttentionAccess, ...]] = {
    UnquantizedKVCacheMethod.name: (
        _plain(_PREFILL, _ANY_BACKEND),
        _plain(_DECODE, _ANY_BACKEND),
    ),
    NVFP4KVCacheMethod.name: (
        _dq_workspace(_PREFILL, _NVFP4_PREFILL_BACKENDS, _NVFP4_SCALE, _FP8_E4M3),
        _native_fp4(_DECODE, _NVFP4_DECODE_BACKENDS, _NVFP4_SCALE, _TORCH_FP4),
    ),
    FP4MXBlock16KVCacheMethod.name: (
        _plain(_PREFILL, _FP4_MX_PREFILL_BACKENDS, _FP4_MX_SCALE, _BF16),
        _plain(_DECODE, _FP4_MX_MHA_BACKENDS, _FP4_MX_SCALE, _BF16),
    ),
}

NVFP4 uses a workspace for prefill and native packed FP4 for decode, which is the right call: prefill reads each entry once, so the extra pass is bounded overhead, while decode reads the whole cache every step, so a workspace pass would give back the traffic you just spent 4 bits to save. It is a registry rather than an if in each backend because torch.float4_e2m1fn_x2 describes packed storage but “does not say whether the recipe is NVFP4 or fp4_mx_block16, nor how scales are interpreted” (fp4_kv_cache_quant_method.py:L15-L34).

Which backend accepts which dtype

Backends advertise a supported_kv_cache_dtypes class variable; the selector turns a mismatch into the reason string "kv_cache_dtype not supported" (vllm/v1/attention/backend.py:L379-L380). The gaps are the interesting part.

Two constraints surprise people. First, FlashAttention refuses fp8_e5m2 unconditionally and accepts fp8_e4m3 only on a narrow set of architecture-plus-version combinations:

def flash_attn_supports_kv_cache_dtype(
    kv_cache_dtype: str = "fp8_e4m3",
    *,
    requires_alibi: bool = False,
    head_size: int | None = None,
    head_size_v: int | None = None,
    has_sinks: bool = False,
) -> bool:
    if kv_cache_dtype == "fp8_e5m2":
        return False
    if current_platform.is_xpu():
        return True
    fa_version = get_flash_attn_version(
        requires_alibi=requires_alibi,
        head_size=head_size,
        head_size_v=head_size_v,
        has_sinks=has_sinks,
    )
    return (fa_version == 3 and current_platform.is_device_capability_family(90)) or (
        fa_version == 4 and current_platform.is_device_capability_family(100)
    )

Which is why --kv-cache-dtype fp8 on an A100 silently routes you off FlashAttention onto FlashInfer or Triton rather than failing. Second, TurboQuant is not an option on an existing backend, it is a backend, because it changes the shape of the cache tensor: it drops the leading K/V dimension and stores one interleaved slot per head per position (turboquant_attn.py:L180-L210), safe only because it never shares cache tensors with any other backend.

Below about 5 bits, per-token scaling stops being enough. If one coordinate of a 128-dimensional key is 20× the others, a uniform quantizer sized to that coordinate assigns every other coordinate to the bottom one or two codes: you spend 4 bits per coordinate and learn almost nothing about 127 of them.

The fix is a change of basis. Let $H$ be orthonormal and $y = Hx$. Three facts, in increasing order of usefulness.

1. Orthogonality preserves everything attention cares about. $\|Hx\| = \|x\|$, and $(Hq)^\top(Hk) = q^\top H^\top H k = q^\top k$, so rotating both query and key leaves every logit numerically identical. Values are even nicer: the output is a convex combination $o = \sum_i p_i v_i$, so storing $Hv_i$ gives $\sum_i p_i (H v_i) = H o$ and one inverse rotation recovers $o$ exactly. Rotation does not change the function being computed.

2. The Hadamard matrix spreads energy maximally and costs almost nothing. Its entries are all $\pm 1/\sqrt{d}$, so $y_j = \tfrac{1}{\sqrt d}\sum_i \pm x_i$ is a signed sum of every input coordinate: a spike is smeared over all $d$ of them, each getting $1/\sqrt d$ of it — an 11.3× reduction at $d_h = 128$. The transform is a butterfly of adds and subtracts, $O(d \log d)$, no multiplies, or as vLLM does it a single GEMM against a cached matrix.

def _build_hadamard(d: int, device_str: str) -> torch.Tensor:
    """Orthonormal Hadamard matrix (Sylvester construction), cached per (d, device).

    Precomputed D×D matrix enables matmul-based WHT — single cuBLAS GEMM
    instead of log2(D) butterfly kernel launches. 64KB for D=128.
    """
    # Normalize device string so "cuda" and "cuda:0" hit the same cache entry.
    return _build_hadamard_cached(d, str(torch.device(device_str)))


@functools.cache
def _build_hadamard_cached(d: int, device_str: str) -> torch.Tensor:
    H = torch.tensor([[1.0]])
    while H.shape[0] < d:
        H = torch.cat([torch.cat([H, H], 1), torch.cat([H, -H], 1)], 0)
    return (H / math.sqrt(d)).to(torch.device(device_str))

3. After rotation the distribution is not just narrower, it is known. This is the part that turns a heuristic into a design. For a unit-norm $x$ and a random orthogonal $H$, each coordinate of $y$ has mean 0 and variance exactly $1/d$, and for $d \gtrsim 64$ the marginal is close to Gaussian. vLLM states this as the premise of its centroid solver:

"""Lloyd-Max optimal scalar quantizer for TurboQuant.

After rotating a d-dimensional unit vector by a random orthogonal matrix,
each coordinate approximately follows N(0, 1/d) for d >= 64.
We solve the Lloyd-Max conditions to find optimal centroids.

Based on: turboquant-pytorch/lloyd_max.py (Zandieh et al.)
"""

For a random sign/Hadamard or Haar-rotation model, concentration results motivate a statistical range near a multiple of $1/\sqrt d$ after normalization. A fixed Hadamard matrix does not guarantee Gaussian coordinates or a hard $3.5/\sqrt d$ maximum. Indeed choose x as one row of H: Hx is a coordinate vector, with maximum magnitude 1. Any clipped range has a tail/error policy, and a stored norm restores scale rather than proving a deterministic bound.

Rotation does not make the vector smaller — the norm is identical by construction. It makes it flat, and flatness is what a scalar quantizer is good at.

Why “random”, and why TurboQuant dropped the randomness

A fixed Hadamard has one weakness: an input can align with one of its rows and come out as spiky as it went in. The fix is the randomized Hadamard transform $\text{RHT} = H D x$, with $D$ a diagonal of arbitrary $\pm 1$ signs. vLLM's INT4 path does this:

# Randomized Hadamard Transform (used by INT4)
# Deterministic ±1 signs for Randomized Hadamard Transform.
# RHT = H × D × x  (sign flip + Hadamard).  Breaks residual structure
# in KV vectors, improving quantization quality.
_RHT_SIGNS_CACHE: dict[tuple[int, int, str], torch.Tensor] = {}


def _get_rht_signs(d: int, round_idx: int, device: torch.device) -> torch.Tensor:
    """Return a cached deterministic ±1 sign vector of length *d*."""
    key = (d, round_idx, str(device))
    if key not in _RHT_SIGNS_CACHE:
        gen = torch.Generator(device="cpu")
        gen.manual_seed(0x9E3779B9 + round_idx * 0x517CC1B7)

The seed is a constant, so the signs are identical on every rank and every restart — “random” means decorrelated from the data, not varying. TurboQuant went the other way and removed its sign flips in PR #40194, with a stated argument:

    def _ensure_on_device(self, layer, device):
        """One-time derivation of TQ buffers (rotation matrix, midpoints).

        The Hadamard rotation is shared across all layers: random sign
        flips do not improve Lloyd-Max quantization quality because the
        quantizer is symmetric around zero (sign-flipping a coordinate
        maps it to the mirror centroid with identical distortion).
        """

Two payoffs from dropping $D$: the pure Sylvester Hadamard is symmetric, $H = H^\top$, so the inverse rotation is the same matmul (turboquant_attn.py:L489-L491), and the transform can be fused as a butterfly inside the decode kernel. Whether the accuracy argument fully holds is an empirical question the PR settled for these presets, not a theorem about rotated quantization in general.

What Lloyd-Max adds on top

A random-rotation model motivates an approximately Gaussian marginal; a fixed Hadamard transform does not guarantee one, and a uniform quantizer is not the best code for a Gaussian — you want levels packed near zero where the density is. The optimal $2^b$-level scalar quantizer satisfies two coupled conditions that Lloyd's algorithm alternates to a fixed point: each boundary is the midpoint of its neighbouring centroids, each centroid the conditional mean of the source over its own cell. vLLM solves them offline, once per $(d, b)$:

    for _ in range(max_iter):
        boundaries = [
            (centroids[i] + centroids[i + 1]) / 2.0 for i in range(n_levels - 1)
        ]
        edges = [lo * 3] + boundaries + [hi * 3]
        new_centroids = []
        for i in range(n_levels):
            a, b = edges[i], edges[i + 1]
            num = _trapz(lambda x: x * pdf(x), a, b)
            den = _trapz(pdf, a, b)
            new_centroids.append(num / den if den > 1e-15 else centroids[i])

        if max(abs(new_centroids[i] - centroids[i]) for i in range(n_levels)) < tol:
            break
        centroids = new_centroids
# ...
@lru_cache(maxsize=32)
def get_centroids(d: int, bits: int) -> torch.Tensor:
    """Get precomputed Lloyd-Max centroids (cached)."""
    centroids, _ = solve_lloyd_max(d, bits)
    return centroids

This is one-dimensional k-means against a known density rather than a sample, and because the rotation fixes the density the result is a compile-time constant: the whole codebook for $d = 128$, $b = 4$ is sixteen fp32 numbers, shared by every layer and every request. Encoding is a binary search over the 15 midpoints, $b$ iterations, entirely in registers (vllm/v1/attention/ops/triton_turboquant_store.py:L278-L292). Rotation is the large win; Lloyd-Max harvests what is left. Skip the rotation and 4-bit keys collapse; use uniform levels instead of Lloyd-Max ones and you lose a fraction of a bit.

TurboQuant

Four presets, frozen parameter sets rather than knobs:

# Named TQ presets: each maps to frozen config parameters.
# key_quant_bits: 8 = FP8 keys, 3-4 = MSE (Lloyd-Max) quantized keys.
# value_quant_bits: 3-4 = uniform quantized values.
TQ_PRESETS: dict[str, dict] = {
    "turboquant_k8v4": {
        "key_quant_bits": 8,
        "value_quant_bits": 4,
        "norm_correction": False,
    },
    "turboquant_4bit_nc": {
        "key_quant_bits": 4,
        "value_quant_bits": 4,
        "norm_correction": True,
    },
    "turboquant_k3v4_nc": {
        "key_quant_bits": 3,
        "value_quant_bits": 4,
        "norm_correction": True,
    },
    "turboquant_3bit_nc": {
        "key_quant_bits": 3,
        "value_quant_bits": 3,
        "norm_correction": True,
    },
}

Note what is not symmetric. Keys get rotation, Lloyd-Max centroids and a stored norm; values get plain asymmetric uniform quantization with a per-(token, head) fp16 scale and zero and no rotation, value_packed_size = ceil(head_dim * bits / 8) + 4 (config.py:L149-L156). That is the K/V asymmetry of §5 turned into a memory layout. In k8v4 the keys are not even narrow: plain fp8, rotation machinery bypassed (config.py:L94-L97).

norm_correction is the third ingredient, and the one with no analogue in weight quantization. Each Lloyd-Max centroid is a conditional mean, and a mean is a contraction, so a rotated unit vector reconstructed coordinate-by-coordinate has norm systematically below 1. The fix is to renormalise the reconstructed direction before multiplying by the stored $\|k\|$; the docstring prices it at “improving PPL by ~0.8% at 4-bit” (config.py:L83-L85).

The fourth ingredient is not an algorithm. Aggressive presets skip the first and last two attention layers, leaving them at native dtype:

        For dense models, skips first N and last N attention layers.
        Empirically required for aggressive presets (k3v4_nc, 3bit_nc)
        — without it GSM8K drops ~30 points on Qwen3-4B.
        """

Thirty GSM8K points from four layers out of thirty-six. Sensitivity is not uniform across depth, and the mechanism is a first-class flag: --kv-cache-dtype-skip-layers takes layer indices or type names such as sliding_window (vllm/config/cache.py:L134-L136).

INT4 per-token-head

The other rotated scheme is smaller in scope and cleverer in the kernel. Its docstring is a complete specification:

"""Sub-byte packed (INT4) per-token-head KV cache mode.

INT4 packs two 4-bit values per cache byte, pre-rotates with a single RHT,
and hides a 4-bit zero-point in the scale's low mantissa bits — too
different from the core kernel to share it. Owns the whole mode: nibble
pack/unpack, the reshape (write) kernel, the split-dot attention (read)
kernel, the RHT transform, and the public ``reshape_and_cache_int4`` /
``unified_attention_int4`` entry points.
"""

Three details worth pulling out. The zero-point is steganographed into the scale. Rather than store a separate zero-point byte, the write kernel clears the bottom 4 mantissa bits of the fp32 scale and puts the 4-bit zero-point there: k_scale_packed = ((k_scale_bits & -16) | (k_zp_int & 0xF)) (int4_per_token_head.py:L157-L161), and the read kernel splits them apart again (L549-L553). Cost: 4 bits of scale mantissa, leaving 19 — more precision than an fp16 scale would have. Benefit: the scale slot in the page stays exactly one fp32 wide, so INT4 reuses the same inline-scale page layout as INT8 and FP8.

The dequantisation is algebraic, not materialised. With $\hat k = s(q_k - z_k)$, $\langle q, \hat k\rangle = s(\langle q, q_k\rangle - z_k \sum_j q_j)$. The kernel computes the raw integer dot product and then subtracts the correction term, never building a dequantized key:

        # Score: split-dot across the 2 INT4 streams; fused
        # softmax_scale * per-(token, head) k_scale in one mul.  INT4
        # subtracts the ``zp * sum(Q)`` correction term.
        S = tl.zeros(shape=(BLOCK_M, TILE_SIZE), dtype=tl.float32)
        raw_dot = tl.dot(Q_s0, K_s0) + tl.dot(Q_s1, K_s1)
        S += (raw_dot - Q_sum[:, None] * k_zp[None, :]) * (
            scale * k_token_head_scales[None, :]
        )

The same trick on the value side subtracts $\sum_i p_i z_v$ from the accumulator once per tile (L604-L608). Two “streams” appear because the low and high nibbles are unpacked and dotted separately — a split-dot, which is why this could not fold into the shared unified-attention kernel.

The rotation is applied to the query and undone on the output — fact 1 above, cashed in:

    """Paged attention over the INT4 packed cache, writing into *out*.

    The forward RHT has norm ``sqrt(head_size)``, so ``softmax_scale`` is
    divided by ``head_size`` and the inverse RHT divides the output by
    ``head_size`` as well.
    """
    q_orig_dtype = q.dtype
    q = single_rht(q.float()).to(q_orig_dtype)
    head_size = q.shape[2]
    softmax_scale = softmax_scale / head_size

The transform used here is the unnormalised Walsh-Hadamard (L1073-L1074), which scales by $\sqrt{d}$ rather than dividing by it. Two of those appear in $\langle Hq, Hk\rangle$, hence the $/\,d$ on softmax_scale, and one appears in the output, hence the second $/\,d$ at the end. Getting one of those two divisions wrong produces attention that looks plausible and is silently wrong — a good thing to know before you go reading the kernel.

Llama-3-8B, layer 7, decode step, --kv-cache-dtype fp8_e4m3, FlashAttention on an H100. Follow one key vector, $d_h = 128$, KV head 3.

1. Startup. CacheConfig._validate_cache_dtype (vllm/config/cache.py:L293-L311) logs “Using fp8_e4m3 data type to store kv cache… it may cause accuracy drop without a proper scaling factor”.
2. Weight load. BaseKVCacheMethod.process_weights_after_loading (vllm/model_executor/layers/quantization/kv_cache.py:L74-L152) finds layer.k_scale < 0 because the checkpoint carries no calibrated scales, falls back to 1.0, and emits the warning that matters most in practice: “Using KV cache scaling factor 1.0 for fp8_e4m3…”. It copies the value into layer._k_scale, the fp32 device tensor the kernels read, plus _k_scale_float and _k_scale_cpu mirrors.
3. Page sizing. AttentionSpec.page_size_bytes (vllm/v1/kv_cache_interface.py:L254-L260) evaluates num_heads * storage_block_size * state_content_size_bytes. With kv_quant_mode = FP8_PER_TENSOR the spec dtype is one byte, so state_content_size_bytes = (128 + 128) * 1 = 256 against 512 for bf16. The profiler divides the same pool by half the page size and allocates twice the blocks. That is the entire capacity win, and it lands before a single token is served.
4. Write. After QKV projection and RoPE, the attention layer calls reshape_and_cache_flash with layer._k_scale. The kernel divides by the scale, clamps to the e4m3 range, casts, and scatters into the slot given by slot_mapping[token] — a slot chosen by §2.2's block allocator, which is unaware the bytes are quantized.
5. Read. FlashAttentionImpl.forward reinterprets the uint8 cache as fp8 (flash_attn.py:L1003-L1006), expands the scalar scale to (num_seqs, num_kv_heads), and passes q_descale, k_descale, v_descale into flash_attn_varlen_func (flash_attn.py:L1016-L1023). The query was already quantized to fp8 in the attention layer, so no dequantisation happens anywhere: the $QK^\top$ MMA consumes fp8 operands and the descales multiply the fp32 accumulator.

Those same 128 bytes are then re-read at every subsequent decode step for the life of the sequence — the read amplification behind the crossover in §3, and the reason a rounding applied once in step 4 is seen thousands of times.

Swap FlashAttention for Triton and only step 5 changes: _cast_kv_tile multiplies the tile by the scale in registers and the MMA runs in bf16. Swap in int8_per_token_head and steps 2 and 4 change instead — no checkpoint scale to load (kv_cache.py:L82-L91 sets both to 1.0 and deletes the parameters), and the write kernel computes k_scale = max(abs(k_head)) / 127 per token per head and stores it inline (triton_reshape_and_cache_flash.py:L207-L214).

I have no GPU, so there is no measured accuracy table here, and you should distrust any book that hands you one without a configuration. What exists is published work and the repositories' CI gates.

Published. KIVI (arXiv:2402.02750) establishes the asymmetric treatment — per-channel keys, per-token values. KVQuant (arXiv:2401.18079) analyses the RoPE interaction and the pre- versus post-RoPE choice. For the rotation line, vLLM's module docstring is unusually careful about attribution and hands you the reading list: DRIVE (Vargaftik et al., NeurIPS 2021), EDEN (Vargaftik et al., ICML 2022), HIGGS (arXiv:2411.17525) and “Cache Me If You Must” (arXiv:2501.19392), all of which predate the TurboQuant paper (Zandieh et al., ICLR 2026) the feature is named for (turboquant/__init__.py:L3-L22). The repository gives no arXiv identifier for TurboQuant itself.

Repository CI gates. Not measurements — the floor below which maintainers consider the feature broken — but the closest thing to a published number that ships with the code:

The rightmost column is quoted from turboquant/config.py:L71-L75. Three of the four ratios reproduce exactly from slot_size_aligned arithmetic; k3v4_nc does not.

What to measure yourself: measure perplexity or token-level log loss together with representative exact-match/generative tasks at deployment context lengths. Include rare long-context cases, multiple seeds where sampling is used, and the same engine/backend baseline. Neither average likelihood nor one accuracy score alone establishes acceptability; no unmeasured thirty-point drop is implied.

Watch the block count change — the capacity lever made visible. Run the same model twice and diff the startup line reporting GPU KV cache size:

vllm serve meta-llama/Meta-Llama-3-8B-Instruct --max-model-len 8192 2>&1 | grep -i "KV cache"
vllm serve meta-llama/Meta-Llama-3-8B-Instruct --max-model-len 8192 --kv-cache-dtype fp8 2>&1 | grep -i "KV cache"

With a fixed byte pool and identical layout overhead, halving payload bytes gives approximately twice the token capacity. Block floors, reserved null blocks, scale metadata, alignment, backend selection, and newly profiled overhead can prevent an exact 2.00 ratio. Compare allocated usable blocks as well as nominal payload bytes.

# Left of the crossover: N = 2304 tokens. Expect TPOT unchanged to within noise.
vllm bench serve --model meta-llama/Meta-Llama-3-8B-Instruct --dataset-name random \
  --random-input-len 2048 --random-output-len 256 --max-concurrency 1 --num-prompts 32

# Right of it: N = 32 x 8192 = 262144 tokens, above the 1.23e5 bf16 crossover.
vllm bench serve --model meta-llama/Meta-Llama-3-8B-Instruct --dataset-name random \
  --random-input-len 8000 --random-output-len 192 --max-concurrency 32 --num-prompts 256

For accuracy, vLLM ships the exact eval the TurboQuant presets are gated on:

pytest -s -v tests/evals/gsm8k/test_gsm8k_correctness.py \
    --config-list-file=configs/models-small.txt

Point --config-list-file at a file listing Qwen3-4B-TQ-t4nc.yaml and its siblings to reproduce the thresholds in §10. Lab 06-quantization-tradeoff is the fuller version of this exercise.

Let W be weight bytes and K be KV bytes per step. Halving KV saves $K/[2(W+K)]$ of total bytes. Write $r=K/W$: the saving is $r/[2(1+r)]$, equal to 0.50% at r=0.01, 16.7% at r=0.5, 25% at parity, and approaching 50% only when KV dominates. There is no discontinuity at the crossover. Actual time savings additionally depend on conversion cost, other traffic, and whether the step is bandwidth-limited. The 128k example uses Llama-3 tensor shapes hypothetically; use a natively long-context checkpoint such as Llama-3.1 for an actual experiment.

• KIVI: A Tuning-Free Asymmetric 2bit Quantization for KV Cache — arXiv:2402.02750. The per-channel-key / per-token-value asymmetry, with the outlier analysis that motivates it.
• KVQuant: Towards 10 Million Context Length LLM Inference with KV Cache Quantization — arXiv:2401.18079. Pre- versus post-RoPE key quantization, and why RoPE makes the key cache hard.
• Pushing the Limits of LLM Quantization via the Linearity Theorem (HIGGS) — arXiv:2411.17525, named by vLLM's TurboQuant module as mathematically equivalent to its scalar case.
• Cache Me If You Must: Adaptive Key-Value Quantization — arXiv:2501.19392, cited in turboquant/__init__.py as the first application of rotation-plus-scalar-quantization to the KV cache.
• vLLM PR #38479 — TurboQuant: KV cache compression with 4x capacity, the backend that owns its own cache shape.
• vLLM PR #40194 — TurboQuant: remove redundant random signs, add prior art attribution. Read this one for the argument about why the sign flips were dropped.
• vLLM PRs #40835 and #38378 — Triton INT4 per-token-head and the INT8/FP8 per-token-head foundation it builds on; #45187 for nvfp4_4over6.
• SGLang PRs #21954 and #21601 — the NVFP4 quantization-strategy abstraction that introduced the access-kind registry, and the FP4 KV design document.
• vLLM docs — docs/features/quantization/quantized_kvcache.md, on which scale granularities need llm-compressor and which backends accept them.