PagedAttention: fragmentation, blocks, block tables

Take the reference machine from §0.4: one H100 SXM, 80 GB. Serve Llama-3-8B in bf16 with --max-model-len 32768. §2.1 did this budget in full — 79.65 GiB of HBM, times the default gpu_memory_utilization of 0.92, minus 14.96 GiB of bf16 weights, minus a 6.0 GiB allowance for activations, CUDA graph pools (§8.1) and workspace — and landed on a 52.32 GiB KV pool. Take that as given here. One token of Llama-3-8B KV costs

where $L$ is layer count, $h_{kv}$ the number of KV heads, $d_h$ the head dimension and $b$ the bytes per element. Carved into block_size=16 blocks, 52.32 GiB is 26,785 blocks — 428,560 token-slots (§2.1 works the two floor divisions). That is a lot of tokens. The question is how many requests it is.

For the fixed maximum-length contiguous-reservation baseline, each sequence gets max_model_len slots at admission. At 32768 tokens and 128 KiB/token this is 4 GiB, so 13 such reservations fit a 52.32 GiB pool. This is one allocator design, not a necessity for every contiguous virtual-address layout. Use a context-capable Llama-3.1 checkpoint for the long-context commands below; the original Llama-3 checkpoint's native limit is 8192.

Now count what those thirteen sequences actually use. A typical chat turn is a 512-token prompt and 100 generated tokens: 612 tokens, 76.5 MiB. Thirteen of them is 0.97 GiB of live KV inside a 52.32 GiB pool.

Count both prompt and output state: a 512-token prompt plus 100 outputs uses 612 of 32768 reserved slots, so the waste is 1-612/32768=98.13%. Counting only the 100 output tokens would incorrectly omit the prompt KV.

The waste has three distinct sources and they behave differently, so split them before fixing them. All three numbers below are derived arithmetic on the setup above, not measurements.

Kwon et al. report that existing serving systems waste 60–80% of KV cache memory to these effects. That figure is theirs, not mine — the systems compared and the methodology are in §1 and Figure 2 of arXiv:2309.06180, and I have not re-measured it. The derivation above only shows how a number that large is arithmetically possible.

The fix is the one operating systems landed on in the 1960s: stop allocating a contiguous extent and start allocating a list of fixed-size pieces. A sequence no longer owns a range of memory; it owns a list of block indices. Logical position 4,097 in the sequence and physical position 4,097 in the pool have nothing to do with each other, and the mapping between them is an array lookup.

That kills two of the three wastes outright — over-reservation, because a block is allocated only when the previous one fills, and external fragmentation, because every free unit is the same size and therefore satisfies any request — and introduces the third, the unwritten tail of the last block, in a bounded form.

686 is a memory ceiling, not a throughput promise — --max-num-seqs and the compute roofline will bind first. The point is that memory stops being the binding constraint, which is the whole game.

The mapping to virtual memory is exact enough to be useful: a sequence's token positions are virtual addresses, physical blocks in the pool are page frames, the per-sequence block table is a page table, and the block pool's free list is the kernel's free-frame list. Keep that. But four things are different, and every one of them is a design decision someone made deliberately.

Blocks are appended in order, never inserted randomly. A page table takes arbitrary VA→PA insertions in any order; a block table only ever grows at the tail, because a sequence only ever extends forward. vLLM leans on this hard enough to write it into a comment justifying a correctness shortcut:

    NOTE #1: We currently don't de-duplicate the blocks in the cache,
    meaning that if a block becomes full and is cached, we don't check
    if there is already an identical block in the cache. This is because
    we want to make sure the allocated block IDs won't change so that
    block tables are append-only.

Because the table is append-only, the worker-side update is a single strided numpy write with no reshuffling and no invalidation, and the device-side row can be memcpy'd wholesale. An OS would need TLB shootdowns for the equivalent.

There is no demand paging. No fault handler exists. The engine never over-commits and then resolves the fault on access — it checks capacity before the step and, if the check fails, does not run the request at all. allocate_slots returns None and the scheduler preempts somebody; the policy for who is §1.4's. The consequence is that a KV block is never absent when the kernel reaches for it, which is why the attention kernel can be a straight gather with no fallback path.

The page size is chosen for kernel efficiency, not for an MMU. An OS page size is fixed by hardware. Here it is a free parameter, and the two engines pick opposite ends of the range — vLLM defaults to 16 tokens, SGLang to 1. Neither number has anything to do with address translation; both are about how the attention backend wants to read memory. §2.2.7 below.

Sharing is content-addressed at prefix granularity. Two processes share a frame because the OS mapped the same physical frame into both address spaces — an explicit act. Two sequences share a block because their token prefixes happen to be identical, discovered by hashing (§2.3). The allocator never decides to share; it only maintains the reference counts that make sharing safe.

