KV-Cache Quantization

The one-sentence version

Quantizing the KV cache trades numerical precision for context capacity and decode memory bandwidth: FP8 halves storage with negligible metadata and runs natively on Hopper Tensor Cores, while 4-bit formats require dequantization or specialized scaling that can reduce throughput unless memory capacity is the binding constraint.

---

1. Opening Workload Decision: The Three Pillars

Before deriving formulas or surveying kernel implementations, consider an original, concrete serving deployment.

The Production Scenario

An infrastructure team serves Llama-3.1-8B (L = 32 layers, H_kv = 8 KV heads, D = 128 head dimension) on an enterprise datacenter GPU.

• Workload Target: B = 16 concurrent streams, each operating at an extended context length of S = 8,192 tokens.
• Dedicated KV Memory Pool: An explicitly hypothetical 20.0 GiB (21,474,836,480 bytes) pool of VRAM dedicated purely to KV cache storage after model weights and runtime workspace buffers are allocated.

The team must decide whether to deploy standard BF16, FP8 (E4M3FN), or experimental 4-bit KV caching. The decision rests on three visible pillars:

┌─────────────────────────────────────────────────────────────────────────────┐
│                       THE THREE QUANTIZATION PILLARS                        │
├───────────────────────┬─────────────────────────┬───────────────────────────┤
│   1. MEMORY SAVED     │    2. QUALITY CHECK     │  3. RUNTIME & DEQUANT     │
│  Exact GiB per stream │  Falsifiable evaluation │   Direct Tensor Core vs   │
│  and static capacity  │  and hardware stability │   ALU unpack overhead     │
└───────────────────────┴─────────────────────────┴───────────────────────────┘

Pillar 1: Memory Saved and Static Slot Capacity

In standard 16-bit precision (BF16, b = 2.0 bytes/element), the KV cache for a single 8,192-token session consumes:

Memory per session = 2 × 32 × 8 × 128 × 8192 × 2.0 B = 1,073,741,824 B = **1.0 GiB**

Within the 20.0 GiB memory pool:

• BF16 (2.0 B/elem): ⌊20.0 GiB / 1.0 GiB⌋ = 20 raw slots. The 16 concurrent streams consume 16.0 GiB (80% of the pool), leaving only 4 slots of headroom.
• FP8 E4M3FN (1.0 B/elem): Storage drops to 0.5 GiB/session (8.0 GiB total for 16 streams). The pool supports ⌊20.0 GiB / 0.5 GiB⌋ = 40 raw slots, doubling static slot capacity.
• 4-bit NVFP4 Storage (0.5625 B/elem with block-16 scale): Storage drops to 0.28125 GiB/session (4.5 GiB total for 16 streams), supporting ⌊20.0 GiB / 0.28125 GiB⌋ = 71 raw slots (a 3.56× capacity expansion over BF16).

Pillar 2: Quality Verification and Hardware Accumulator Bounds

Doubling slot capacity is worthless if the model’s output degenerates:

• Short-to-Medium Context (S ≤ 32k): FP8 E4M3FN preserves retrieval and reasoning quality within statistical noise on standard benchmarks when evaluated with paired bootstrap confidence intervals.
• Ultra-Long Context (S > 100k) on Hopper SM90: On NVIDIA H100 GPUs using FlashAttention-3, hardware FP32 accumulators suffer numerical truncation over long contraction loops. On a 128k Needle-in-a-Haystack (NIAH) retrieval test, unpatched FP8 retrieval collapsed from 91% to 13%. A software two-level accumulation algorithm restored accuracy to 89%, but introduces an ≈1.6× prefill latency penalty for models with head dimension D = 256 (though negligible for D = 128).
• 4-bit Precision Degradation: Experimental algorithmic 4-bit schemes (such as vLLM’s TurboQuant) preserve most accuracy on standard benchmarks, but sub-4-bit and 3-bit variants experience substantial accuracy drops (up to ~20 points on challenging benchmarks like AIME25 and LiveCodeBench).

Pillar 3: Runtime Dequantisation and Throughput Costs

Halving memory traffic does not automatically double serving throughput:

• Path A: Native FP8 Tensor Core Execution (Hopper SM90 + FlashAttention-3): Keys and Values reside in FP8 in HBM, load directly via the Tensor Memory Accelerator (TMA) into SRAM, and execute natively on Tensor Cores using wgmma instructions. No dequantization to BF16 occurs for cached tokens. In measured vLLM benchmarks on H100, the single-stream decode Inter-Token Latency (ITL) slope drops by 45.8%, while multi-stream throughput increases by 14.9% (because weight reads and prefill bursts still share the memory bus and compute units).
• Path B: Dequantize-on-the-Fly (Ampere SM80 or 4-bit TurboQuant in vLLM): Ampere lacks native FP8 Tensor Cores, and current 4-bit engines dequantize cached blocks back to BF16 in SRAM/registers before attention matrix multiplication. While this relieves HBM read bandwidth, it burns ALU cycles. In vLLM serving tests, TurboQuant 4-bit modes reduce throughput to 66–80% of BF16 and add 10–68% per-token latency, trading raw generation speed for memory capacity under tight VRAM constraints.

