Block table: mapping logical tokens to physical KV blocks

源码版本v0.25.1

Responsibilities

The core idea of PagedAttention is to slice the KV cache into fixed-size blocks. Which block and which slot each token lands in is looked up via the "block table". vLLM v1 implements the block table as the BlockTable class: an int32 tensor of shape (max_num_reqs, max_num_blocks_per_req). Each row corresponds to one request, and the j-th element of a row is the physical block_id for that request's j-th logical block. MultiGroupBlockTable packages the block tables of multiple KV cache groups together, so that hybrid models (e.g. full attention + sliding window) each have their own copy. BlockTable also holds a slot_mapping tensor — directly mapping each query token to a physical slot id in the KV cache, which the attention kernel uses to read and write KV directly.

The block table's lifecycle is tied to the request: at InputBatch.add_request, block_table.add_row(request.block_ids, req_index) writes the list of block_ids returned by the manager into that row (gpu_input_batch.py:336-379). After each scheduling step, newly allocated block_ids are appended to the row tail via block_table.append_row(new_block_ids, req_index) (gpu_model_runner.py:1425-1428). commit_block_table copies the numpy array to the GPU (block_table.py:166-167); compute_slot_mapping runs a Triton kernel to compute each token's slot id (block_table.py:141-164). During the forward pass, the attention backend reads this tensor directly from common_attn_metadata.block_table_tensor (backend.py:420-421).

Design motivation

Why does the worker maintain its own copy of the block table, instead of using the manager's req_to_blocks directly?

• numpy + GPU double buffering: block_table uses CpuGpuBuffer to hold three copies at once (numpy / CPU / GPU). It writes on CPU (numpy) and reads on GPU (int32 tensor), so commit_block_table(num_reqs) only needs one H2D copy (block_table.py:70-77). Under CUDA graph capture the size is fixed and the GPU tensor is reused directly.
• kernel block_size ≠ manager block_size: Some attention kernels (e.g. TRTLLM) use block_size 16, while the KV cache manager's block_size is 32. map_to_kernel_blocks expands one manager block into blocks_per_kv_block=2 kernel blocks (block_table.py:47-68) (block_table.py:173-201).
• slot_mapping as a kernel: compute_slot_mapping calls the Triton kernel _compute_slot_mapping_kernel. Each program handles one request, computing block_index = pos // block_size from positions, then looking up block_table[req_idx, block_index] to get block_number, and finally slot_id = block_number * block_size + offset (block_table.py:325-380). This is faster than a numpy implementation and also runs inside a CUDA graph.
• DCP / PCP interleaving: Three constexprs — TOTAL_CP_WORLD_SIZE / TOTAL_CP_RANK / CP_KV_CACHE_INTERLEAVE_SIZE — let the kernel, under decode context parallelism, write real values only to slots belonging to the current rank and write PAD_ID to the rest (block_table.py:368-379).
• CUDA graph padding: When num_reqs < num_reqs_padded, _get_block_table fills padding rows with NULL_BLOCK_ID (gpu_model_runner.py:2279-2295). The attention kernel skips rows that see NULL_BLOCK_ID.
• Multi-group builder cache reuse: Builders with supports_update_block_table=True do not rebuild on the second group with the same spec — they only update_block_table (gpu_model_runner.py:2477-2493), saving redundant metadata construction.
• Row-level maintenance: append_row / clear_row / move_row / swap_row treat the block table as a dynamic array indexed by request row. In scenarios like condense (request rearrangement) and preempt (request swap-out), they update in place without rebuilding the whole thing (block_table.py:102-139).

Key files

