KV Cache — Technologies & Tricks

Updated July 2026 with 2025–2026 SOTA additions — new entries marked ★. Algorithm names link to their papers (arXiv / project page).

July 2026 · Updated Edition

Contents

1. What Is the KV Cache?
2. KV Cache Memory Math
3. Attention-Head KV Sharing: MHA · MQA · GQA · MLA
4. KV Cache Quantization
5. Paged KV Cache (vLLM)
6. Sliding Window + Sink Tokens
7. KV Cache Eviction and Token Pruning
8. Sparse Attention with Cache Efficiency
9. Prefix Caching and Cross-Request Reuse
10. Continuous Batching
11. Speculative Decoding and the KV Cache
12. Long-Context Position Encoding and KV
13. Distributed KV Cache
14. Hardware and Kernel Tricks
15. KV Cache for Multimodal Models
16. Cache Compression Beyond Quantization
17. Streaming and Long Generation
18. Production Patterns
19. Failure Modes
20. Production Stack 2026

Appendix A: Twenty-Five Things to Know
Appendix B: Decision Tree — "How to Cap KV?"
Appendix C: Year-by-Year KV Cache Milestones

1. What Is the KV Cache?

1.1 The autoregressive inference problem

Generating one token at a time, naive attention recomputes \(K, V\) for every previous token at every step:

\[\text{step } t:\ \ \text{recompute } K_{1:t}, V_{1:t} \in \mathbb{R}^{t\times d}.\]

Total cost: \(O(n^2)\) per step \(\to O(n^3)\) for an \(n\)-token generation. Disaster.

1.2 Caching to the rescue

Cache \(K, V\) at every step; reuse from step \(t-1\) in step \(t\):

\[K_t = \mathrm{Concat}(K_{t-1}, k_t), \quad V_t = \mathrm{Concat}(V_{t-1}, v_t),\]

where \(k_t, v_t \in \mathbb{R}^d\) are computed only for the new token. Per-step cost drops to \(O(t\,d_h)\) for the matmul (reading the cache); generation cost becomes \(O(n^2)\) total.

1.3 The two phases of LLM inference

Prefill (compute-bound): process the entire prompt in one forward pass; build the initial cache. FLOPs scale with prompt length squared.

Decode (memory-bound): generate one token at a time, each reading the entire KV cache. Memory bandwidth dominates because each step does just one matmul vector \(\times\) matrix.

Key

Prefill is FLOP-bound; decode is bandwidth-bound. Different optimizations apply: prefill benefits from tensor cores / FP8 matmul; decode benefits from KV cache compression and quantization.

1.4 Why this matters

KV cache size dominates inference memory at long context. For 70B-class models at 100K context, KV cache often exceeds the model weights. Every long-context system optimization is, in some way, KV-cache optimization.

2. KV Cache Memory Math

2.1 The bytes-per-token formula

Key

\[\text{bytes per token} = 2 \cdot L \cdot H_{kv} \cdot d_h \cdot \text{bytes}_{dtype}\] where \(L\) = number of layers, \(H_{kv}\) = number of KV heads (less than \(H\) for GQA/MQA), \(d_h\) = per-head dim, and the 2 is for both \(K\) and \(V\).

2.2 Worked examples

2.3 Total memory at long context

For Llama-3-70B at 100K context, single request:

\[\text{KV}_{\text{mem}} = 100{,}000 \cdot 320\ \text{KB} = 32\ \text{GB}.\]

Model weights: 140 GB BF16 (or 70 GB FP8). KV cache approaches weight cost.

For batch of 32 at 100K context: \(32 \cdot 32 = 1024\) GB. Doesn't fit on any single \(8\times\text{H100}\) node.

2.4 KV vs weights ratio

At long context + large batch, KV memory often exceeds model weights. This flips the optimization priorities:

weight quantization helps less than KV compression.

2.5 Per-token bandwidth cost (decode)

Decode reads the full cache every step. For Llama-3-70B at 32K context:

\[\text{bytes/decode-step} = 32{,}000 \cdot 320\ \text{KB} \approx 10\ \text{GB}.\]

At H100 HBM3 bandwidth (3 TB/s): \(10/3000 = 3.3\) ms per step just for KV read. Sets a hard ceiling on decode speed.

2.6 Why Multi-Latent Attention is so impactful

DeepSeek-V3 cache: \(\sim 70\) KB/token vs Llama-3-70B's 320 KB. \(\sim 5\times\) less KV for a model with \(\sim 10\times\) more total parameters. This is what makes 671B-active-37B serving viable.

3. Attention-Head KV Sharing: MHA · MQA · GQA · MLA

3.1 Multi-Head Attention (MHA)

\(H\) separate query, key, and value heads. Per-head KV: \(H \cdot d_h\). Largest cache; best modeling capacity. Standard pre-2023.