---

2. Verified Memory Calculus with Metadata Overhead

2.1 The General KV Cache Formula

For an autoregressive transformer generating across B concurrent sequences of context length S, the raw memory footprint in bytes is:

KV_bytes = 2 × L × H_kv × D × S × B × b

Where:

• 2: Accounts for the two distinct cache tensors: Key (K) and Value (V).
• L: Number of transformer layers (32 for Llama-3.1-8B).
• H_kv: Number of KV attention heads (8 for Llama-3.1-8B with Grouped-Query Attention; 32 for standard Multi-Head Attention).
• D: Head dimension (128).
• S: Sequence length in tokens.
• B: Batch size (number of active concurrent sequences).
• b: Effective bytes per element, including raw data payload and scaling metadata.

Total Elements = 2 (K & V) × 32 (Layers) × 8 (KV Heads) × 128 (Head Dim) × 8192 (Tokens) × 16 (Batch)
               = 8,589,934,592 elements = 2³³ elements

2.2 Metadata Overhead by Precision Format

Quantization formats require scaling factors to map integer or floating-point codes back to numerical representations. Depending on scale granularity, metadata overhead ranges from mathematically negligible to significant:

KV Cache Precision Formats: Payload, Metadata, and Compression

Format	Payload Bits	Scaling Scheme	Metadata Overhead	Effective Bytes/Elem (b)	Effective Compression
BF16 / FP16	16 bits	None	0 bytes	2.0 B	1.00× (Baseline)
FP8 (E4M3FN)	8 bits	Per-tensor or per-head	1 FP32 scale (4 B) per tensor/head	~1.00000003 B	2.00×
INT8 (symmetric)	8 bits	Per-token / per-head	1 FP16/FP32 scale per token vector	1.0156 B – 1.0312 B	~1.94× – 1.97×
4-bit NVFP4	4 bits	Block-16 microscaling	8-bit E4M3 scale per 16 elements	0.5625 B (+12.5%)	3.56×
4-bit MXFP4	4 bits	Block-32 microscaling	8-bit E8M0 scale per 32 elements	0.53125 B (+6.25%)	3.76×
vLLM TurboQuant 4-bit	4 bits	Algorithmic polar / norm correction	Vector scale + norm correction	Algorithmic layout	~2.4× – 3.4× capacity

Metadata Derivations

1. FP8 Per-Tensor: One FP32 scale factor (4 bytes) per layer per cache tensor. Across 32 layers, total metadata is 2 × 32 × 4 B = 256 bytes for the entire model cache. Across 8,589,934,592 elements, metadata adds 256 / 2^33 ≈ 2.98 × 10^-8 B/elem, giving an effective b ≈ 1.00000003 B (negligible).
2. 4-bit NVFP4 Block-16: One 8-bit scale byte per 16 elements (8 bytes of 4-bit nibbles). Metadata overhead is 1 / 16 = 0.0625 B/elem. Adding the 0.5 B payload yields b = 0.5625 B/elem (12.5% overhead). A theoretical 4× compression reduces to 2.0 / 0.5625 = 3.56×.
3. 4-bit MXFP4 Block-32: One 8-bit scale byte per 32 elements. Metadata overhead is 1 / 32 = 0.03125 B/elem, yielding b = 0.53125 B/elem (6.25% overhead) and 2.0 / 0.53125 = 3.76× compression.
4. vLLM TurboQuant: An algorithmic software quantization scheme using polar coordinates and norm correction (turboquant_4bit_nc and turboquant_k8v4). It achieves 3.4× (4-bit keys and values) and 2.4× (8-bit keys and 4-bit values) memory capacity in vLLM benchmarks, distinct from hardware microscaling storage formats.

2.3 Worked Footprint Comparison: Llama-3.1-8B (B=16, S=8192)

Total Elements = 2^33 = 8,589,934,592 elements