• BlockTable.__init__:18-100 — Double-buffer construction; handles the hybrid block mode where kernel_block_size != block_size.
• BlockTable.append_row/add_row:102-122 — Append to row tail / rewrite the entire row of block_ids.
• BlockTable.clear/move/swap_row:124-139 — Row-level maintenance supporting condense and preempt rearrangement.
• BlockTable.compute_slot_mapping:141-164 — Calls the Triton kernel to convert positions into the slot_mapping tensor.
• BlockTable.commit_block_table/clear:166-171 — H2D copy / zeroing.
• BlockTable.map_to_kernel_blocks:173-201 — Expands the manager block_id list into the kernel block_id list.
• BlockTable.get_device_tensor:203-213 — Returns the GPU tensor slice for the attention backend.
• MultiGroupBlockTable:223-322 — Multi-group wrapper; broadcasts add_row / append_row / compute_slot_mapping to each group's BlockTable.
• _compute_slot_mapping_kernel:325-380 — Triton kernel: per-request program computing slot_id = block_table[req, pos // block_size] * block_size + offset; under DCP filtered by TOTAL_CP_RANK.
• InputBatch.block_table init:172-180 — MultiGroupBlockTable construction point.
• InputBatch.add_request:336-379 — Calls block_table.add_row(request.block_ids, req_index) when a new request arrives.
• _get_block_table:2279-2298 — Per-step preparation of block_table_tensor with NULL_BLOCK_ID padding.
• append_row on new blocks:1425-1428 — Newly allocated block_ids appended to the existing row tail.
• commit_block_table:1933-1933 — After scheduling, copies the numpy array H2D to GPU.
• CommonAttentionMetadata:2378-2396 — Where block_table_tensor is attached to common metadata.
• update_block_table:2477-2493 — Groups with the same spec reuse already-built metadata.
• block_table_tensor field:420-421 — The block_table field in CommonAttentionMetadata.
• FlashAttentionMetadata.block_table:234-234 — The block_table tensor held by the Flash backend.

Data flow

In each scheduling step, gpu_model_runner._prepare_inputs translates the SchedulerOutput into input-batch state: for each existing request it calls block_table.append_row(new_block_ids, req_index) to append new blocks to the row tail; for empty slots after condense it uses move_row / swap_row to tidy up (gpu_input_batch.py:629-758). Then commit_block_table(num_reqs) H2D-copies the numpy array to GPU, and compute_slot_mapping runs the Triton kernel to compute slot_mapping:

@triton.jit(do_not_specialize=["num_tokens", "max_num_tokens"])
def _compute_slot_mapping_kernel(
    num_tokens,
    max_num_tokens,
    query_start_loc_ptr,  # [num_reqs + 1], int32
    positions_ptr,  # [num_tokens], int64
    block_table_ptr,  # [max_num_reqs, max_num_blocks_per_req], int32 (flat)
    block_table_stride,  # max_num_blocks_per_req
    block_size,
    slot_mapping_ptr,  # [max_num_tokens], int64
    TOTAL_CP_WORLD_SIZE: tl.constexpr,
    TOTAL_CP_RANK: tl.constexpr,
    CP_KV_CACHE_INTERLEAVE_SIZE: tl.constexpr,
    PAD_ID: tl.constexpr,
    BLOCK_SIZE: tl.constexpr,
):
    req_idx = tl.program_id(0)

    if req_idx == tl.num_programs(0) - 1:
        # Pad remaining slots for CUDA graph compatibility.
        for i in range(num_tokens, max_num_tokens, BLOCK_SIZE):
            offsets = i + tl.arange(0, BLOCK_SIZE)
            tl.store(
                slot_mapping_ptr + offsets,
                PAD_ID,
                mask=offsets < max_num_tokens,
            )
        return

(block_table.py:325-352) The last program is dedicated to writing PAD_ID into padding slots, preventing leftover data from contaminating attention when CUDA graphs are reused. Then _get_block_table(kv_cache_gid) fetches this group's GPU tensor (gpu_model_runner.py:2279-2295) and stuffs it into CommonAttentionMetadata.block_table_tensor (gpu_model_runner.py:2378-2396). Flash attention / Triton / ROCm backends each read attn_metadata.block_table (flash_attn.py:234-234) (triton_attn.py:208-234), and combined with slot_mapping complete the read / write of KV cache. The block_ids come from KVCacheManager allocated through Coordinator and BlockPool.

Boundaries and failure modes

• kernel_block_size does not divide block_size: block_size % kernel_block_size != 0 directly raise ValueErrors (block_table.py:58-62).
• Appending an empty list: append_row([]) returns directly without updating num_blocks_per_row (block_table.py:107-108).
• clear_row does not write zeros before clearing num_blocks_per_row: clear_row zeros block_table.np[row, :num_blocks] first, then zeros the count (block_table.py:124-128).
• CUDA graph padding: Rows in num_reqs:num_reqs_padded are filled with NULL_BLOCK_ID (gpu_model_runner.py:2292-2294); the kernel skips on seeing NULL_BLOCK_ID.
• EncoderOnlyAttentionSpec uses a zero tensor: Encoder-only KV cache has no block_table. _get_block_table returns a (num_reqs_padded, 1) all-zero tensor as a placeholder (gpu_model_runner.py:2282-2287).
• PCP / DCP not initialized: get_pcp_group() / get_dcp_group() may not be initialized in tests; a try/except falls back to world_size=1 (block_table.py:86-99).
• Mamba reuses block_table_tensor: Mamba / GDN / Linear attention backends use mamba_get_block_table_tensor to treat block_table as state_indices; they need a gather in cache mode (utils.py:925-965).
• max_num_blocks aligned for TRTLLM: MultiGroupBlockTable.__init__ aligns max_num_blocks to a multiple of 128 // block_size, as required by some attention backends (TRTLLM) (block_table.py:260-265).

Summary

BlockTable organizes the block_id list allocated by KVCacheManager / Coordinator and BlockPool into a GPU tensor, computes each token's physical slot via the slot_mapping Triton kernel, and stuffs it into CommonAttentionMetadata for all attention backends to use. It is decoupled from model weight loading (see DefaultModelLoader) and only couples with the KV cache management layer.

See official docs: vLLM docs · README