3.2 Multi-Query Attention (MQA, Shazeer)

One KV head shared by all \(H\) query heads.

\[H_{kv} = 1 \Rightarrow \text{cache size} \div H.\]

\(H\times\) smaller cache. Quality cost: noticeable on some tasks. Used in PaLM, Falcon.

3.3 Grouped-Query Attention (GQA)

\(H_{kv} < H\) KV heads, each shared by \(H/H_{kv}\) query heads. Standard \(H_{kv} = 8\).

\[\text{cache size} \div (H/H_{kv}).\]

Quality \(\approx\) MHA, cache \(\sim 4\text{–}8\times\) smaller. Standard since Llama 2; Llama 3, Mistral, Mixtral, Qwen all use GQA.

3.4 Multi-Latent Attention (MLA, DeepSeek-V2/V3)

Project \(KV\) to a low-rank latent cached in place of explicit \(K, V\):

\[c_t = W_{DKV}\, h_t \in \mathbb{R}^{d_c}, \quad d_c \ll d.\]

Cache only \(c_t\). At attention time:

\[K_t = W_{UK}\, c_t, \quad V_t = W_{UV}\, c_t.\]

Up-projections fold into \(W_Q\) and \(W_O\) at inference (no extra matmul):

\[q^\top k = (W_Q^{(i)} h_q)^\top (W_{UK}^{(i)} c_t) = h_q^\top \big(W_Q^{(i)\top} W_{UK}^{(i)}\big) c_t.\]

DeepSeek-V3: \(d_c = 512\) vs \(d = 7168\), \(H = 128\). KV cache reduced \(\sim 14\times\) vs MHA, no quality loss.

3.5 Decoupled RoPE in MLA

RoPE doesn't compose with the absorbed up-projection trick (rotation depends on position). Solution: split each head into a non-RoPE latent component (cached in low rank) and a small RoPE component (cached separately).

The RoPE component is the only "MHA-like" part.

3.6 Comparison table

Key

GQA is the 2024 standard; MLA is the 2025–26 frontier. New decoder LLMs starting from scratch should consider MLA; retrofits from MHA \(\to\) GQA are well-trodden via UPCYCLE recipes.

4. KV Cache Quantization

4.1 Why KV quantization

KV often dominates memory; reducing per-element bytes is direct savings. KV quantization is post-training: the model weights remain FP / BF16, only the cache is quantized.

4.2 Where the outliers live

Activations (and thus \(K, V\)) have channel-wise outliers — a few channels with \(\sim 100\times\) the typical magnitude.

Per-tensor quantization stretches the scale to fit outliers, ruining precision for everyone else. Per-channel scaling fixes this.

4.3 Per-channel KV quantization

\[\hat{k}_{i,c} = \mathrm{round}(k_{i,c}/s_c)\cdot s_c, \quad s_c = \frac{\max_i |k_{i,c}|}{2^{b-1} - 1}.\]

Each channel \(c\) has its own scale. INT8 nearly lossless; INT4 with care.

4.4 KIVI (Liu et al. 2023)

• Per-channel for keys (channel-wise outliers).
• Per-token for values (token-wise stable).
• INT2 / INT4 viable; perplexity loss minimal.

4.5 KV-Cache-INT4, KIVI-2, ZipCache

Variants on the same theme: 4-bit KV with carefully chosen scaling axis. Standard 2024 default for memory-constrained deployments.

4.6 FP8 KV cache

E5M2: wider range, lower precision; good for KV with outliers. E4M3: narrower range, higher precision; usually used for matmul forward.

H100 / B200 native FP8 KV. \(2\times\) cache reduction vs FP16; less quality risk than INT4.

4.7 Mixed-precision KV

• Keys in INT4; values in INT8 (Q has more impact through dot-product).
• Or: outlier channels in FP16, rest in INT4.
• Or: recent tokens in FP16, older in INT4 (recency-aware).

4.8 Calibration

KV quantization rarely needs calibration data: the cache observes activations during runtime, and per-channel/per-token scales adapt online.

4.9 Quality results (typical)

★ 2026 SOTA update — Sensitivity-aware mixed-precision KV

• KVTuner: offline multi-objective search picks hardware-friendly per-layer mixed-precision KV pairs (applied online with no per-token overhead), achieving nearly lossless ~3.25-bit KV for Llama-3.1-8B and 4-bit for sensitive models like Qwen2.5-7B (ICML 2025).

5. Paged KV Cache (vLLM)

5.1 The fragmentation problem

Naive contiguous KV allocation per request:

• Pre-allocate max length per request.
• Variable lengths waste memory (pad).
• Hard to insert / preempt requests.
• Prefix sharing is impossible (each request has its own block).

5.2 The PagedAttention idea (Kwon et al. 2023, vLLM)