Format	Effective b	Total Bytes (2^33 × b)	Exact Binary (GiB)	Decimal (GB)	Per-Stream (S=8192)
BF16	2.0 B	17,179,869,184 B	16.0 GiB	17.18 GB	1.0 GiB
FP8 (E4M3FN)	~1.0 B	8,589,934,592 B	8.0 GiB	8.59 GB	0.5 GiB
INT8 (per-token)	~1.0156 B	8,724,152,320 B	8.125 GiB	8.72 GB	0.508 GiB
NVFP4 (block-16)	0.5625 B	4,831,838,208 B	4.5 GiB	4.83 GB	0.281 GiB
MXFP4 (block-32)	0.53125 B	4,563,402,752 B	4.25 GiB	4.56 GB	0.266 GiB

Avoid decimal versus binary unit conflation

In memory capacity planning, 16.0 GiB = 17.18 GB (10^9 bytes). Confusing GiB with decimal GB will lead you to undersize your VRAM allocation by 7.4%. Similarly, an 8.0 GiB FP8 cache is 8.59 GB; it is not an 8.59 GiB cache.

2.4 The GQA Diminishing Returns Principle

Grouped-Query Attention (GQA) architecturally compresses the KV cache by sharing key/value heads across query head groups (H_q = 32, H_kv = 8, a 4× reduction):

Format	GQA (H_kv = 8) Total Cache	MHA (H_kv = 32) Equivalent	Absolute Memory Saved by GQA
BF16	16.0 GiB	64.0 GiB	48.0 GiB
FP8	8.0 GiB	32.0 GiB	24.0 GiB
NVFP4	4.5 GiB	18.0 GiB	13.5 GiB

The Architectural Takeaway: On legacy Multi-Head Attention models, quantizing from BF16 to FP8 freed an immense 32.0 GiB of VRAM. On modern GQA models, the exact same percentage reduction saves only 8.0 GiB. While percentage memory reduction remains identical (50%), the absolute memory pressure relieved by sub-4-bit quantization is smaller on GQA models, raising the bar before complex 4-bit dequantization overhead is justified.

---

3. Key vs. Value Quantization Asymmetry

Keys and Values serve fundamentally different roles in attention, resulting in distinct numerical distributions under low-precision representation.

Key Cache Tensor [S, H_kv, D]         Value Cache Tensor [S, H_kv, D]
Tokens (S) ------------------------>  Tokens (S) ------------------------>
┌──┬──┬──┬──┬──┬──┐                   ┌──┬──┬──┬──┬──┬──┐
│  │██│  │  │██│  │                   │  │  │  │  │  │  │
│  │██│  │  │██│  │                   │█████████████████│ → Delimiter/Outlier Token
│  │██│  │  │██│  │                   │  │  │  │  │  │  │
│  │██│  │  │██│  │                   │  │  │  │  │  │  │
└──┴──┴──┴──┴──┴──┘                   └──┴──┴──┴──┴──┴──┘
     ▲     ▲
     Persistent Outlier Channels      Quantization: Per-Token Granularity
     Quantization: Per-Channel        (Isolates outlier tokens to single rows)

3.1 Empirical Distribution Differences

Research in sub-4-bit integer quantization (such as KIVI and KVQuant) demonstrates that Keys and Values exhibit different outlier geometry:

1. Key Cache (K): Persistent Channel Outliers

   • A tiny subset of hidden feature channels (typically 1–2 channels per head) exhibit magnitudes up to 20× larger than surrounding channels.
   • These outlier channels persist across sequence positions S.
   • Integer Prototype Solution: Per-channel scaling (along head dimension D) isolates outlier channels into their own quantization bins.
   • RoPE Mixing Interaction: Rotary Position Embeddings (RoPE) apply a 2D rotation matrix across channel pairs: Kpost-rope = RΘ,m K_pre-rope. This rotation mixes channel magnitudes across token positions m. In the KVQuant prototype, Key quantization was performed Pre-RoPE to preserve channel outlier sparsity before rotation.

2. Value Cache (V): Localized Token Outliers

   • Outliers in the Value cache do not concentrate in specific feature channels. Instead, specific tokens (such as initial tokens, line breaks, or delimiter punctuation) exhibit elevated activation norms across all channels.
   • Integer Prototype Solution: Per-token scaling (along sequence dimension S) isolates high-magnitude tokens without corrupting preceding or succeeding context tokens. Values are not rotated by RoPE.

3.2 Production Reality: Why FP8 Uses Per-Tensor Scaling

If Keys suffer from persistent channel outliers, why do production engines like vLLM, SGLang, and TensorRT-LLM successfully deploy uniform per-tensor or per-head FP8 scaling?

The answer lies in the dynamic range of floating-point formats:

• Standard INT8 uses a uniform linear grid of 256 steps. An outlier 20× larger than normal values forces the uniform scale factor S = max(|x|) / 127 to expand, crushing non-outlier channels into quantization bins -1, 0, 1.
• FP8 E4M3FN features a logarithmic dynamic range spanning [-448, 448] with 3 mantissa bits and 4 exponent bits. The exponent bits allocate smaller step sizes to small numbers and larger step sizes to large numbers. This dynamic range accommodates channel outlier variance at 8-bit precision for standard sequence lengths (S ≤ 32k) without requiring complex per-channel memory layouts.
• For 4-bit and sub-4-bit integer precision, uniform scaling fails. Frameworks must either adopt specialized algorithmic formulations (such as TurboQuant’s polar coordinates and norm correction) or integrate Pre-RoPE per-channel PagedAttention kernels.

---

4. Physical GPU Execution Paths: Direct FP8 vs. Dequantization

Serving engines do not execute quantized attention through a single universal pathway. The execution flow depends directly on GPU microarchitecture.

═════════════════════════════════════════════════════════════════════════════════
PATH A: DIRECT FP8 ATTENTION (NVIDIA Hopper SM90 with FlashAttention-3)
═════════════════════════════════════════════════════════════════════════════════
Prefill:  Linear Projection (BF16) ──> Quantize K, V to FP8 ──> Write FP8 to HBM
Decode:   1. Read FP8 K, V blocks via TMA directly into SRAM (50% HBM traffic saved)
          2. Query q_t cast to FP8 in SRAM
          3. Tensor Cores execute S = q_t · K^T natively in FP8 (wgmma instructions)
          4. Accumulate partial sums in FP32; evaluate Softmax
          5. Tensor Cores execute O = Softmax(S) · V natively in FP8 (wgmma instructions)
          6. Rescale by (s_q × s_k × s_v) and cast output to BF16
          * NO DEQUANTIZATION PASS TO BF16 OCCURS FOR CACHED TOKENS *

═════════════════════════════════════════════════════════════════════════════════
PATH B: DEQUANTIZE-ON-THE-FLY (NVIDIA Ampere SM80 / vLLM TurboQuant 4-Bit)
═════════════════════════════════════════════════════════════════════════════════
Prefill:  Linear Projection (BF16) ──> Quantize K, V to Low-Bit ──> Write to HBM
Decode:   1. Read Quantized K, V blocks + Scale Factors from HBM into SRAM
          2. ALU unpack & dequantize in registers: K_bf16 = K_quant × scale_k
          3. Tensor Cores execute S = q_t · K_bf16^T using standard BF16 mma.sync
          4. Evaluate Softmax in FP32
          5. ALU unpack & dequantize in registers: V_bf16 = V_quant × scale_v
          6. Tensor Cores execute O = Softmax(S) · V_bf16 using standard BF16 mma.sync
          * REDUCES HBM BANDWIDTH, BUT ADDS ALU OVERHEAD AND ZERO TENSOR CORE SPEEDUP *

4.1 Path A: Direct FP8 Attention (Hopper SM90 + FlashAttention-3)

On NVIDIA Hopper (H100/H200 SM90) architectures running the FlashAttention-3 backend:

1. Asynchronous Memory Transfer: The Tensor Memory Accelerator (TMA) reads FP8 cache blocks directly from global memory (HBM) into shared memory (SRAM), saving 50% of memory bus traffic.
2. Native Tensor Core Math: The kernel issues asynchronous warp-group matrix multiply-accumulate (wgmma.mma_async) instructions directly on FP8 operands.
3. Register Accumulation: Partial products accumulate into 32-bit floating-point registers.
4. Softmax & Value Projection: Softmax is computed in registers; the normalized probabilities are multiplied by FP8 Value blocks directly in Tensor Cores.
5. No Intermediate Dequantization: Cached tokens never expand to 16-bit precision in SRAM or registers. (Note: this direct execution is verified specifically for FlashAttention-3 on Hopper; execution semantics for other backends such as FlashInfer require distinct qualification).

4.2 Path B: Dequantize-on-the-Fly (Ampere SM80 & 4-Bit Engines)

On NVIDIA Ampere (A100 SM80) or with experimental 4-bit formats:

1. Hardware Limitation: Ampere Tensor Cores lack native FP8 arithmetic instructions. Similarly, modern Tensor Cores do not support direct 4-bit floating-point attention matrix multiplication.
2. ALU Conversion: Compressed cache blocks loaded from HBM into SRAM must be explicitly unpacked and multiplied by scale factors in registers, converting them back to BF16 or FP16.
3. Standard GEMM: The attention matmul executes on standard BF16 Tensor Cores (mma.sync).
4. The Performance Tradeoff: This path eliminates memory bus saturation during decode, but consumes additional ALU cycles and register file space. It provides no compute acceleration.

Serving Engine KV Cache Implementation Support

