vLLM — PagedAttention

Overview

vLLM introduces PagedAttention, an attention algorithm inspired by operating-system virtual memory and paging, that fundamentally changes how LLM serving systems manage the key-value (KV) cache. By dividing the KV cache into fixed-size blocks that need not be contiguous in physical GPU memory, PagedAttention achieves near-zero memory waste and enables flexible memory sharing across sequences—delivering 2–4× throughput improvements over prior systems such as FasterTransformer and Orca.

Figure 1 (from paper): Left: Memory layout when serving an LLM with 13B parameters on NVIDIA A100. The parameters (gray) persist in GPU memory throughout serving. The KV cache memory (red) is (de)allocated per serving request. Right: vLLM smooths out the rapid growth curve of KV cache memory seen in existing systems, leading to a notable boost in serving throughput.

The Memory Problem

In autoregressive LLM serving, each request requires storing key and value tensors for every previously generated token—the KV cache. For a 13B-parameter model (OPT-13B), a single token's KV cache occupies 800 KB (2 × 5120 × 40 × 2 bytes), and a full 2048-token sequence can require up to 1.6 GB. On an A100 GPU with 40GB RAM, roughly 65% goes to model parameters, leaving only ~30% for KV cache—the critical bottleneck.

Figure 2 (from paper): Average percentage of memory wastes in different LLM serving systems during the basic sampling experiment.

Figure 3 (from paper): KV cache memory management in existing systems. Three types of memory wastes -- reserved, internal fragmentation, and external fragmentation -- prevent other requests from fitting into the memory.

Three Sources of Memory Waste

Measured Memory Waste (OPT-13B on ShareGPT)

Orca (Max): only 20.4% of KV cache memory stores actual token states. vLLM achieves 96.3%.

System Architecture

vLLM adopts a centralized scheduler with distributed GPU workers. The KV Cache Manager manages memory via block tables, analogous to how an OS kernel manages page tables. GPU workers only see physical block IDs—they never need to reason about memory sharing or fragmentation.

Figure 4 (from paper): vLLM system overview with centralized scheduler and distributed GPU workers connected through a KV cache manager.

PagedAttention Algorithm

Traditional attention requires keys and values stored contiguously in memory. PagedAttention partitions the KV cache into blocks, each holding B tokens' keys and values. During attention computation, the kernel fetches blocks from non-contiguous physical locations via the block table, computing attention block-by-block.

Mathematical Formulation

For a sequence of length i, denote the key block Kj = (k(j-1)B+1, ..., kjB) and value block Vj similarly. The attention computation becomes block-wise:

// Block-wise attention computation (Equation 4 from the paper)
// B = block size (default 16 tokens per block)

// For each key block j = 1 ... ceil(i/B):
Aij = exp(qiT Kj / sqrt(d)) / sum(exp(qiT Kt 1 / sqrt(d)), t=1..ceil(i/B))

// Final output is weighted sum over value blocks:
oi = sum(Vj AijT, j=1..ceil(i/B))

// Key difference: kernel identifies and fetches DIFFERENT KV blocks
// from non-contiguous physical memory using the block table.
// Each GPU warp processes one block for coalesced memory access.

Key Mechanisms

1. Block-Level Memory Management

2. Copy-on-Write (CoW)

3. FCFS Scheduling with Preemption

4. Application to Decoding Scenarios

Code Concepts

vLLM is implemented in ~8.5K lines of Python and ~2K lines of C++/CUDA. The block manager and scheduler are in Python; the PagedAttention kernel and block operations are fused CUDA kernels.

Block Table Data Structure

# Conceptual block table structure (per sequence)
# Maps logical block index -> (physical_block_id, num_filled, ref_count)

class BlockTable:
    """Maps logical KV blocks to physical blocks in GPU DRAM."""

    def __init__(self, block_size: int = 16):
        self.block_size = block_size  # tokens per block (B)
        self.entries: List[BlockEntry] = []

    def append_token(self, allocator: BlockAllocator):
        """Called for each new token generated."""
        if not self.entries or self.entries[-1].is_full():
            # Last block is full -> allocate a NEW physical block
            phys_block = allocator.allocate()
            self.entries.append(BlockEntry(phys_block, num_filled=1))
        else:
            # Room in current block -> just increment counter
            self.entries[-1].num_filled += 1

    def get_physical_ids(self) -> List[int]:
        """Returns physical block IDs for the attention kernel."""
        return [e.physical_block_id for e in self.entries]

Copy-on-Write Logic

# When a sequence writes to a block shared by others (ref_count > 1)

def cow_write(block_entry: BlockEntry, allocator: BlockAllocator):
    """Copy-on-write: fork a new physical block if shared."""
    if block_entry.ref_count > 1:
        # Shared block -> copy to new physical block
        old_phys = block_entry.physical_block_id
        new_phys = allocator.allocate()
        cuda_block_copy(src=old_phys, dst=new_phys)  # fused kernel
        block_entry.physical_block_id = new_phys
        block_entry.ref_count = 1
        allocator.dec_ref(old_phys)  # decrement old block's ref
    # ref_count == 1 -> write directly, no copy needed

Fused CUDA Kernels

// Conceptual structure of the PagedAttention CUDA kernel
// Three fused operations minimize kernel launch overhead:

// 1. Fused Reshape & Block Write
//    New KV cache is split into blocks, reshaped for coalesced reads,
//    and saved at positions specified by the block table.

// 2. Fused Block Read & Attention
//    Each GPU warp reads one KV block from the block table,
//    computes partial attention scores, then reduces across warps.
__global__ void paged_attention_kernel(
    const float* query,           // [num_heads, head_dim]
    const float* key_cache,       // [num_blocks, block_size, num_heads, head_dim]
    const float* value_cache,     // [num_blocks, block_size, num_heads, head_dim]
    const int*   block_tables,    // [num_seqs, max_blocks_per_seq]
    const int*   seq_lens,        // [num_seqs]
    float*       output           // [num_seqs, num_heads, head_dim]
) {
    // Each warp handles one KV block for one attention head
    int block_idx = block_tables[seq_id * max_blocks + logical_block];
    // Fetch K, V from physical block_idx (non-contiguous!)
    // Compute softmax(Q * K^T / sqrt(d)) * V per block
    // Reduce partial results across blocks
}

// 3. Fused Block Copy (for CoW)
//    Batches all pending copy operations into a single kernel launch.

Three Core API Methods

# vLLM supports all decoding via three sequence operations:

fork(seq)    # Create new sequence sharing parent's KV blocks
              # Used for: parallel sampling, beam search branching

append(seq)  # Add new token to sequence (may allocate block or trigger CoW)
              # Used for: every autoregressive decode step

free(seq)    # Release sequence's blocks (decrement ref counts)
              # Used for: completed sequences, pruned beam candidates

Performance Results

Throughput Improvement (Basic Sampling)

Memory Utilization

Memory Savings from Sharing

Batch Size Improvement

Kernel Overhead

Block Size Selection

Limitations for Agentic AI Systems

While PagedAttention is foundational, its design targets single-request serving within a single session. Modern agentic AI systems expose several gaps:

Impact & Legacy

PagedAttention is arguably the single most influential systems contribution to LLM serving. Its impact extends across three dimensions:

Why PagedAttention is Foundational

Key Numbers at a Glance

Repository & Resources

Implementation Summary