Treat KV as virtual memory:

• Physical KV blocks of fixed size (e.g., 16 tokens).
• Each request has a block table mapping logical block index \(\to\) physical block.
• Allocate physical blocks on demand.

5.3 Benefits

• No fragmentation: arbitrary lengths fit; fully utilize memory.
• Prefix sharing: multiple requests with same prompt prefix share physical blocks; CoW on divergence.
• Higher concurrency: more requests in same memory.
• Preemption: swap blocks out to CPU on overload.

5.4 The PagedAttention kernel

Custom CUDA kernel reads KV blocks via the per-request block table during attention. Indirect-indexed memory access; designed to maintain throughput despite the indirection.

5.5 Block size trade-off

• Small (4–8): fine-grained sharing; more page-table overhead.
• Standard (16): vLLM default.
• Large (32+): less overhead; more internal fragmentation.

5.6 CoW for branching / speculative decoding

When two requests share a prefix and diverge: the divergent suffix copies the affected block; prefix blocks remain shared. Critical for tree-style speculative decoding (Medusa, EAGLE).

5.7 vLLM's automatic prefix caching

The page table can be hashed; identical prefixes across requests automatically detected and reused, no application code needed.

5.8 Beyond vLLM

Paged KV is now standard:

• TensorRT-LLM: paged KV manager.
• SGLang: RadixAttention (radix tree of cached prefixes).
• TGI (HuggingFace): paged KV.
• LMDeploy: blocked KV.

6. Sliding Window + Sink Tokens

6.1 Sliding window attention

Each token attends only to the last \(w\) tokens. Cache size capped at \(w\) regardless of generation length:

\[\text{cache} \le 2 \cdot L \cdot H_{kv} \cdot d_h \cdot w \cdot \text{bytes}.\]

Used in Mistral 7B (\(w = 4096\)), Mistral Nemo, Llama 3.

6.2 Why naive sliding window breaks long generation

At step \(t > w\), dropping the early tokens causes catastrophic perplexity spikes. The first few tokens act as

"attention sinks" — the softmax dumps unwanted attention mass on them. Without sinks, mass redistributes onto random tokens and meaning collapses.

6.3 Attention sinks (StreamingLLM, Xiao et al.)

Always keep the first \(k\) (\(\sim 4\)) tokens in cache, plus a sliding window of \(w\) recent tokens:

\[\text{cache} \le 2 \cdot L \cdot H_{kv} \cdot d_h \cdot (k + w).\]

Effective infinite context with constant cache. Standard in Mistral, StreamingLLM, many production stacks.

6.4 Why sinks work (intuition)

Pre-trained Transformers learn to use a few tokens as "do nothing" anchors. Removing them forces the softmax to redistribute, which corrupts the attention pattern. Keeping a few sinks preserves the trained behavior.

6.5 Cyclic / rolling buffer

• Allocate a fixed-size buffer; overwrite oldest slots when full.
• Combined with sinks: sinks pinned in fixed slots; rolling window cycles.
• Used in Mistral's Rolling Buffer Cache.

6.6 Local attention with global tokens

• Local tokens: sliding window (cheap).
• Global tokens: attended to and from by every position (expensive but scarce).
• Examples: Longformer, BigBird (encoder-side).

7. KV Cache Eviction and Token Pruning

7.1 The premise

Not every token in the cache contributes equally. Drop the "unimportant" ones to keep cache small while preserving quality.

7.2 H2O (Heavy-Hitter Oracle, Zhang et al. 2023)

Insight: a small fraction of tokens ("heavy hitters") receive most of the attention mass. Keep recent tokens + heavy hitters; evict the rest.

• Compute per-token attention score sum across past steps.
• Eviction policy: keep top-k scoring + recent window.
• Cache reduced by \(\sim 5\times\) with \(\sim 1\%\) perplexity loss.

7.3 Scissorhands

Attention sparsity is persistent: tokens unimportant once tend to remain unimportant. Tracks per-token importance over time; evicts low-persistence tokens.

7.4 SnapKV (Li et al. 2024)

For long-prompt scenarios:

• During prefill, examine the last few tokens' attention to all prior tokens.
• Pool importance scores; keep top-k.
• Discard the rest before decode starts.
• Massive reduction (e.g., 100K \(\to\) 4K tokens) with minimal quality loss on QA / summarization.

7.5 Pyramid KV (Cai et al. 2024)

Lower layers retain more tokens; higher layers fewer. Insight: deeper layers concentrate attention on fewer tokens.

Memory savings without uniform pressure.

7.6 FastV (Chen et al. 2024)

Vision-token specific: in VLMs, the LLM rarely attends to vision tokens after a few layers. Drop most vision tokens after layer \(K\). \(\sim 50\%\) FLOP reduction; minimal quality loss.

7.7 Quest (Tang et al. 2024)

