PagedAttention: Virtual Memory Guide for LLM Inference
The Problem: The High Cost of Memory Fragmentation in LLM Serving
In high-throughput Large Language Model (LLM) serving, GPU High Bandwidth Memory (HBM) is the most scarce and expensive resource. While static model weights occupy a fixed footprint (e.g., ~140 GB for an FP16 70B parameter checkpoint), the remaining VRAM must accommodate the dynamically expanding Key-Value (KV) cache across concurrent requests.
In traditional serving systems (such as early Hugging Face pipelines, Megatron-LM, and FasterTransformer), the KV cache for every incoming request was allocated as a contiguous chunk of physical GPU memory. Because autoregressive token generation has an unpredictable sequence length, serving systems were forced to pre-allocate memory buffers based on the maximum possible sequence length (e.g., 2,048 or 8,192 tokens).
This contiguous allocation strategy causes catastrophic memory waste across three distinct dimensions:
1. Internal Fragmentation (Over-Allocation)
Because the system cannot predict whether a user prompt will trigger a 50-token response or a 2,000-token response, it reserves memory for the maximum capacity ($L_{\max}$). If a query finishes in 200 tokens, the remaining 90% of the allocated buffer sits completely idle—unusable by any other concurrent request. In real-world enterprise workloads, internal fragmentation wastes between 60% and 80% of total GPU memory.
2. Reservation Waste (The Lookahead Penalty)
Even when memory is dynamically resized in geometric increments, servers must hold reservations for tokens that have not yet been generated. If a batch of 16 requests is currently at token step 100, the system must hold buffer capacity for step 101 through 200, stranding gigabytes of physical memory while awaiting future decoding cycles.
3. External Fragmentation (Memory Holes)
Because concurrent requests arrive and depart at different times, dynamic memory allocators (like CUDA’s cudaMalloc and PyTorch’s caching allocator) create fragmented holes across physical memory. Even if 20 GB of total free VRAM exists, a new long-context request requiring an 8 GB contiguous buffer will fail with an Out-of-Memory (OOM) error if no single contiguous block of 8 GB is available.
The consequence: serving engines historically ran out of memory when only 20% to 40% of physical GPU memory was actually storing active token activations, severely constraining server batch capacity and dramatically inflating the cost per token.
The Solution: Operating System Virtual Memory Meets Transformer Attention
Introduced in the landmark vLLM research (Kwon et al., SOSP 2023), PagedAttention resolves memory fragmentation by applying the classic principles of Operating System Virtual Memory with Paging directly to the transformer KV cache.
Instead of demanding large, contiguous physical memory spans, PagedAttention decouples the logical representation of a sequence’s tokens from its physical storage in GPU HBM:
- Logical Token Blocks: A sequence’s Key-Value activations are partitioned into fixed-size logical blocks containing a constant number of tokens (typically $B = 16$ or $B = 32$).
- Physical Memory Frames: GPU HBM is divided into uniform physical frames of identical size. These frames do not need to be contiguous.
- The Block Table (Page Table): A lightweight management table maintained on the host CPU and GPU that maps each logical block index of a request to its corresponding physical frame address.
The Mathematics of Block Translation
Let $B$ denote the block size (number of tokens per block). For a sequence currently containing $t$ tokens, token index $i \in [0, t-1]$ is mapped to its logical block index $b_{\text{logical}}$ and intra-block offset $o$ via integer arithmetic:
$$b_{\text{logical}} = \lfloor i / B \rfloor, \quad o = i \pmod B$$
The serving engine inspects the request’s Block Table to retrieve the physical frame address:
$$\text{Address}(K_i) = \text{BlockTable}[b_{\text{logical}}] \times (\text{Frame Size}) + o \times (2 \times n_{\text{heads}} \times d_{\text{head}} \times \text{bytes})$$
Because physical blocks are allocated strictly on demand, internal fragmentation is confined exclusively to the final partial block of a sequence. The maximum memory wasted per request is strictly bounded by:
$$\text{Maximum Waste per Request} < B \times \text{Token Cache Size}$$
For a standard configuration with $B = 16$, total memory waste drops from over 70% to under 4%, recovering up to 96% of previously stranded VRAM.
Key Architectural Mechanics
PagedAttention powers several critical runtime optimizations that are impossible under contiguous allocation schemes:
1. Demand Paging during Autoregressive Decoding
When a new request arrives, vLLM allocates only the exact number of physical frames required to store the initial prompt. During subsequent token generation:
- The model generates token $t$.
- The engine calculates the logical block index $b = \lfloor t / B \rfloor$.
- If $t \pmod B \ne 0$, the new token is written into the existing physical frame at offset $t \pmod B$.
- When $t \pmod B = 0$ (the current block is full), the memory manager pops a single free frame from the pre-allocated physical frame pool and appends the new mapping to the Block Table.
No memory reallocation, memory moving, or tensor reallocation occurs.
2. Zero-Copy Branching via Copy-on-Write (CoW)
Advanced generative workflows—such as parallel sampling ($n > 1$ completions from a single prompt), beam search, and Monte Carlo tree search (MCTS)—traditionally duplicated the prompt’s KV cache across all candidate branches.
With PagedAttention, parallel branches achieve true zero-copy sharing:
- When a user requests 4 candidate completions for a 1,000-token prompt, the prompt’s logical blocks point to the exact same physical frames.
- Each physical frame maintains a
ref_counttracking how many active requests point to it. - When branch $A$ generates divergent tokens that spill into a new block, the memory manager allocates a new physical frame only for that divergent block.
- If an existing shared block must be modified, the engine triggers Copy-on-Write (CoW): duplicating only that single 16-token frame while preserving shared references to all prior prompt frames.
This cuts memory usage during parallel sampling and reasoning rollouts by up to 55%, enabling massive multi-candidate generation without memory exhaustion.
3. PagedAttention CUDA Kernel Execution
Standard attention kernels (like classic FlashAttention) expect contiguous Key and Value tensors formatted as (Batch, SeqLen, Heads, HeadDim). Feeding non-contiguous memory blocks into traditional kernels would require gathering blocks into a temporary buffer—incurring prohibitive memory bandwidth overhead.
The PagedAttention CUDA kernel solves this by accepting the Block Table directly as an input parameter:
// Conceptual signature of the PagedAttention CUDA kernel
__global__ void paged_attention_kernel(
const scalar_t* __restrict__ q, // [num_seqs, num_heads, head_dim]
const scalar_t* __restrict__ k_cache, // [num_blocks, num_heads, head_dim/x, block_size, x]
const scalar_t* __restrict__ v_cache, // [num_blocks, num_heads, head_dim, block_size]
const int* __restrict__ block_tables, // [num_seqs, max_num_blocks_per_seq]
const int* __restrict__ seq_lens, // [num_seqs]
scalar_t* __restrict__ out, // [num_seqs, num_heads, head_dim]
const int block_size
); During execution, each warp of threads queries the block_tables array to determine the physical memory pointers for the required blocks, loading them directly into on-chip GPU SRAM. Computation and address resolution occur seamlessly in a single fused kernel.
Memory Management Comparison Matrix
| Dimension | Static Contiguous Allocation | Dynamic Reallocation | PagedAttention (vLLM) |
|---|---|---|---|
| Physical Storage | Contiguous virtual & physical | Contiguous virtual & physical | Non-contiguous physical frames |
| Allocation Timing | Upfront for $L_{\max}$ | Reallocates on capacity breach | Demand paging per block ($B=16$) |
| Average Memory Waste | 60% – 80% (Internal fragmentation) | 25% – 40% (Overhead & copies) | < 4% (Confined to last block) |
| Memory Defragmentation | Impossible without dropping requests | Expensive host-device sync | Zero defragmentation needed |
| Parallel Sampling (n=4) | 4x prompt KV cache duplication | 4x prompt KV cache duplication | Zero-copy shared prompt blocks |
| Effective Batch Capacity | Baseline (1.0x) | 1.3x – 1.6x | 2.0x – 4.0x higher throughput |
| Prefix Caching Compatibility | Fragile hash matching | Fragile hash matching | Native chunk-level prefix reuse |
Reference Implementation: Block Table & Paged Cache Manager
Below is a complete, educational Python reference implementation demonstrating how a PagedAttention memory manager tracks physical frames, dynamically maps logical blocks, and implements Copy-on-Write for parallel generation.
from typing import List, Dict, Optional
import math
class PhysicalBlockPool:
"""
Manages the allocation, freeing, and reference counting of GPU physical blocks.
"""
def __init__(self, num_blocks: int, block_size: int = 16):
self.num_blocks = num_blocks
self.block_size = block_size
self.free_blocks: List[int] = list(range(num_blocks))
self.ref_counts: Dict[int, int] = {i: 0 for i in range(num_blocks)}
def allocate(self) -> int:
if not self.free_blocks:
raise MemoryError("Out of GPU Physical Blocks (VRAM Full)!")
block_id = self.free_blocks.pop(0)
self.ref_counts[block_id] = 1
return block_id
def add_reference(self, block_id: int):
self.ref_counts[block_id] += 1
def free(self, block_id: int):
if self.ref_counts[block_id] <= 0:
return
self.ref_counts[block_id] -= 1
if self.ref_counts[block_id] == 0:
self.free_blocks.append(block_id)
@property
def free_count(self) -> int:
return len(self.free_blocks)
class SequenceGroupMemoryManager:
"""
Tracks logical-to-physical block tables across requests and parallel forks.
"""
def __init__(self, pool: PhysicalBlockPool, block_size: int = 16):
self.pool = pool
self.block_size = block_size
# Maps request_id -> list of physical_block_ids
self.block_tables: Dict[str, List[int]] = {}
# Maps request_id -> sequence length in tokens
self.seq_lens: Dict[str, int] = {}
def init_sequence(self, request_id: str, prompt_tokens: int):
"""Allocates initial physical blocks for a new prompt."""
num_blocks = math.ceil(prompt_tokens / self.block_size)
allocated = [self.pool.allocate() for _ in range(num_blocks)]
self.block_tables[request_id] = allocated
self.seq_lens[request_id] = prompt_tokens
def append_token(self, request_id: str):
"""
Appends a newly generated token, allocating a new physical block only when needed.
Handles Copy-on-Write (CoW) if the current tail block is shared across branches.
"""
cur_len = self.seq_lens[request_id]
table = self.block_tables[request_id]
cur_block_id = table[-1]
# Check if we are crossing into a brand new block
if cur_len % self.block_size == 0:
new_block_id = self.pool.allocate()
table.append(new_block_id)
else:
# Check Copy-on-Write: If the block is shared by another branch, copy it!
if self.pool.ref_counts[cur_block_id] > 1:
new_block_id = self.pool.allocate()
# Emulate copying memory: decrease old ref, assign new unique block
self.pool.free(cur_block_id)
table[-1] = new_block_id
self.seq_lens[request_id] = cur_len + 1
def fork_sequence(self, parent_id: str, child_id: str):
"""
Zero-copy branching: Fork a sequence for parallel sampling.
Increments reference count on all existing blocks without duplicating memory!
"""
parent_table = self.block_tables[parent_id]
for b_id in parent_table:
self.pool.add_reference(b_id)
self.block_tables[child_id] = list(parent_table)
self.seq_lens[child_id] = self.seq_lens[parent_id]
def free_sequence(self, request_id: str):
"""Frees all blocks associated with a completed request."""
if request_id in self.block_tables:
for b_id in self.block_tables[request_id]:
self.pool.free(b_id)
del self.block_tables[request_id]
del self.seq_lens[request_id]
# Demonstration of Zero-Copy Forking and Copy-on-Write
if __name__ == "__main__":
pool = PhysicalBlockPool(num_blocks=100, block_size=16)
manager = SequenceGroupMemoryManager(pool=pool, block_size=16)
# 1. Start Prompt of 32 tokens (occupies exactly 2 blocks: Frame 0, Frame 1)
manager.init_sequence(request_id="req_alpha", prompt_tokens=32)
print(f"req_alpha Block Table: {manager.block_tables['req_alpha']}")
print(f"Free blocks remaining: {pool.free_count} (Allocated: 2)")
# 2. Fork into two parallel branches (zero-copy memory sharing)
manager.fork_sequence(parent_id="req_alpha", child_id="req_beta")
print(f"req_beta Block Table: {manager.block_tables['req_beta']}")
print(f"Free blocks after fork: {pool.free_count} (Zero extra frames used!)")
print(f"Frame 0 ref count: {pool.ref_counts[0]}, Frame 1 ref count: {pool.ref_counts[1]}")
# 3. Append token to req_alpha (triggers allocation of block 3 because length was 32)
manager.append_token(request_id="req_alpha")
print(f"req_alpha Block Table after 1 token: {manager.block_tables['req_alpha']}")
print(f"req_beta Block Table unchanged: {manager.block_tables['req_beta']}")
print(f"Free blocks remaining: {pool.free_count}")
# 4. Clean up
manager.free_sequence("req_alpha")
manager.free_sequence("req_beta")
print(f"All requests completed. Free blocks: {pool.free_count} / 100") Production Deployment and Ecosystem Integration
PagedAttention has graduated from an experimental optimization into the foundational memory substrate of modern AI serving infrastructure:
1. Sizing the Block Dimension ($B=16$ vs. $B=32$)
When configuring serving engines like vLLM and SGLang, choosing block size $B$ represents a delicate hardware trade-off:
- Smaller Blocks ($B=16$): Minimizes tail memory waste to $<16$ tokens per sequence. Ideal for workloads dominated by short conversational turns or heavy parallel branching.
- Larger Blocks ($B=32$ or $B=64$): Improves memory coalescing and GPU memory bandwidth utilization during CUDA kernel memory reads. Recommended for uniform long-context document synthesis ($>16\text{k}$ tokens) on NVIDIA Hopper (H100/H200) and Blackwell GPUs.
2. Native Synergy with RadixAttention and Prompt Caching
Modern inference systems layer prefix-caching algorithms (such as SGLang’s RadixAttention) directly on top of PagedAttention. Because blocks are immutable once filled, identical multi-turn chat prefixes or static system prompts map directly to cached physical frames, enabling instant multi-turn retrieval with zero recomputation.
3. Local LLM Acceleration in aiFetchly
For engineering teams utilizing AiFetchly to orchestrate local open-weight models via Ollama or vLLM backends, PagedAttention provides a critical operational advantage. By running local models with paged KV allocation, desktop workstations can comfortably serve 32k-context research extraction loops and multi-agent web intelligence sweeps without triggering desktop VRAM thrashing or system crashes.