★ If you read Section II.06 (virtual memory), this file will feel inevitable rather than clever. That’s the point — PagedAttention is operating systems applied to KV cache.
1. What is it?#
Storing the KV cache in fixed-size non-contiguous blocks with a per-sequence block table mapping logical positions to physical blocks.
CONTIGUOUS (before):
Sequence A: [████████████████░░░░░░░░░░░░░░░░] reserved 2048, used 700
Sequence B: [████░░░░░░░░░░░░░░░░░░░░░░░░░░░░] reserved 2048, used 180
↑ wasted (internal fragmentation)
+ external fragmentation between allocations
PAGED (after):
Physical blocks: [0][1][2][3][4][5][6][7][8][9]...
A's block_table: [3, 7, 1, 9] ← 4 blocks × 16 tokens = 64 tokens
B's block_table: [5, 2]
Free list: [0, 4, 6, 8, ...]
No reservation. No fragmentation. Blocks allocated one at a time as needed.Diagram — Logical KV blocks map to physical blocks through a block table#
flowchart LR
subgraph LOG["Logical view - per request"]
direction TB
A0["Req A block 0"]
A1["Req A block 1"]
A2["Req A block 2"]
B0["Req B block 0"]
B1["Req B block 1"]
end
subgraph TBL["Block tables"]
direction TB
TA["A: 7, 1, 4"]
TB2["B: 3, 5"]
end
subgraph PHY["Physical KV pool - fixed-size blocks"]
direction TB
P1["block 1"]
P3["block 3"]
P4["block 4"]
P5["block 5"]
P7["block 7"]
PF["free blocks ..."]
end
A0 --> TA
A1 --> TA
A2 --> TA
B0 --> TB2
B1 --> TB2
TA --> P7
TA --> P1
TA --> P4
TB2 --> P3
TB2 --> P5
class A0,A1,A2,P7,P1,P4 compute
class B0,B1,P3,P5 memory
class TA,TB2 queue
class PF neutral2. Why does it exist?#
Because the contiguous KV cache wastes 60-80% of GPU memory, and GPU memory is your capacity.
The vLLM paper measured actual utilization in existing systems:
System KV memory actually holding useful tokens
Orca (contiguous) 20.4% - 38.2%
vLLM (paged) 96.3%A 2.5-4x capacity increase from a memory-management change. No new hardware, no quality tradeoff, no different model.
The three wastes it eliminates#
1. INTERNAL FRAGMENTATION (reservation waste)
You must reserve for max_tokens because you don't know the output length.
Request generates 80 tokens, reserved 2048 → 96% of that allocation wasted.
2. EXTERNAL FRAGMENTATION
After many variable-size allocate/free cycles, free memory exists but not
contiguously. New requests fail despite available memory.
3. NO SHARING
Ten requests with the same 2000-token system prompt store ten copies.3. Simple analogy#
A hotel that rents rooms instead of floors.
Old system: when a guest checks in, you reserve an entire floor because they might bring a large party. Most bring two people. 95% of the hotel sits empty while you turn people away.
New system: rent one room at a time as the party grows. A shared conference room can be used by several parties at once. Occupancy goes from 30% to 96%.
The block table is the front desk’s ledger: “party A is in rooms 3, 7, 1, 9.” The rooms need not be adjacent, and the guest never notices.
4. Tiny example#
A working block manager in ~60 lines:
// paged.go — the block manager at the heart of PagedAttention.
package main
import (
"errors"
"fmt"
)
type BlockManager struct {
blockSize int
free []int // stack of free physical block ids
refcount map[int]int // block id -> number of sequences using it
}
func NewBlockManager(numBlocks, blockSize int) *BlockManager {
m := &BlockManager{blockSize: blockSize, refcount: map[int]int{}}
for i := numBlocks - 1; i >= 0; i-- {
m.free = append(m.free, i)
}
return m
}
var ErrNoBlocks = errors.New("out of KV blocks")
func (m *BlockManager) Allocate(n int) ([]int, error) {
if len(m.free) < n {
return nil, fmt.Errorf("%w: need %d, have %d", ErrNoBlocks, n, len(m.free))
}
blocks := append([]int{}, m.free[len(m.free)-n:]...)
m.free = m.free[:len(m.free)-n]
for _, b := range blocks {
m.refcount[b] = 1
}
return blocks, nil
}
func (m *BlockManager) Free(blocks []int) {
for _, b := range blocks {
if m.refcount[b]--; m.refcount[b] == 0 { // only free when nobody references it
delete(m.refcount, b)
m.free = append(m.free, b)
}
}
}
// Fork shares blocks copy-on-write: used by beam search, parallel sampling, prefix caching.
func (m *BlockManager) Fork(blocks []int) []int {
for _, b := range blocks {
m.refcount[b]++
}
return append([]int{}, blocks...) // same physical blocks, new logical sequence
}
func (m *BlockManager) BlocksNeeded(tokens int) int { return (tokens + m.blockSize - 1) / m.blockSize }
type Sequence struct {
mgr *BlockManager
BlockTable []int // logical block index -> physical block id
Length int
}
func NewSequence(m *BlockManager, promptLen int) (*Sequence, error) {
blocks, err := m.Allocate(m.BlocksNeeded(promptLen))
return &Sequence{m, blocks, promptLen}, err
}
func (s *Sequence) AppendToken() error {
s.Length++
if s.mgr.BlocksNeeded(s.Length) > len(s.BlockTable) {
b, err := s.mgr.Allocate(1) // grow by ONE block
if err != nil {
return err
}
s.BlockTable = append(s.BlockTable, b...)
}
return nil
}
// PhysicalIndex: where token `pos` of this sequence actually lives in the KV pool.
func (s *Sequence) PhysicalIndex(pos int) int {
return s.BlockTable[pos/s.mgr.blockSize]*s.mgr.blockSize + pos%s.mgr.blockSize
}
func main() {
mgr := NewBlockManager(10, 16)
a, _ := NewSequence(mgr, 32) // 2 blocks
b, _ := NewSequence(mgr, 32) // 2 blocks
c, _ := NewSequence(mgr, 32) // 2 blocks
d, _ := NewSequence(mgr, 32) // 2 blocks → 8 used, 2 free
mgr.Free(b.BlockTable) // → 4 free, but scattered
mgr.Free(d.BlockTable)
e, err := NewSequence(mgr, 64) // needs 4 blocks — SUCCEEDS despite scattering
fmt.Println("e's blocks:", e.BlockTable, err)
fmt.Println("token 40 of e lives at physical slot", e.PhysicalIndex(40))
_, _ = a, c
}Compare PhysicalIndex to an OS page table walk. It’s the same operation:
(virtual_page → physical_frame) × page_size + offset. That’s not an analogy; it’s the same
algorithm.
Demonstrate the fragmentation fix:
// From main() above:
mgr := NewBlockManager(10, 16)
a, _ := NewSequence(mgr, 32) // 2 blocks
b, _ := NewSequence(mgr, 32) // 2 blocks
c, _ := NewSequence(mgr, 32) // 2 blocks
d, _ := NewSequence(mgr, 32) // 2 blocks → 8 used, 2 free
mgr.Free(b.BlockTable) // → 4 free, but scattered
mgr.Free(d.BlockTable)
e, _ := NewSequence(mgr, 64) // needs 4 blocks — SUCCEEDS despite scattering
fmt.Println("e's blocks:", e.BlockTable)A contiguous allocator would fail here. This one doesn’t.
5. Technical explanation#
Block size — the one real tuning parameter#
Small blocks (8 tokens):
✓ less internal fragmentation (avg waste = block_size/2 per sequence)
✗ longer block tables → more metadata, more indirection in the kernel
✗ finer-grained allocation → more scheduler work
Large blocks (128 tokens):
✓ short block tables, better kernel memory access
✗ more waste (avg 64 tokens wasted per sequence)
Standard: 16 tokens. vLLM's default. For an 8B model that's 16 × 128 KiB = 2 MB per block.Internal fragmentation with block size B: on average B/2 tokens wasted per sequence, so with
B=16 and 64 concurrent sequences, ~512 tokens wasted total — negligible versus the 60-80% the
contiguous scheme wasted.
Copy-on-write sharing#
The refcount enables genuine sharing:
Parallel sampling (n=4 completions from one prompt):
prompt blocks: [3, 7, 1] refcount 4 ← ONE copy of the prompt KV
seq 0 extra: [9]
seq 1 extra: [12]
seq 2 extra: [4]
seq 3 extra: [8]
Memory: 3 + 4 blocks instead of 4 × (3 + 1) = 16 blocks. 56% saved.For beam search with 4 beams and a 2000-token prompt, the saving is larger still.
Copy-on-write: when a shared block needs to be written (because two sequences diverge
within a block), copy it first, then write. Same mechanism as fork() in an OS.
The kernel side#
The attention kernel now takes a block table:
// Simplified
__global__ void paged_attention_kernel(
const float* q, // [num_seqs, num_heads, head_dim]
const float* k_cache, // [num_blocks, block_size, num_kv_heads, head_dim]
const float* v_cache,
const int* block_tables, // [num_seqs, max_blocks_per_seq]
const int* context_lens, // [num_seqs]
float* out, ...)
{
int seq = blockIdx.x, head = blockIdx.y;
int ctx = context_lens[seq];
const int* table = block_tables + seq * max_blocks;
for (int blk = 0; blk < ceil_div(ctx, BLOCK_SIZE); blk++) {
int phys = table[blk]; // ← the indirection
const float* k_blk = k_cache + phys * BLOCK_SIZE * num_kv_heads * head_dim;
// ... compute q·k for this block, online softmax accumulate ...
}
}Cost of the indirection: one extra load per block (amortized over 16 tokens × head_dim values) and slightly worse coalescing at block boundaries. Measured: 2-8% slower kernel. In exchange for 2.5-4x more concurrency. An easy trade.
The layout, more precisely#
vLLM’s cache shapes:
key_cache: [num_blocks, num_kv_heads, head_dim/x, block_size, x] where x = 16/sizeof(dtype)
value_cache: [num_blocks, num_kv_heads, head_dim, block_size]The odd key layout exists so that the kernel’s per-thread reads are 16-byte vectorized and
coalesced across the warp (Section IV.07). K and V differ because the kernel accesses them in
different orders: K is reduced over head_dim (for q·k), V over block_size (for p·v).
This is a good example of layout being chosen by the kernel’s access pattern, not by elegance.
The relationship to prefix caching#
Once KV is in shareable blocks with refcounts, prefix caching (file 11) is nearly free to add:
Hash the token ids of each block.
Keep a hash → block_id map.
When a new sequence's first N blocks hash to existing blocks, just point at them
(increment refcount) instead of computing them.PagedAttention is the enabling mechanism for prefix caching. That’s a large part of why it matters beyond fragmentation.
6. Under the hood#
Memory accounting in a real deployment:
vLLM startup on an 80 GB H100 with an 8B model:
gpu_memory_utilization = 0.90 → 72 GB usable
model weights 16.1 GB
activation peak (profiled) 2.4 GB
CUDA graphs 1.2 GB
─────────────────────────────────────
KV cache pool 52.3 GB
block size 16, 128 KiB/token → 2 MiB/block
→ 26,777 blocks = 428,432 tokens of KV capacityvLLM prints this at startup ("# GPU blocks: 26777"). Read that line every time you start a server — it’s your capacity, stated directly. Divide by your average context length for concurrency.
7. Performance implications#
Metric Contiguous Paged Change
KV memory utilization 20-38% 96%+ 2.5-4x
Max concurrent sequences baseline 2.5-4x ↑↑
Attention kernel speed baseline 0.92-0.98x ↓ slightly
End-to-end throughput baseline 2-4x ↑↑
Memory allocation failures common rare ↑↑
Sharing (parallel sampling) none 55%+ saved ↑↑The vLLM paper reports 2-4x throughput over Orca at the same latency, and up to 24x over HuggingFace’s naive pipeline (which combines paging with continuous batching).
8. Production implications#
- Use an engine with paged KV. vLLM, SGLang, TGI, TensorRT-LLM all have it now.
- Read the block count at startup. It’s your capacity number.
- Tune
gpu_memory_utilizationcarefully. 0.90 is the default; 0.95 gives more KV but risks OOM from activation spikes on long prefills. Test with your longest realistic prompt. - Monitor block-pool utilization, not GPU memory. GPU memory is always ~full (the pool is preallocated); block utilization is the real signal.
- Block size rarely needs tuning. 16 is right for almost everything.
- Paging enables prefix caching — turn it on (file 11).
9. Common mistakes#
Thinking PagedAttention is a faster attention kernel. It’s slightly slower per kernel. The win is memory management.
Setting gpu_memory_utilization to 0.98. OOM on an unlucky long prefill.
Monitoring GPU memory instead of block utilization. The pool is preallocated; GPU memory tells you nothing about load.
Setting max_model_len unnecessarily high. It affects the block table size per sequence and
the memory profiling, and can prevent startup.
Assuming blocks are freed when a client disconnects. Only if abort is wired up (file 04).
10. Hands-on exercise#
A. Build it. Implement the BlockManager and Sequence classes from section 4, then extend
to support: copy-on-write forking, a prefix hash table, and eviction. This is Project 09.
B. Demonstrate fragmentation. Write a simulator comparing a contiguous allocator and a paged allocator under a realistic request stream (lognormal lengths, Poisson arrivals). Measure: peak memory utilization, allocation failures, and max concurrency. Reproduce the 20-38% vs 96% finding.
C. Block size study. In your simulator, sweep block size over {1, 8, 16, 32, 64, 128}. Plot internal fragmentation and block-table length. Where’s the optimum, and does it match vLLM’s choice?
D. Read the real thing. Read vllm/core/block_manager_v2.py (or the current equivalent) and
csrc/attention/attention_kernels.cu. Find: the block table construction, the refcounting, and
the indirection in the kernel.
E. Sharing. Implement parallel sampling (n=4) with and without block sharing. Measure the memory difference for a 2000-token prompt.
11. Interview questions#
- What problem does PagedAttention solve? Quantify it.
- Explain the analogy to OS virtual memory. Which concept maps to which?
- What are internal and external fragmentation in a KV cache? Give an example of each.
- How does the attention kernel change, and what does the indirection cost?
- How does copy-on-write sharing work, and which features use it?
- How would you choose the block size?
- Why is PagedAttention a prerequisite for efficient prefix caching?
- Your engine reports 26,777 GPU blocks. What does that tell you about capacity?
12. Further reading#
- [ESTABLISHED] Kwon et al., “Efficient Memory Management for Large Language Model Serving with PagedAttention” (SOSP 2023) — read the whole paper
- [REFERENCE] vLLM
vllm/core/block_manager*.pyandcsrc/attention/ - [FUNDAMENTAL] Operating Systems: Three Easy Pieces, paging chapters — the ancestor
- Next: 11 — Prefix and prompt caching