Query-aware token selection: at each query, retrieve only the top-k most-relevant past KV blocks. Combines paged KV with importance-based retrieval. Long-context speed-up with quality preserved.

7.8 StreamingLLM eviction policy

Combine sinks + sliding window. The implicit eviction is just "oldest non-sink token." Simple, robust.

7.9 Adaptive / dynamic eviction

• Per-layer different keep-fractions.
• Per-head different policies.
• Schedule-based (more aggressive late in generation).

7.10 Eviction trade-offs

• Aggressive eviction \(\to\) memory wins; quality risk.
• Importance-based \(\to\) keeps quality; computation overhead.
• Recency-only (sliding window) \(\to\) predictable; loses long-range context.
• Combine: sinks + recency + importance.

★ 2026 SOTA update — Query-agnostic KV eviction

• KVzip: scores KV-pair importance by how well the LLM can reconstruct the original context from the cache, producing a single compressed cache reusable across diverse queries; 3-4x smaller KV and ~2x faster FlashAttention decode with negligible loss (LLaMA3.1/Qwen2.5/Gemma3, up to 170K).

★ 2026 SOTA update — Adaptive-budget KV compression

• Ada-KV: allocates the eviction budget non-uniformly across attention heads based on each head's attention concentration, improving budget use; plugs into SnapKV/PyramidKV (NeurIPS 2025).
• RocketKV: two-stage training-free scheme, coarse permanent eviction then fine-grained top-k hybrid sparse attention, reaching up to 400x compression and 3.7x decode speedup with negligible accuracy loss.

8. Sparse Attention with Cache Efficiency

8.1 Native Sparse Attention (DeepSeek NSA)

Hardware-friendly hybrid sparse training pattern:

• Compressed branch: down-sampled global tokens.
• Selected branch: top-k blocks of past KV (importance-driven).
• Sliding branch: local window.

End-to-end trained with sparsity from the start. KV reads \(\sim 10\times\) less than dense at long context.

8.2 Mixture-of-Attention (MoBA)

Per-token selection of which past block to attend to. Trained sparse from scratch.

8.3 Mamba / SSM

Recurrent state of fixed size \(h_t\) replaces KV cache entirely:

\[h_t = \bar{A}\, h_{t-1} + \bar{B}\, x_t.\]

Cache size \(O(d^2)\) per layer, independent of sequence length. Strong at very long context; quality gap vs attention at moderate context.

8.4 RWKV / RetNet (linear attention)

Linear-attention recurrence with constant per-token state:

\[S_t = \gamma\, S_{t-1} + k_t^\top v_t.\]

\(S_t \in \mathbb{R}^{d\times d}\) regardless of \(t\). Train in parallel like Transformer; run as RNN at inference.

8.5 Hybrid: linear + softmax

MiniMax-01, Jamba (Mamba + Transformer), Zamba, Hymba: alternate linear/recurrent layers with softmax attention layers. Best of both: long-context efficiency + softmax expressiveness.

8.6 Sliding-window Mamba layers

Some hybrid models use sliding-window softmax for short range and Mamba for long range.

★ 2026 SOTA update — DeepSeek Sparse Attention (V3.2)

• DeepSeek Sparse Attention (DSA): a lightning indexer scores all preceding tokens with a ReLU-gated dot product, then top-k selection cuts core attention from \(O(L^2)\) to \(O(Lk)\); ships in DeepSeek-V3.2 for ~3-6x cheaper 128K-context serving with near-identical quality.

★ 2026 SOTA update — Block-sparse prefill kernels

• XAttention: uses the sum of antidiagonal values as a cheap proxy for block importance to prune non-essential attention blocks, giving training-free block-sparse prefill with up to 13.5x attention speedup at full-attention accuracy on RULER/LongBench.

★ 2026 SOTA update — Hybrid linear-attention KV cut

• Kimi Linear: layerwise hybrid of Kimi Delta Attention (a gated DeltaNet variant) and MLA that outperforms full attention while reducing KV cache by up to 75% and delivering up to 6x decode throughput at 1M context.

9. Prefix Caching and Cross-Request Reuse

9.1 The opportunity

Many production workloads share long prefixes:

• System prompts (same for every request).
• Multi-turn conversation (cumulative history).
• Few-shot prompts (same exemplars).
• RAG (same retrieved docs across batched queries).

9.2 Prefix cache

Pre-compute KV for the shared prefix once; reuse across requests. Speedup proportional to prefix length / total length.

9.3 Anthropic prompt caching

Mark a prefix in the request; Anthropic's API caches its KV for \(\sim 5\) minutes; subsequent requests with the same prefix get up to 90% cost discount on cached tokens. Granular checkpointing: multiple cache breakpoints per request.

9.4 OpenAI prompt caching