Underneath all four: there is no MMU. Translation is neither free nor hidden — every attention kernel invocation is handed the block table as a tensor and does the lookup itself, in software, per query. That cost is why block size is a tuning parameter at all.

Everything above is implemented by exactly three things in vLLM. Read them in this order and the rest of the file makes sense.

1. The block pool

The pool owns every block in the KV cache and one free list. A block is 40-odd bytes of Python metadata — not the KV data itself, which lives in per-layer GPU tensors indexed by block_id:

@dataclass(slots=True)
class KVCacheBlock:
    """KV-cache block metadata."""

    # Block ID, ranging from 0 to num_gpu_blocks - 1.
    block_id: int
    # Reference count.
    ref_cnt: int = 0
    # The hash key (block hash + group id) of the block, only available
    # when the block is full and cached.
    _block_hash: BlockHashWithGroupId | None = None
    # Number of prefix tokens covered by _block_hash. For full blocks this is
    # the full block boundary; partial entries can end inside a cache block.
    _block_hash_num_tokens: int | None = None

    # Used to construct a doubly linked list for free blocks.
    # These two attributes should only be manipulated by FreeKVCacheBlockQueue.
    prev_free_block: "KVCacheBlock | None" = None
    next_free_block: "KVCacheBlock | None" = None

    # Whether the block is a null block that should never be cached.
    is_null: bool = False

The free list is intrusive: prev_free_block/next_free_block live on the block itself, so FreeKVCacheBlockQueue allocates no Python objects while manipulating the list, and removing a block from the middle — which happens whenever a cached block gets a prefix hit — is O(1). The pool is constructed once, at startup, with one KVCacheBlock per physical block:

        assert isinstance(num_gpu_blocks, int) and num_gpu_blocks > 0
        self.num_gpu_blocks = num_gpu_blocks
        self.enable_caching = enable_caching
        self.hash_block_size = hash_block_size
        # All kv-cache blocks.
        self.blocks: list[KVCacheBlock] = [
            KVCacheBlock(idx) for idx in range(num_gpu_blocks)
        ]
        # Free block queue that constructs and manipulates a doubly linked
        # list of free blocks (including eviction candidates when caching is
        # enabled).
        self.free_block_queue = FreeKVCacheBlockQueue(self.blocks)

        # Cache for block lookup
        self.cached_block_hash_to_block: BlockHashToBlockMap = BlockHashToBlockMap()
        self.cached_block_hashes_by_block: dict[int, set[BlockHashWithGroupId]] = {}

        # To represent a placeholder block with block_id=0.
        # The ref_cnt of null_block is not maintained, needs special care to
        # avoid freeing it.
        self.null_block = self.free_block_queue.popleft()
        self.null_block.is_null = True

Allocation is a pop from the head, plus an eviction of whatever prefix-cache entry that block used to hold:

    def get_new_blocks(self, num_blocks: int) -> list[KVCacheBlock]:
        """Get new blocks from the free block pool.

        Note that we do not check block cache in this function.

        Args:
            num_blocks: The number of blocks to allocate.

        Returns:
            A list of new block.
        """
        if num_blocks > self.get_num_free_blocks():
            raise ValueError(f"Cannot get {num_blocks} free blocks from the pool")

        ret: list[KVCacheBlock] = self.free_block_queue.popleft_n(num_blocks)

        # In order to only iterate the list once, we duplicated code a bit
        if self.enable_caching:
            for block in ret:
                self._maybe_evict_cached_block(block)
                assert block.ref_cnt == 0
                block.ref_cnt += 1

Freeing is where the interesting policy sits, and it is the mechanism behind the finding §1.4 reported about preemption victims resuming with a near-total prefix-cache hit:

    def free_blocks(self, ordered_blocks: Iterable[KVCacheBlock]) -> None:
        """Free a list of blocks. The blocks should be ordered by their
        eviction priority, where the first block will be evicted first.

        Args:
            ordered_blocks: A list of blocks to free ordered by their eviction
                priority.
        """
        # Identify blocks with hash (LRU cache) and without it (never match APC)
        blocks_to_evict_last = []
        blocks_to_evict_first = []
        for block in ordered_blocks:
            block.ref_cnt -= 1
            if block.ref_cnt == 0 and not block.is_null:
                if block.block_hash is None or not self.enable_caching:
                    # LIFO reuse of non-cached blocks for better GPU locality.
                    blocks_to_evict_first.append(block)
                else:
                    # FIFO reuse of cached blocks for LRU eviction behavior.
                    blocks_to_evict_last.append(block)

        # Blocks to reuse first are prepended to the front of the free queue.
        self.free_block_queue.prepend_n(blocks_to_evict_first)
        # Blocks to reuse last are appended to the end of the free queue.
        self.free_block_queue.append_n(blocks_to_evict_last)