Serving Engine	Precision Format	Key / Value Granularity	Execution Pathway	Maturity Status
vLLM (v0.19+)	FP8 (E4M3FN)	Per-tensor or per-head	Path A on Hopper (FA3); Path B on Ampere	Production Standard
vLLM (Preview)	TurboQuant 4-bit	Algorithmic polar / norm correction	Path B (Dequantize to BF16 for attention)	Experimental Preview
SGLang	FP8 (E4M3FN)	Per-tensor	Path A on Hopper (FA3)	Production Standard
TensorRT-LLM	FP8 (E4M3FN)	Per-head (XQA kernels)	Native FP8 Tensor Core execution	Production Standard
KIVI (Prototype)	2-bit INT2	Per-channel K / per-token V	Path B (Custom CUDA dequantization)	Research Prototype
KVQuant (Prototype)	3-bit / 4-bit	Pre-RoPE per-channel K / per-token V	Path B (Custom CUDA dequantization)	Research Prototype

---

5. Long-Context Reliability, Accumulator Drift, and Layer Skipping

5.1 The Hopper Tensor Core Accumulator Precision Loss

During long-context evaluations of FlashAttention-3 with FP8 KV caches on NVIDIA Hopper GPUs, engineers discovered an unexpected failure mode:

• Symptom: On a 128k-token Needle-in-a-Haystack (NIAH) retrieval test, retrieval accuracy collapsed from 91% in BF16 down to 13% in FP8.
• Root Cause: The failure was not mathematical outlier corruption or RoPE mixing. Instead, it was an FP32 accumulator precision loss inside the Hopper Tensor Core hardware. When summing millions of small FP8 products across contraction loops exceeding N > 100k tokens, the hardware summation tree inside the Hopper matrix engine suffers catastrophic truncation drift.
• The Software Solution: FlashAttention developers implemented a two-level accumulation algorithm (FlashAttention PR #104 and #122). The kernel divides the sequence contraction into blocks, periodically accumulating partial sums into separate register buffers before performing a master reduction.
• The Tradeoff: Two-level accumulation restored NIAH accuracy to 89%. However, for models with large head dimensions D = 256, the additional register allocation induces an approximately 1.6x prefill latency penalty due to register spilling. For standard D = 128 models like Llama-3.1-8B, the prefill latency penalty is negligible (< 3%).
• Blackwell SM100: NVIDIA Blackwell architectures resolve this truncation issue in hardware with redesigned accumulator trees.

5.2 Layer-Skipping for Hybrid Attention Architectures

In models combining full global attention with Sliding-Window Attention (SWA), quantizing all layers uniformly yields diminishing returns while introducing numerical error:

• SWA layers maintain a fixed token buffer (e.g., 4,096 tokens). Their memory footprint does not grow with context length S.
• Quantizing these small, fixed-size buffers introduces noise into local attention without freeing context-proportional VRAM.
• Mitigation: Production engines support selective layer skipping (e.g., vLLM --kv-cache-dtype-skip-layers 0,2,4...) to retain sensitive initial layers and sliding-window layers in full BF16 precision.

---

6. Serving Performance: Single-Stream Latency vs. Batched Throughput

Serving engineers must distinguish single-stream latency from batched system throughput when evaluating KV cache quantization.

6.1 Isolated Benchmark: vLLM on NVIDIA H100 SXM

The following table reports isolated, single-variable benchmark results from vLLM’s evaluation report (“The State of FP8 KV-Cache and Attention Quantization in vLLM”, April 2026) for Llama-3.1-8B on an NVIDIA H100 GPU using FlashAttention-3:

Metric	BF16 Baseline	FP8 KV Cache	Delta	Isolated Benchmark Conditions
Decode ITL Scaling Slope (m)	4.37 × 10^-5 ms/tok	2.37 × 10^-5 ms/tok	-45.8% (54% of base)	Concurrency = 1, S ≤ 20k, H100 SXM, FlashAttention-3
Median Prefill TTFT	763.6 ms	742.8 ms	-2.7%	Concurrency = 8, variable prompt lengths, H100 SXM
Output Token Throughput	450.3 tok/s	517.5 tok/s	+14.9%	Concurrency = 8, ShareGPT trace, Poisson arrival, H100 SXM

6.2 Why Throughput Gains (+14.9%) Lag Memory Bandwidth Halving (-45.8%)

Why does halving KV cache memory traffic reduce the single-stream decode latency slope by 45.8%, yet only increase multi-tenant serving throughput by 14.9%?

Decode Memory Bus Traffic per Token Step:
┌────────────────────────────────────────────────────────┐
│  Model Weights Read (~16 GB for 8B model in BF16)      │  → Constant per step
├────────────────────────────────────────────────────────┤
│  KV Cache Read (B × S × bytes_per_token)               │  → Scales with B and S
└────────────────────────────────────────────────────────┘

1. Model Weight Traffic Dominance: In autoregressive decode, the GPU must stream the entire model’s weights from HBM to on-chip SRAM for every single token generated. For Llama-3.1-8B in BF16, reading model weights moves ≈16.0 GB of data per step. At moderate context lengths, weight streaming dominates memory bus utilization.
2. Arithmetic Intensity Shifts: As concurrent batch size B increases, the model reuses weight tensors across multiple tokens, increasing compute arithmetic intensity and shifting execution toward compute-bound regimes.
3. Prefill Phase Interference: In dynamic serving systems, newly arriving requests undergo compute-heavy prefill phases that interleave with decode iterations, consuming shared streaming multiprocessor (SM) execution cycles.

Consequently, halving KV cache size eliminates memory bandwidth only for the cache portion of the bus load, translating to a modest +14.9% aggregate throughput improvement under real-world multi-tenant traffic.

---

7. Bounded Worked Capacity Scenario

To ground capacity planning, we examine an explicitly hypothetical capacity scenario.

7.1 Setup Parameters

• Hardware Memory Budget: A dedicated KV cache memory pool of 20.0 GiB (21,474,836,480 bytes) on a serving GPU.
• Model: Llama-3.1-8B (L = 32, H_kv = 8, D = 128).
• Target Workload Context: S = 8,192 tokens per sequence.

7.2 Step-by-Step Slot Derivation

1. Element Count per Sequence:

   Elements/seq = 2 × 32 × 8 × 128 × 8192 = 536,870,912 elements = 2^29 elements

2. Per-Sequence Memory Consumption:

   • BF16 (b = 2.0 B):

     536,870,912 × 2 = 1,073,741,824 bytes = **1.0 GiB / session**

   • FP8 (b = 1.0 B):

     536,870,912 × 1 = 536,870,912 bytes = **0.5 GiB / session**

   • NVFP4 Block-16 (b = 0.5625 B):

     536,870,912 × 0.5625 = 301,989,888 bytes = **0.28125 GiB / session**

3. Theoretical Raw Slot Capacity in 20.0 GiB Pool:

   • BF16: ⌊20.0 / 1.0⌋ = 20 raw slots
   • FP8: ⌊20.0 / 0.5⌋ = 40 raw slots
   • NVFP4: ⌊20.0 / 0.28125⌋ = 71 raw slots

7.3 Static Slots vs. Dynamic Serving Concurrency

In production serving frameworks (such as vLLM or SGLang), actual schedulable concurrency differs from raw static slots due to:

• PagedAttention Block Fragmentation: Memory is allocated in discrete block tables (e.g., 16 tokens/block), leaving partially filled blocks at sequence tails.
• Runtime Headroom Reserves: Frameworks configure gpu_memory_utilization (typically 0.90) to preserve memory for intermediate activation spikes during prefill.
• Variable Length Arrival: When requests have varying prompt lengths, dynamic schedulers pack sequences into memory blocks adaptively, avoiding worst-case static allocation.

---

8. Quality Verification Protocol & Hypothesis Testing

Never deploy KV cache quantization based on arbitrary percentage rules of thumb (e.g., “accuracy drop must be ≤ 2%”). Rigorous evaluation requires paired statistical hypothesis testing.

8.1 The Evaluation Suite

1. Reasoning & Code: GSM8K (8-shot chain-of-thought) and HumanEval (0-shot Pass@1).
2. Long-Context Retrieval: Needle-in-a-Haystack evaluated across context depths (16k, 32k, 64k, 128k).
3. Structured Format Adherence: JSON schema compliance rate on structured tool-calling benchmarks.

8.2 Statistical Protocol

1. Run inference across N ≥ 500 evaluation instances for both the unquantized BF16 baseline and the candidate quantized setup under greedy decoding (temperature = 0.0).
2. Compute the paired difference: Δ_i = Metric_quantized,i - Metric_BF16,i.
3. Generate 95% bootstrap confidence intervals over 1,000 resamples.
4. Falsification Rule: Reject the quantized configuration if the lower bound of the 95% confidence interval for accuracy loss exceeds your service’s defined risk tolerance, or if structured schema valid generation drops below 99.0%.

---

9. Testable Learner Verification: Exact Arithmetic Script

Learners can verify the exact tensor byte counts and memory calculus using this self-contained PyTorch script (runnable on CPU or GPU without external serving frameworks):

Llama-3.1-8B

import torch

def verify_kv_cache_arithmetic():
    layers = 32
    kv_heads = 8
    head_dim = 128
    seq_len = 8192
    batch_size = 16

    # Calculate total elements: factor of 2 accounts for Key and Value caches
    total_elements = 2 * layers * kv_heads * head_dim * seq_len * batch_size
    print(f"Total KV cache elements: {total_elements:,}")
    assert total_elements == 8_589_934_592, "Element count mismatch!"

    # Simulate single-layer cache tensors to measure element byte sizes
    k_bf16 = torch.empty((batch_size, kv_heads, seq_len, head_dim), dtype=torch.bfloat16)
    k_fp8 = torch.empty((batch_size, kv_heads, seq_len, head_dim), dtype=torch.float8_e4m3fn)

    bytes_per_elem_bf16 = k_bf16.element_size()  # 2 bytes
    bytes_per_elem_fp8 = k_fp8.element_size()    # 1 byte

    # Total theoretical memory
    total_bf16_bytes = total_elements * bytes_per_elem_bf16
    total_fp8_bytes = total_elements * bytes_per_elem_fp8

    gib_divisor = 1024 ** 3
    bf16_gib = total_bf16_bytes / gib_divisor
    fp8_gib = total_fp8_bytes / gib_divisor

    print(f"BF16 KV Cache: {total_bf16_bytes:,} bytes = {bf16_gib:.4f} GiB")
    print(f"FP8  KV Cache: {total_fp8_bytes:,} bytes = {fp8_gib:.4f} GiB")
    print(f"Compression Ratio: {total_bf16_bytes / total_fp8_bytes:.2f}x")

    assert bf16_gib == 16.0, "BF16 GiB must equal exactly 16.0 GiB!"
    assert fp8_gib == 8.0, "FP8 GiB must equal exactly 8.0 GiB!"
    print("✓ Memory calculus verified successfully.")

if __name__ == "__main__":
    verify_kv_cache_arithmetic()

---

Key takeaways

• KV cache memory scales as 2 × L × H_kv × D × S × B × b. For Llama-3.1-8B at B=16, S=8192, total cache elements equal 8,589,934,592 (2^33), consuming 16.0 GiB in BF16 and 8.0 GiB in FP8.
• FP8 per-tensor scaling requires only 256 bytes of metadata across 32 layers (~1.00000003 B/elem). Sub-4-bit formats carry real metadata overhead: NVFP4 block-16 adds 12.5% (0.5625 B/elem), reducing theoretical 4× compression to 3.56×.
• Grouped-Query Attention (GQA) reduces KV cache by 4× architecturally. Quantizing an 8B GQA model from BF16 to FP8 saves 8.0 GiB, whereas on an MHA equivalent it would save 32.0 GiB, reducing the absolute return of sub-4-bit quantization.
• Path A (Hopper SM90 with FlashAttention-3) executes FP8 attention directly on Tensor Cores (wgmma) with zero dequantization to BF16 for cached tokens. Path B (Ampere SM80 or 4-bit engines) dequantizes back to BF16 in SRAM/registers before attention.
• In research prototypes, Keys exhibit persistent channel outliers (requiring per-channel or Pre-RoPE isolation), while Values exhibit localized token outliers. Production FP8 operates with per-tensor or per-head scaling because E4M3FN dynamic range absorbs standard variance.
• Hopper FP8 Tensor Cores suffer from FP32 accumulator precision loss over long contraction loops (N > 100k), causing unpatched Needle-in-a-Haystack accuracy to drop from 91% to 13%. FlashAttention-3 two-level accumulation restores accuracy to 89%.
• Decode memory traffic consists of model weights plus KV cache reads. Cutting KV cache memory traffic in half reduced H100 single-stream ITL slope by 45.8%, but increased batched throughput by only 14.9% because weight reads and prefill bursts share the bus.
• Experimental 4-bit modes (vLLM TurboQuant) achieve 2.4×–3.4× capacity, but dequantize to BF16, reducing throughput to 66–80% of BF16 and adding 10–68% per-token latency. They trade speed for capacity when VRAM is exhausted.

---

Try it yourself

1. 01Calculate the exact KV cache memory footprint in GiB for Llama-3.1-70B (L = 80, H_kv = 8, D = 128) at batch size B = 32 and sequence length S = 4096 across both BF16 and FP8 E4M3FN. What is the absolute memory saved in GiB?

   Show a hint

   Total elements = 2 × 80 × 8 × 128 × 4096 × 32 = 21,474,836,480 elements (= 20 × 2^30). BF16 (2 B/elem) consumes 42,949,672,960 bytes = 40.0 GiB. FP8 (1 B/elem) consumes 21,474,836,480 bytes = 20.0 GiB. Absolute memory saved = 20.0 GiB.

2. 02A serving team observes that deploying FP8 KV cache on an H100 GPU reduces single-stream decode ITL slope by ~46%, but batched throughput at concurrency 16 increases by only ~15%. Explain the architectural reason for this discrepancy.

   Show a hint

   At concurrency 1, decode is purely memory-bandwidth bound by reading the KV cache as context grows. At concurrency 16, model weight reads (~16 GB per step for an 8B model) remain a large fixed component of memory bus traffic, arithmetic intensity increases, and newly arriving prefill bursts consume shared SM execution cycles.

3. 03Explain why sub-4-bit integer quantization prototypes (like KIVI) require per-channel scaling for Keys but per-token scaling for Values, whereas production FP8 operates reliably with simple per-tensor scaling.

   Show a hint

   Keys feature persistent channel outliers along the head dimension, while Values feature localized token outliers along the sequence dimension. In low-bit uniform integer grids, large outliers collapse non-outlier bins to zero. FP8 E4M3FN features a logarithmic floating-point dynamic range ([-448, 448]) that accommodates standard channel variance without custom per-channel layout.

4. 04Describe the root cause of the Hopper SM90 long-context retrieval failure with FP8 FlashAttention-3, the software algorithm that resolved it, and its associated latency tradeoff.

   Show a hint

   Over contraction dimensions N > 100k tokens, Hopper Tensor Core FP32 accumulator registers suffer from truncation drift, collapsing NIAH retrieval from 91% to 13%. FlashAttention-3 two-level accumulation chunks the contraction and accumulates intermediate sums, restoring accuracy to 89% at the cost of a ~1.6× prefill latency penalty on D = 256 models due to register pressure.

5. 05Using the 20.0 GiB memory pool scenario from the lesson, derive the static slot capacity for 4-bit NVFP4 block-16 storage (b = 0.5625 B) at S = 8192 on Llama-3.1-8B. Compare this with the BF16 slot capacity.

   Show a hint

   Elements per sequence = 536,870,912. Bytes per session = 536,870,912 × 0.5625 B = 301,989,888 bytes ≈ 0.28125 GiB. Raw slot capacity = floor(20.0 / 0.28125) = 71 slots, compared to 20 slots for BF16 (a 3.56× expansion).

---

Sources and scope

Topic	Primary source	Scope used here
FP8 KV Cache Benchmarks	vLLM Blog: “The State of FP8 KV-Cache and Attention Quantization in vLLM” (2026-04-22); vLLM Documentation	ITL slope reduction (-45.8% at C=1), throughput gain (+14.9% at C=8), and native FP8 execution on Hopper SM90 via FlashAttention-3.
Hopper Accumulator Drift	vLLM FP8 Blog (2026-04-22); FlashAttention PR #104 & #122	NIAH retrieval drop (91% → 13%), two-level accumulation restoration to 89%, and 1.6× prefill latency penalty on $D=256$.
FP8 Floating Point Specs	OCP 8-bit Floating Point Specification (OFP8) v1.0 §5.1, Table 2	E4M3FN max value (448.0) and bit format; dynamic range comparison against linear integer grids.
Microscaling Formats	OCP Microscaling Formats (MX) Spec v1.0 Table 1; NVIDIA NVFP4 Blog	NVFP4 block-16 (0.5625 B/elem) and MXFP4 block-32 (0.53125 B/elem) storage arithmetic.
vLLM TurboQuant 4-Bit	vLLM TurboQuant Technical Report (May 11, 2026)	Experimental algorithmic status, measured capacity (3.4× for 4bit_nc, 2.4× for k8v4), dequantization to BF16 (Path B), and throughput reduction (66–80% of BF16).
Sub-4-Bit Integer Prototypes	KIVI: arXiv:2402.02750 (ICML 2024); KVQuant: arXiv:2401.18079 (NeurIPS 2024)	Per-channel Key vs per-token Value asymmetry, Pre-RoPE quantization in KVQuant, and LLaMA-7B 1M/10M token context scaling bounds.

Estimates and Bounds: Single-stream decode slope and multi-stream throughput deltas reflect specific vLLM v0.19.1 runs on H100 SXM with FlashAttention-3; they are not universal across all kernels or GPU generations. The 20.0 GiB memory pool is an explicitly hypothetical capacity baseline. FlashInfer direct FP8 execution semantics remain unverified in primary docs.

---

Related lessons

• Inference Optimisation & AI Infrastructure · Quantize Without Losing the ResultLow-bit LLM Weights and Activations15 min
• Inference Optimisation & AI Infrastructure · Measure and Explain Inference PerformanceWhere Inference Time and Memory Go13 min
• Inference Optimisation & AI Infrastructure · Schedule Requests and Manage KV MemoryPagedAttention: Virtual Memory for the KV Cache35 min