Similar pattern: GPT-4o and beyond cache prefixes automatically; cache hit gives \(\sim 50\%\) discount. No application code needed; automatic.

9.5 Gemini context caching

Explicit CachedContent resource; created once, referenced by future requests. Pricing model with TTL.

9.6 vLLM's automatic prefix caching

• Compute hash of every block during prefill.
• Insert into hash table; subsequent requests look up by hash.
• LRU eviction when memory pressure.

No code changes required; cross-user sharing handled by hashing.

9.7 SGLang's RadixAttention

Radix tree of cached prefixes; request finds the longest prefix match. More sophisticated than hash-based matching for branching dialogs.

9.8 KV cache hot/cold tiering

• Hot tier: GPU HBM; recently-used / popular prefixes.
• Warm tier: CPU RAM; less-used.
• Cold tier: disk / object storage; very long-tail.
• Background prefetch on cache miss.

9.9 Persistent KV across sessions

• Per-user persistent cache for system + memory.
• Multi-tenant isolation.
• Storage cost vs recompute trade-off.

★ 2026 SOTA update — Non-prefix cross-request KV reuse

• CacheBlend: fuses precomputed KV caches of arbitrary (non-prefix) chunks and selectively recomputes a small subset of tokens, matching full-prefill quality for RAG where retrieved docs are not a shared prefix.
• EPIC: position-independent caching whose LegoLink algorithm repairs the per-chunk attention-sink effect in \(O(kN)\), enabling modular reuse of chunk KV under varying prefixes for few-shot and RAG serving (ICML 2025).

10. Continuous Batching

10.1 The problem with static batching

Static batching: wait for \(B\) requests, run together; output all when slowest finishes. Latency dominated by slowest; resources idle while waiting.

10.2 Continuous batching (Orca, vLLM)

At each token step:

• Accept new requests (insert into batch).
• Drop completed requests (remove from batch).
• Run one decode step on the current batch.

Per-token batching at the kernel level; throughput up \(5\text{–}10\times\) vs static.

10.3 Per-request KV state

Each request has its own KV cache; decode kernel handles per-request access via paged KV. Variable lengths within batch handled cleanly by PagedAttention.

10.4 Prefill + decode mixing

Some requests in prefill phase (compute-bound), others in decode (memory-bound). Mixing improves utilization but complicates scheduling.

10.5 Disaggregated prefill / decode (DistServe)

Run prefill and decode on separate clusters:

• Prefill cluster optimized for compute (high-FLOP GPUs, no memory pressure).
• Decode cluster optimized for memory bandwidth (HBM3e, high concurrency).
• KV transferred from prefill node to decode node (RDMA / NVLink).

Improves both prefill latency and decode throughput. Used in production at Anthropic, OpenAI, Mooncake.

10.6 Mooncake (Moonshot AI)

Disaggregated serving with dedicated KV cache pool:

• KV store as a separate service.
• Prefill node \(\to\) KV store \(\to\) decode nodes.
• Cache reuse across nodes.
• Open-sourced production stack.

10.7 LMCache

KV cache layer for LLM serving; shares prefixes across vLLM instances; CPU/disk/remote tiering. Open-source.

11. Speculative Decoding and the KV Cache

11.1 The setup

Draft model proposes \(k\) tokens; target verifies all in one forward pass. Target's KV cache must accommodate the verification.

11.2 Linear speculative

Single chain of \(k\) candidate tokens. KV cache extended speculatively; on rejection, truncate.

11.3 Tree attention (Medusa, EAGLE)

Propose a tree of candidate continuations. Single target forward pass evaluates the entire tree via custom causal mask:

\[M_{ij} = 0 \text{ if } j \text{ ancestor of } i, \quad -\infty \text{ otherwise.}\]

KV cache stores all tree positions; longest accepted prefix committed; rest discarded.

11.4 KV management for tree decoding

• Each tree node has its own KV slot.
• Reject branches: free their KV.
• Accept path: commit; extend root.
• Paged KV's CoW handles branching cleanly.

11.5 EAGLE-2 / EAGLE-3

Dynamic tree construction based on draft model confidence. Larger trees in high-uncertainty regions; smaller (or chain) in confident regions. Adapts KV usage per step.

11.6 Lookahead decoding

Maintain a verification window of \(W\) tokens generated \(W\) steps ago + an n-gram pool. Verify multiple positions in parallel. KV cache used for both lookahead positions and committed prefix.

12. Long-Context Position Encoding and KV

12.1 The extrapolation problem

A model trained at \(L_{train}\) context degrades at \(L_{eval} > L_{train}\). Need position-encoding tricks that extrapolate; KV cache must store positions correctly.

12.2 Position Interpolation (PI)

Linearly compress positions during inference: \(p \to p \cdot L_{train}/L_{eval}\). KV cache stores compressed positions. Cheap; needs short fine-tune to recover.