A freed block does not lose its hash. It goes to the tail of the free queue if it carries one and to the head if it does not, so hashed blocks are the last to be recycled and remain matchable in cached_block_hash_to_block the whole time. That single branch is why a preempted request can come back and hit its own blocks. It is also the reason a request's blocks are freed in reverse order — SingleTypeKVCacheManager.free calls free_blocks(reversed(...)) so the tail of the sequence, which is the least valuable prefix, is evicted first.

2. The block table

Per sequence, the logical→physical map is a plain Python list on the scheduler side — self.req_to_blocks: defaultdict[str, list[KVCacheBlock]] in SingleTypeKVCacheManager.__init__ (vllm/v1/core/single_type_kv_cache_manager.py:L91-L94). Position $p$ lives in block req_to_blocks[req_id][p // block_size] at offset p % block_size. That is the whole abstraction.

On the worker side that Python list becomes a rectangular int32 tensor with a CPU staging copy:

        self.block_table = self._make_buffer(
            self.max_num_reqs, self.max_num_blocks_per_req, dtype=torch.int32
        )
        self.num_blocks_per_row = np.zeros(max_num_reqs, dtype=np.int32)

        self.slot_mapping = self._make_buffer(
            self.max_num_batched_tokens, dtype=torch.int64
        )

Note the shape. max_num_blocks_per_req is max_model_len / block_size — the table itself is still sized for the worst case, because it must be a fixed-shape tensor for CUDA graphs. But it costs 4 bytes per logical block instead of 2 MiB of all-layer KV payload: at max_num_reqs=256 and a 32k window, 256 × 2048 × 4 = 2 MiB of metadata to index 52.32 GiB of KV. That ratio — 4 bytes of index per 2 MiB of all-layer KV data (16 tokens times 128 KiB/token) — is what makes the indirection affordable, and it is exactly the ratio block_size controls.

Appending is the append-only write promised above:

    def append_row(
        self,
        block_ids: list[int],
        row_idx: int,
    ) -> None:
        if not block_ids:
            return

        if self.use_hybrid_blocks:
            block_ids = self.map_to_kernel_blocks(
                np.array(block_ids), self.blocks_per_kv_block, self._kernel_block_arange
            )

        num_blocks = len(block_ids)
        start = self.num_blocks_per_row[row_idx]
        self.num_blocks_per_row[row_idx] += num_blocks
        self.block_table.np[row_idx, start : start + num_blocks] = block_ids

3. The slot mapping

The block table tells the kernel where to read past K/V. The slot mapping tells it where to write this step's new K/V: one int64 per scheduled token, giving a flat index into the block-major KV storage. It is recomputed every step by a Triton kernel, and the arithmetic is four lines (shown with context-parallel world size 1, where the virtual/local distinction collapses):

            block_indices = (
                virtual_block_indices * BLOCKS_PER_KV_BLOCK
                + local_block_offsets // block_size
            )
            block_numbers = tl.load(
                block_table_ptr + row_offset + block_indices,
                mask=mask & is_local,
                other=0,
            ).to(tl.int64)
            slot_offsets = local_block_offsets % block_size
            slot_ids = block_numbers * block_size + slot_offsets
            slot_ids = tl.where(is_local, slot_ids, PAD_ID)
            tl.store(slot_mapping_ptr + offsets, slot_ids, mask=mask)

Strip the parallelism and it is slot = block_table[req][pos // block_size] * block_size + pos % block_size. That the KV storage is genuinely block-major, so this index is valid, is asserted elsewhere in the codebase: "Block-major backing storage: block i owns the contiguous byte range [i * page_size, (i + 1) * page_size)" (vllm/v1/worker/utils.py:L585-L586).

KVCacheManager.allocate_slots is the single entry point, and its docstring carries an ASCII diagram of the token ranges it reasons about that is worth reading before the code:

        Blocks layout:
        ```
        ----------------------------------------------------------------------
        | < comp > | < new_comp > | < ext_comp >  | < new >  | < lookahead > |
        ----------------------------------------------------------------------
                                                  |   < to be computed >     |
        ----------------------------------------------------------------------
                                  |            < to be allocated >           |
        ----------------------------------------------------------------------
                                  | < to be cached (roughly, |
                                  | details below)>          |
        ----------------------------------------------------------------------

The capacity test is a subtraction, not a search — which is the whole payoff of uniform block size:

        # Keep `reserved_blocks` free for other in-flight sequences, and an
        # additional watermark of headroom for waiting/preempted admissions.
        available_blocks = self.block_pool.get_num_free_blocks() - reserved_blocks
        required_blocks = num_blocks_to_allocate + watermark_blocks
        if required_blocks > available_blocks:
            # Cannot allocate new blocks
            return None

Under a contiguous allocator this test would be "is there a free extent of at least N contiguous bytes", which is a first-fit or best-fit scan and can fail while N bytes are free. Here it is one integer comparison and it never lies.

Concretely: request r7 is at row 3 of the persistent batch, has 4,096 committed tokens of Llama-3-8B context, and the scheduler gives it one more token this step. block_size is 16, so its 4,096 tokens fill logical blocks 0–255 exactly — the tail block is full and the next write has nowhere to go.

1. Scheduler asks. Scheduler.schedule() reaches the running-request loop and calls self.kv_cache_manager.allocate_slots(request, num_new_tokens=1, ...) (vllm/v1/core/sched/scheduler.py:L637-L643) inside a while True that retries after each preemption.
2. Sizing. get_num_blocks_to_allocate computes num_required_blocks = cdiv(4097, 16) = 257. The request already holds num_req_blocks = 256. It is a running request, so the fast path fires: return max(257 - 256, 0) = 1 (vllm/v1/core/single_type_kv_cache_manager.py:L191-L197).
3. Capacity. 1 + watermark_blocks is compared against get_num_free_blocks() - reserved_blocks. Suppose 8,412 blocks are free; it passes. Had it failed, allocate_slots would return None and the scheduler's loop would evict somebody and come back.
4. Pop. allocate_new_blocks recomputes num_new_blocks = 257 - 256 = 1 and calls self.block_pool.get_new_blocks(1). That pops the head of free_block_queue — say block 1913 — calls _maybe_evict_cached_block on it (if 1913 still carried a hash from a previous tenant, that prefix-cache entry dies here), asserts ref_cnt == 0 and sets it to 1.
5. Append. req_blocks.extend([block 1913]). The scheduler-side row is now 257 long; req_to_blocks["r7"][256].block_id == 1913.
6. Cross the process boundary. The new ids travel as CachedRequestData.new_block_ids, documented as "For request ids not in resumed_req_ids, new_block_ids will be appended to the request's block IDs" (vllm/v1/core/sched/output.py:L131-L142).
7. Worker table. The model runner calls self.input_batch.block_table.append_row(new_block_ids, req_index) (vllm/v1/worker/gpu_model_runner.py:L1526-L1528). With row_idx = 3 and num_blocks_per_row[3] == 256, this writes block_table.np[3, 256] = 1913 and bumps the row length to 257. Nothing else in the row moves. Then commit_block_table(num_reqs) does the H2D copy.
8. Slot index. compute_slot_mapping launches the Triton kernel with this step's positions. For our token, pos = 4096, so block_indices = 4096 // 16 = 256, block_numbers = block_table[3, 256] = 1913, slot_offsets = 4096 % 16 = 0, and slot_ids = 1913 * 16 + 0 = 30,608.
9. Write. Every layer's KV-cache-write op stores this token's K and V at flat row 30,608 of its block-major cache tensor. All 32 layers use the same slot index; only the tensor differs.

Internal fragmentation is bounded by block_size − 1 tokens per sequence per KV cache group. For Llama-3-8B at 128 KiB/token that is up to 1.875 MiB per sequence at block_size=16, 15.875 MiB at 128. At 686 concurrent sequences those are 1.3 GiB and 10.6 GiB of the 52.32 GiB pool respectively — 2.4% versus 20%. Small blocks are strictly better on memory.

They are not better on everything else. Larger blocks mean a narrower block table (fewer int32 loads per query in the attention kernel), longer contiguous runs for the KV gather, fewer entries to hash for prefix caching, and fewer free-list operations per step. Some backends will not accept small blocks at all — FlashAttentionBackend.get_preferred_block_size raises the floor to 64 on XPU (vllm/v1/attention/backends/flash_attn.py:L96-L99), and the platform layer lets any backend override the default before the pool is sized.

vLLM's default is a constant with a backend override hook:

class CacheConfig:
    """Configuration for the KV cache."""

    DEFAULT_BLOCK_SIZE: ClassVar[int] = 16

    block_size: int = Field(default=None, gt=0)  # type: ignore[assignment]
    """Size of a contiguous cache block in number of tokens.
    Accepts None (meaning "use default"). After construction, always int."""

SGLang's is a resolution pass that returns 1 on everything except MUSA and a specific ROCm KV layout:

    if is_hip() and envs.SGLANG_AITER_KV_CACHE_LAYOUT.get().lower() == "vectorized_5d":
        logger.info(
            "Setting page_size=64 as default for "
            "SGLANG_AITER_KV_CACHE_LAYOUT=vectorized_5d."
        )
        return {"page_size": 64}
    if not is_musa():
        return {"page_size": 1}
    return {"page_size": 64}

vLLM decoupled the two pressures rather than picking a compromise. get_block_table_width takes a separate kernel_block_size and splits each allocation block into several kernel blocks, so a 32-token allocation block can feed a kernel that wants 16-token pages (vllm/v1/worker/block_table.py:L29-L49, map_to_kernel_blocks at L238-L266). And prefix_match_unit — the hash_block_size threaded through BlockPool — can be finer than block_size, so prefix-cache hits can land inside a physical block. That last decoupling is what creates the need for copy-on-write.

The trigger is one modulo:

    def _has_partial_local_hit(
        self,
        new_computed_blocks: Sequence[KVCacheBlock],
        num_local_computed_tokens: int,
    ) -> bool:
        # The local prefix-cache hit ends inside one of this manager's
        # blocks: the shared tail block needs CoW.
        return (
            len(new_computed_blocks) > 0
            and num_local_computed_tokens % self.block_size != 0
        )

If a hit covers 4,100 tokens with block_size=16, logical block 256 is shared but only 4 of its 16 slots are valid for this request. The request cannot write into slots 4–15 of a block another request also references. add_local_computed_blocks records the offending index — block_idx = num_local_computed_tokens // self.block_size, with the shared tail block, at vllm/v1/core/single_type_kv_cache_manager.py:L280-L286 — and allocate_new_blocks then pulls one extra block and redirects:

        cow_blocks: list[KVCacheBlock] = []
        if request_id in self._partial_hit_reqs:
            # Partial hit: redirect the shared tail to a private CoW block.
            # Replacing in place keeps the length-based allocation below
            # correct; the extra block was reserved by
            # get_num_blocks_to_allocate.
            block_idx, source_block = self._partial_hit_reqs.pop(request_id)
            cow_block = self.block_pool.get_new_blocks(1)[0]
            self._apply_cow(request_id, block_idx, source_block, cow_block)
            self.new_block_ids.append(cow_block.block_id)
            cow_blocks.append(cow_block)

        """Redirect a partial prefix-cache hit to a private CoW block.

        Both copy endpoints stay retained until the copy has run on the worker,
        so a same-step free cannot recycle them: ``source_block`` keeps its
        hit-ref, ``cow_block`` takes an extra ref beyond the one handed to the
        request.
        """
        req_blocks = self.req_to_blocks[request_id]
        assert block_idx < len(req_blocks)
        assert req_blocks[block_idx] is source_block
        assert not source_block.is_null and source_block.ref_cnt > 0
        req_blocks[block_idx] = cow_block
        self._pending_cow_copies.append((source_block, cow_block))
        cow_block.ref_cnt += 1

The row is patched in place — the only non-append write to a block table anywhere in the allocator — and the copy is deferred to the worker. It travels as SchedulerOutput.kv_cache_block_copies, a list of (src_block_id, dst_block_id) pairs, and runs before the forward pass reads the cache:

        # Apply copy-on-write block copies for partial prefix-cache hits, after
        # zeroing new blocks and before the forward pass reads them.
        if scheduler_output.kv_cache_block_copies:
            copy_kv_cache_blocks_inplace(
                self.kv_caches,
                self.kv_cache_config.num_blocks,
                scheduler_output.kv_cache_block_copies,
            )

The copy itself is a single fancy-index assignment over the block-major storage — not a per-layer loop, because the layers alias the same backing tensor:

    device = storage_tensors[0].device
    indices_np = np.array(kv_cache_block_copies, dtype=np.int64)
    indices = async_tensor_h2d(indices_np, device=device)
    src_indices, dst_indices = indices.unbind(dim=1)

    for tensor in storage_tensors:
        assert tensor.device == device
        blocks = torch.empty(0, dtype=torch.uint8, device=device)
        blocks.set_(tensor.untyped_storage())
        # Block-major backing storage: block i owns the contiguous byte range
        # [i * page_size, (i + 1) * page_size).
        assert blocks.numel() % num_blocks == 0
        blocks = blocks.view(num_blocks, -1)
        blocks[dst_indices] = blocks[src_indices]

The cost is one block copy — 2 MiB across all Llama-3-8B layers — per partial hit. Full-block sharing needs no copy at all: both sequences hold refs via BlockPool.touch, which increments ref_cnt and pulls the block out of the free queue if it was sitting there as an eviction candidate (vllm/v1/core/block_pool.py:L702-L717). Divergence is then automatic, because each sequence appends into its own newly-allocated tail. Reference counting does the work; CoW only patches the one case where a shared block is partially valid.

Everything above is scheduler-side bookkeeping. It becomes the kernel's problem through two tensors on a metadata struct. The backend-agnostic carrier is CommonAttentionMetadata:

    query_start_loc: torch.Tensor
    query_start_loc_cpu: torch.Tensor
    """(batch_size + 1,), the start location of each request in query Tensor"""

    seq_lens: torch.Tensor
    """(batch_size,), the number of computed tokens for each request"""

    num_reqs: int
    """Number of requests"""
    # TODO(lucas): rename to num_tokens since it may be padded and this is misleading
    num_actual_tokens: int
    """Total number of tokens in batch"""
    max_query_len: int
    """Longest query in batch"""
    max_seq_len: int
    """Longest context length (may be an upper bound)"""

    block_table_tensor: torch.Tensor
    slot_mapping: torch.Tensor

Each backend's builder turns that into its own struct; FlashAttention's carries the same two fields verbatim:

    num_actual_tokens: int  # Number of tokens excluding padding.
    max_query_len: int
    query_start_loc: torch.Tensor
    max_seq_len: int
    seq_lens: torch.Tensor
    block_table: torch.Tensor
    slot_mapping: torch.Tensor

block_table is [num_reqs, max_num_blocks_per_req] int32; slot_mapping is [num_tokens] int64. The kernel writes new K/V at slot_mapping[i] and gathers past K/V by walking row r of block_table for ceil(seq_lens[r] / block_size) entries. How it does that efficiently — the tiling, the split-K, the paged gather — is Part 3.

SGLang solves the same problem with a different granularity and a different data structure, and the difference is genuine rather than cosmetic. Its default page_size is 1, which means the default allocator is not paged at all:

    def clear(self):
        # The padded slot 0 is used for writing dummy outputs from padded tokens.
        self.free_pages = torch.arange(
            1, self.size + 1, dtype=torch.int64, device=self.device
        )
        self.is_not_in_free_group = True
        self.free_group = []
        self.release_pages = torch.empty((0,), dtype=torch.int64, device=self.device)

    def available_size(self):
        # To avoid minor "len(free_pages) * 1" overhead
        return len(self.free_pages) + len(self.release_pages)

    def alloc(self, need_size: int):
        if self.need_sort and need_size > len(self.free_pages):
            self.merge_and_sort_free()

        if need_size > len(self.free_pages):
            return None

        select_index = self.free_pages[:need_size]
        self.free_pages = self.free_pages[need_size:]
        return select_index

Compare structure, not just size. vLLM's free list is a CPU-side intrusive doubly linked list of Python objects, traversed and mutated by the scheduler process. SGLang's is a torch.Tensor on the GPU; allocation is a tensor slice and freeing is a torch.cat. Nothing is ever popped one at a time, and there are no per-block Python objects to garbage-collect. The tradeoff is that removing an arbitrary block from the middle is not O(1) — it is not supported — and reuse order degrades until merge_and_sort_free re-sorts.

The counterpart of vLLM's block table is ReqToTokenPool.req_to_token, and it is a token-level table:

        self.size = size
        # +1 padding row at index 0: cuda-graph padded batches default
        # req_pool_indices to 0, so dummy reads/writes land here harmlessly.
        self._alloc_size = size + 1
        self.max_context_len = max_context_len
        self.device = device
        with memory_saver_adapter.region(GPU_MEMORY_TYPE_KV_CACHE):
            self.req_to_token = torch.zeros(
                (self._alloc_size, max_context_len), dtype=torch.int32, device=device
            )
        self.free_slots = list(range(1, self._alloc_size))
        self.req_generation = torch.zeros(self._alloc_size, dtype=torch.int64)

One int32 per token, not per block. That is the entire trade, in one shape. At 256 requests and a 32,768-token window: vLLM's block table is 256 × 2,048 × 4 B = 2 MiB; SGLang's req_to_token is 257 × 32,768 × 4 B = 32.1 MiB, and it is allocated on the GPU inside the KV-cache memory-saver region, so it competes with the KV cache itself. In exchange, SGLang has zero internal fragmentation and its prefix cache can share at exact token boundaries with no copy-on-write machinery at all — the radix tree (§2.4) refcounts tree nodes rather than blocks.

SGLang does have a paged allocator, selected when the backend demands page_size > 1 or DCP is on:

                    elif (
                        get_schedule().page_size == 1 and not get_parallel().dcp_enabled
                    ):
                        token_to_kv_pool_allocator = TokenToKVPoolAllocator(
                            sizes.max_total_num_tokens,
                            dtype=self.kv_cache_dtype,
                            device=self.device,
                            kvcache=token_to_kv_pool,
                            need_sort=need_sort,
                        )
                    else:
                        token_to_kv_pool_allocator = PagedTokenToKVPoolAllocator(
                            sizes.max_total_num_tokens * get_parallel().attn_dcp_size,
                            page_size=get_schedule().page_size
                            * get_parallel().attn_dcp_size,
                            dtype=self.kv_cache_dtype,
                            device=self.device,
                            kvcache=token_to_kv_pool,
                            need_sort=need_sort,
                        )

The paged path is more elaborate than vLLM's precisely because the table stays token-level: a prefill must fill the partial tail page of the prefix, then whole new pages, then a partial final page, which is three cases handled by a Triton kernel (alloc_extend_kernel, dispatched at python/sglang/srt/mem_cache/allocator/paged.py:L197-L204; the readable reference implementation is alloc_extend_naive at L44-L103). vLLM never needs this, because its block table is the page list — appending one block id covers all three cases.

Neither choice is wrong. vLLM optimises index-tensor bandwidth and pays workload-dependent internal fragmentation (15/16 waste for an isolated one-token sequence in a 16-token block); SGLang optimises fragmentation and prefix-sharing granularity and pays 16× the index memory plus per-request page-boundary logic when it does page. The fact that both projects independently defaulted to different ends of this axis is a good sign that the optimum depends on the workload.

The startup error that reads like a memory error but is an over-reservation error. The engine refuses to start if it cannot hold one request at max_model_len:

        raise ValueError(
            f"To serve at least one request with the model's max seq len "
            f"({max_model_len}), ({format_gib(needed_memory)} GiB KV "
            f"cache is needed, which is larger than the available KV cache "
            f"memory ({format_gib(available_memory)} GiB). {estimated_msg}"
            f"Try increasing `gpu_memory_utilization` (which also controls "
            f"CPU memory on the CPU backend) or decreasing `max_model_len` "
            f"when initializing the engine. "

Paging removes most runtime over-reservation, while startup validates the largest supported request. Set max_model_len and memory utilization from the actual workload and measured non-KV headroom. A utilization above 0.95 raises OOM risk but does not inevitably cause it; allocator, graph, workspace, and other-process peaks decide.

The pool exhaustion assertion. get_new_blocks raises ValueError(f"Cannot get {num_blocks} free blocks from the pool") (vllm/v1/core/block_pool.py:L659). Seeing this in production means the capacity check in allocate_slots and the pool disagreed — an accounting bug, not a capacity problem. Reach for it when debugging a new KV cache group type or a connector.

SGLang's version is a warning plus a tree dump. alloc_token_slots and alloc_paged_token_slots_extend both log and then raise, after printing the radix tree:

    if out_cache_loc is None:
        error_msg = (
            f"Out of memory. Try to lower your batch size.\n"
            f"Try to allocate {num_tokens} tokens.\n"
            f"{available_and_evictable_str(tree_cache)}"
        )
        logger.error(error_msg)
        if tree_cache is not None:
            tree_cache.pretty_print()
        raise RuntimeError(error_msg)

The available_and_evictable_str split is the diagnostic that matters: if evictable is large and available is zero, the eviction pass ran too late, not the pool being too small.

Block size is not free to change. Raising --block-size to 32 halves the block table width and the free-list operation count, and doubles worst-case internal fragmentation. It also changes prefix-cache hit granularity, so a workload with many 20-token-boundary shared prefixes can lose hits. Measure hit rate, not just throughput.

Reference counts and preemption interact. Because free_blocks keeps hashes and appends to the tail, a preempted request's blocks are recoverable — but only until enough allocation pressure recycles them. Under sustained pressure the preemption/recompute cycle degrades to full re-prefill. The counter to watch is vllm:num_preemptions alongside the KV usage percentage in the periodic log line (vllm/v1/metrics/loggers.py:L195).

Prove the block accounting to yourself without a GPU-heavy benchmark. Start a server with the pool pinned to a small, known size:

vllm serve meta-llama/Llama-3.1-8B-Instruct \
  --max-model-len 32768 \
  --block-size 16 \
  --num-gpu-blocks-override 2048 \
  --max-num-seqs 64 \
  --no-enable-prefix-caching

2,048 blocks is 32,768 token-slots — exactly one full-length sequence. Send two concurrent 4,000-token prompts. Both are admitted (each needs 250 blocks) even though a contiguous allocator sized for max_model_len could admit only one. Then raise the concurrency until the log line starts reporting preemptions and read the KV usage percentage it prints alongside.

Three things to vary and predict before you run them:

1. --block-size 32 with the block count halved to 1024. Same total token capacity; watch whether steady-state throughput moves. Changes can come from kernels, allocator overhead, prefix-hit granularity, or internal fragmentation even at the same nominal token capacity.
2. Drop --no-enable-prefix-caching and re-send the same prompt twice. The second request should allocate almost no new blocks — the pool's free-queue tail still holds the hashed blocks from the first (§2.3 owns why they match).
3. Set --max-model-len 131072 with the same override and watch the startup check in _check_enough_kv_cache_memory fire with the exact message quoted above.

The SGLang comparison is one flag: python -m sglang.launch_server --model-path meta-llama/Llama-3.1-8B-Instruct --page-size 1 versus --page-size 16. The first uses TokenToKVPoolAllocator, the second PagedTokenToKVPoolAllocator — different code paths, same interface, and the selection is the branch quoted from kv_cache_configurator.py above.

• Kwon et al., Efficient Memory Management for Large Language Model Serving with PagedAttention (arXiv:2309.06180, SOSP 2023) — the original. Read §1–§4 and treat §4.3's copy-on-write as historical.
• vllm#47361 — Delete PagedAttention, the commit that removed the CUDA kernel. Worth reading for the discussion of why general varlen kernels won.
• vLLM V1 alpha release notes — the rewrite that produced vllm/v1/core/block_pool.py and the coordinator/single-type-manager split described here.
• Zheng et al., SGLang: Efficient Execution of Structured Language Model Programs (arXiv:2312.07104) — the token-level allocator and radix tree, from the other side.
• In this book: §2.1 for the sizing arithmetic, §2.3 for block hashing, §2.4 for SGLang's radix tree, §1.4 for what happens when allocate_slots returns None.

The CPU reference allocator stores small integers instead of tensor rows, so every logical token can be inspected. It is a single-group ownership model, not vLLM's or SGLang's allocator. In particular, its fork operation is a pedagogical way to create shared prefixes; it does not imply that the pinned vLLM V1 request path supports beam forking. Run the following after downloading scheduler_sim.py into the same directory as your Python session.

from scheduler_sim import PagedKV

pool = PagedKV(capacity=3, block_size=4)
pool.create("a")
pool.append("a", 10)
pool.fork("a", "b")
source, = pool.pin("a", "transfer-1")
pool.append("b", 20)  # Partial shared tail: copy before writing.
assert pool.read("a") == [10]
assert pool.read("b") == [10, 20]
pool.release("a")    # Cancellation removes ownership, not the transfer pin.
assert source in pool.pages and source not in pool.free
pool.unpin("transfer-1")
assert source in pool.free
pool.check()
pool.release("b")
assert len(pool.free) == 3
print("shared prefix, copy-on-write, and transfer cleanup: OK")

Two kinds of references

An owner reference means a live request's block table names that physical page. A transfer pin means an outstanding operation may still read its storage. Cancellation removes the request's references; it cannot revoke an already-issued read. A page returns to the free heap only when both sets are empty. Transfer handles are unique, and completing an unknown or already-completed transfer raises an error rather than accidentally freeing an unrelated page.

In the example, page zero initially has two owners and one pin. Appending to b allocates a fresh page containing the old token, redirects only b's table, and appends token 20 there. Releasing a leaves page zero with no owners but one pin. Removing that pin finally makes it reusable. The test also covers a sole owner writing a pinned partial page: it must copy too, because the transfer's source snapshot is immutable. Copy-on-write here is synchronous; the pin models an independent outstanding read, not an asynchronous copy-completion engine.

What should every allocator mutation preserve?

• The free set and allocated set are disjoint and together cover every usable physical page exactly once.
• Every page named in a request table exists, and its owner set names that request. Every owner entry has the corresponding reverse table reference.
• Each page is retained by at least one owner or transfer pin. Every pin has a live transfer record, and every transfer record has matching page pins.
• A request table contains no duplicate physical page IDs. Every non-tail page is full; a tail contains between one and block_size tokens.
• A failed allocation leaves the original logical sequence intact. When a one-page pool cannot copy a shared partial tail, the append raises MemoryError; neither branch is silently modified.

PagedKV.check() checks these structural invariants. The randomized test compares every request's decoded page contents with an independent Python-list oracle after every mutation. It also freezes the contents of every transfer-pinned page and checks they never change. Twenty-five fixed seeds execute 200 create, append, fork, pin, unpin or release choices each. Fixed seeds make failures replayable, but do not constitute exhaustive verification or a concurrent race detector.

Tail waste has a distribution

For lengths 1, 15, 16 and 17 with 16-token pages, the unused tail slots are respectively 15, 1, 0 and 15. The observed mean is 7.75 slots only if those four cases are equally likely. In general the expectation is the sum, over workload lengths, of their probability times (B - (length % B)) % B. A uniform residue distribution gives (B - 1) / 2, but real tokenizer and template lengths need not have uniform residues. Shared pages must be counted once physically, not once per logical request.

This allocator has no zero-reference prefix index: an unpinned page with zero owners becomes free immediately. Add an eviction reference or explicit cached state before using it to investigate prefix-cache hit rates. It also excludes heterogeneous cache groups, sliding-window reclamation, asynchronous copies, failures during writes, address-generation fences, distributed transfer acknowledgements and real GPU data. The scheduler model uses the same private-page allocator and adds conservative admission credits and terminal cleanup delays; its reservations do not yet support shared-prefix discounts.