12.3 NTK-aware scaling

Scale RoPE base wavelength so high-frequency dims stay intact, low-frequency dims stretch. KV cache stores RoPE-rotated keys with the new base.

12.4 YaRN

Piecewise rescaling by frequency band + temperature adjustment in attention logits. Achieves \(10\text{–}32\times\) context extension with light fine-tuning.

12.5 LongRoPE

Per-dimension rescaling factors found via evolutionary search. Extends Llama 2 to 2M context.

12.6 Self-Extend (inference only)

Bin positions in groups of \(G\) at distant ranges; keep fine positions only locally. No fine-tune. KV cache stores normal positions; the bucketing is applied on read.

12.7 DCA (Dual Chunk Attention)

Split sequence into chunks; intra-chunk attention is normal; inter-chunk attention uses position-shifted RoPE.

Reduces effective position range without fine-tune.

12.8 Implications for KV cache

• Some tricks just reinterpret positions on read (PI, NTK, Self-Extend, DCA).
• Some require re-rotation of cached \(K\) (rare).
• In practice: store at training-position; apply scaling on read for inference-time tricks.

13. Distributed KV Cache

13.1 Tensor parallel KV

KV split by head across TP devices; each device holds its own head shards. Standard with Megatron-style TP.

13.2 Sequence parallelism

For LayerNorm and dropout, shard activations along sequence. Doesn't change KV cache layout but reduces activation memory.

13.3 Ring attention

For very long sequences across \(P\) devices:

1. Each device holds its \(Q\) shard + initial \(K, V\) shards.
2. Compute attention with own \(K, V\); pass \(K, V\) around the ring.
3. Accumulate via online softmax.
4. Total \(P\) rotations.

KV stays distributed; no single device holds the full cache.

13.4 Striped attention

Variant of ring attention: striping \(K, V\) blocks differently to balance work in causal-mask scenarios.

13.5 Context parallelism (NVIDIA Megatron)

Sequence parallelism specifically for the attention compute, with all-gather of KV at attention time. Different decomposition trade-off than ring.

13.6 Distributed KV across pipeline stages

With pipeline parallelism, each stage holds the KV for its own layers. Total KV memory split across PP stages naturally; no extra communication.

13.7 Cross-node KV transfer (disaggregated serving)

• RDMA / GPUDirect for fast KV transfer.
• NVLink within node (\(\sim 900\) GB/s).
• InfiniBand across nodes (\(\sim 200\) Gb/s).
• Compress KV before transfer (FP8 or INT4) to halve / quarter transfer time.

14. Hardware and Kernel Tricks

14.1 FlashAttention with KV cache

FA2 / FA3 support KV cache via the varlen API: variable sequence lengths in one batched call. PagedAttention is a generalization for non-contiguous KV.

14.2 FlashDecoding / FlashDecoding++

Decode-specific FA variant: parallelize across the KV sequence (large) instead of just over heads (small). Critical for long-context decoding when batch size is small.

14.3 Hopper (H100) async TMA for KV

Tensor Memory Accelerator copies KV blocks asynchronously while the compute units do matmul. Hides KV-load latency behind compute.

14.4 FP8 attention with KV

KV stored FP8; matmul accumulator in FP32. Per-block scaling factors stored alongside KV. Halves memory, doubles bandwidth utilization vs FP16.

14.5 2:4 sparsity in projections

\(Q, K, V, O\) projection matrices can be 2:4-sparsified post-training. Reduces compute; KV cache itself is dense (no sparsification of stored values).

14.6 Per-block scaling

For FP8 / INT4 KV, scaling factors stored per-block (e.g., per 128 elements). Allows decode kernel to apply scale on read without per-token overhead.

14.7 Custom kernels per stack

• vLLM: PagedAttention CUDA kernels.
• SGLang: RadixAttention + tree-decoding kernels.
• TensorRT-LLM: NVIDIA's tuned kernels per architecture.
• ThunderKittens: low-level kernel kit; competitive with FA3.
• Triton implementations: pythonic, customizable.

15. KV Cache for Multimodal Models

15.1 Image tokens

Each image becomes 100s–1000s of tokens. KV cache scales accordingly. For 16 images per prompt \(\times\) 256 tokens each = 4096 image tokens added to text.

15.2 High-res images

LLaVA-NeXT AnyRes tiling: a 1024-px image can become 5–10 tiles \(\times\) 256 tokens each \(= \sim 2500\) tokens of cache.

15.3 Video tokens (the big one)

1-hour video at 1 fps with 256 tokens/frame: \(\sim 920\text{k}\) tokens. KV at 320 KB/token (Llama-3-70B): \(\sim 295\) GB just for KV. Compression mandatory.

15.4 Vision-token KV eviction (FastV)

After layer \(K\) (\(K \sim 2\text{–}4\)), the LLM rarely attends to image tokens. Drop most vision-token KV after layer \(K\).

Massive savings for VLMs:

\[\text{VLM KV} = K \cdot \text{full} + (L - K) \cdot \text{text-only}.\]

15.5 Cross-modal sharing

• Same image used in multiple turns: cache image token KV once.
• Same system prompt across modalities: shared.
• Per-modality cache pools.

15.6 Long-video VLMs

Qwen2.5-VL, Gemini 2.5: 1-hour to multi-hour video native. Strategies:

• Aggressive token compression at the encoder.
• Hierarchical pooling across time.
• Q-Former summarization.
• Selective frame attention.

16. Cache Compression Beyond Quantization

16.1 Low-rank decomposition

Cache the SVD of \(K, V\) instead of raw matrices:

\[K \approx U_K \Sigma_K V_K^\top,\]

cache \(U_K \Sigma_K\) (rank-\(r\)). Per-token \(K\) reconstructed on read. Aggressive savings for redundant content.

16.2 MLA (already covered)

DeepSeek's MLA is a learned low-rank factorization built into the architecture. Cleanest realization of the low-rank-cache idea.

16.3 KV cache distillation

Train a smaller model that achieves same quality with less KV per token. LoRACPP: distill compression layers post-training.

16.4 Token merging in cache

Periodically merge similar adjacent tokens in the cache (cosine similarity). Saves space; some quality cost.

16.5 Compressed long-term memory

Older cache periodically compressed (averaged, pooled, or summarized by an auxiliary model). Recent kept full-resolution.

16.6 Cache fingerprinting + dedup

For repeated content (system prompts, cited docs), detect duplicates by hash; share single canonical KV. Used in vLLM automatic prefix caching.

17. Streaming and Long Generation

17.1 Streaming inference patterns

• Output tokens streamed to user as generated.
• Cache grows during streaming.
• Cap with sliding window for unbounded streams.
• Periodic checkpointing of cache state.

17.2 Cyclic / circular cache

Pre-allocated buffer of fixed size; overwrite oldest slots when full. Combined with sinks (preserve first \(k\) slots).

17.3 Memorizing Transformer (Wu et al.)

External memory bank of past KVs; current attention augmented with kNN retrieval over the memory. Cache stays bounded; effective context unbounded.

17.4 Compressive Transformer

Two memory tiers: short-term (recent KVs) + compressed long-term (pooled / convolved). Trade exactness for capacity.

17.5 RMT (Recurrent Memory Transformer)

Pass small memory tokens between segments; each segment processes its KV normally + reads memory tokens.

Read-write memory, similar to Neural Turing Machine.

17.6 Practical streaming LLM stacks

• StreamingLLM: window + sinks; production-friendly.
• Mistral Rolling Buffer: fixed-size ring + sinks.
• Custom: combine sliding window with periodic summarization.

18. Production Patterns

18.1 Per-tier memory budget

• Per-request: cap context by user tier.
• Per-model: cap concurrent requests by KV memory.
• Auto-scaling triggered by KV memory pressure.

18.2 Preemption

On overload, preempt low-priority requests:

• Swap their KV blocks to CPU memory.
• Resume by swapping back when capacity returns.
• Or recompute prefill on resume.

18.3 KV cache eviction policies (multi-request)

• LRU on cached prefixes.
• Pinned: never evict (system prompts).
• TTL-based (Anthropic-style 5-minute window).
• Cost-aware: prefer evicting cheap-to-recompute prefixes.

18.4 Observability

• Per-request cache hit rate.
• Per-request KV memory usage.
• Cluster-level cache utilization.
• Eviction rate.
• Cache reconstruction latency.

18.5 Cost models

• Cached input tokens: 50–90% discount (provider-dependent).
• Pricing tier by cache hit / miss.
• Pre-warm popular prefixes.
• Track per-tenant cache utilization.

18.6 Multi-tenant isolation

• Per-tenant cache pool (avoid cross-tenant leakage).
• Hash with tenant ID to prevent collision.
• Quota per tenant.
• Audit log of cache access.

19. Failure Modes

19.1 Out-of-memory (OOM)

KV cache exhausts memory at long context or high batch. Fix: paged KV, eviction, GQA/MLA, quantization, swap to CPU.

19.2 Cache thrashing

Frequent eviction-and-reload cycle as requests overlap. Fix: better eviction policy (LRU, importance), preemption with swap.

19.3 Quality loss from compression

Aggressive INT4 / token eviction degrades long-context quality. Fix: hold-out eval per compression setting; tune per-deployment.

19.4 Position drift in eviction

After eviction, position indices break the assumed contiguity. Fix: re-index positions; or use position-tolerant attention (RoPE relative).

19.5 Streaming long-context collapse

Without sinks, sliding window destroys long generation. Fix: always include sinks in streaming setups.

19.6 Cache leakage across tenants

Hash collisions or shared keys can cause prefix sharing across tenants \(\to\) privacy bug. Fix: prefix tenant ID into hash; per-tenant cache pool.

19.7 Prefill / decode imbalance

Disaggregated stack with mismatched compute / memory. Fix: monitor utilization; rebalance prefill / decode capacity.

19.8 Speculative decoding cache mismatch

Tree branches not properly freed on rejection \(\to\) memory leak. Fix: explicit cleanup hooks; CoW correctness tests.

20. Production Stack 2026

Appendix A: Twenty-Five Things to Know

1. KV bytes/token formula: \(2 \cdot L \cdot H_{kv} \cdot d_h \cdot\) bytes dtype.
2. Llama-3-70B GQA-8 in FP16: 320 KB/token.
3. DeepSeek-V3 MLA: \(\sim 70\) KB/token (\(14\times\) smaller than equivalent MHA).
4. Prefill is FLOP-bound; decode is bandwidth-bound.
5. GQA-8 is the 2024 standard; MLA is the 2025–26 frontier.
6. MLA absorbs up-projections into \(W_Q\), \(W_O\) at inference.
7. Decoupled RoPE: split each head into RoPE + non-RoPE components.
8. KV quantization: per-channel for \(K\), per-token for \(V\) (KIVI).
9. INT4 KV typically loses 1–2% quality; INT2 loses 5–10%.
10. FP8 (E5M2) for KV; FP8 (E4M3) for forward.
11. PagedAttention block size: 16 tokens default.
12. Automatic prefix caching: hash blocks; reuse on match.
13. SGLang RadixAttention: prefix tree of cached blocks.
14. Sliding window + sinks (StreamingLLM) for unbounded streaming.
15. Sinks (\(k \sim 4\)) prevent softmax-redistribution collapse.
16. H2O eviction: keep heavy hitters + recent.
17. SnapKV: prefill-time importance pooling.
18. Pyramid KV: deeper layers retain fewer tokens.
19. FastV: drop vision tokens after layer \(K\) in VLMs.
20. Quest: query-aware KV block retrieval.
21. Continuous batching: per-token, not per-request.
22. DistServe / Mooncake: disaggregated prefill / decode.
23. Anthropic prompt cache: 90% discount on cached prefix.
24. Tree attention (Medusa/EAGLE) needs CoW paged KV.
25. Mamba / RWKV / RetNet replace KV with constant-size state.

Appendix B: Decision Tree — "How to Cap KV?"

1. Designing a new model from scratch? \(\to\) MLA (DeepSeek pattern). Best long-term cache efficiency.
2. Retrofitting an existing model? \(\to\) GQA via continued pretraining (Llama 3 pattern).
3. Need streaming / unbounded context, OK with bounded effective context? \(\to\) Sliding window + sinks (Mistral / StreamingLLM).
4. Long context, full attention needed but memory tight? \(\to\) INT4 KV (KIVI) or FP8 KV.
5. Cloud serving with shared prefixes? \(\to\) Paged KV + automatic prefix caching (vLLM / SGLang).
6. Many requests, mixed lengths? \(\to\) Continuous batching (vLLM, SGLang, TGI).
7. Long prompts, short outputs (RAG, document QA)? \(\to\) SnapKV prefill-time compression.
8. VLM with many image tokens? \(\to\) FastV drop vision tokens after early layers.
9. Throughput at scale? \(\to\) Disaggregated prefill/decode + Mooncake-style KV pool.
10. Edge / on-device? \(\to\) GQA + INT4 KV + sliding window.

Appendix C: Year-by-Year KV Cache Milestones

• 2017: Original Transformer; KV cache implicit in autoregressive decoding.
• 2019–2020: MQA (Shazeer); Multi-Query as first cache reduction.
• 2022: ALiBi enables length extrapolation without re-rotating KV; Mistral introduces sliding window.
• 2023: GQA standardized via Llama 2; vLLM (PagedAttention) debuts; StreamingLLM (sinks); H2O eviction.
• 2024: KIVI (INT4 KV); SnapKV; FastV (vision-token drop); MLA in DeepSeek-V2; SGLang RadixAttention; Anthropic prompt caching API; OpenAI prompt cache; DistServe; Mooncake released.
• 2025: DeepSeek-V3 (MLA + FP8 KV at 671B); Qwen3 GQA + long context; tree-decoding production (EAGLE-3); native sparse attention (DeepSeek NSA); MiniMax-01 hybrid attention; Quest dynamic block retrieval.
• 2026: MLA / fine-grained KV-sharing standard; FP8 / INT4 KV mainstream; disaggregated serving widespread; multimodal KV optimization (FastV-style) standard in VLMs.