How vLLM CUDA Kernels Write and Read the Paged KV Cache

A request sees its tokens in order. The GPU may store their keys and values in physical cache blocks scattered across a shared pool. How does attention find the right data without first copying the request into one contiguous buffer?

The connection is easiest to follow through two pieces of metadata: slot_mapping tells the cache-write kernel where to put each new token; block_table tells attention where to read a request's logical blocks. They describe the same storage from opposite directions.

This is Part 2 of the vLLM internals series. Part 1 follows a request from the API to GPU execution. You can follow the example below independently.

The kernel details here refer to vLLM source commit 50ac1c7bab47f14d56d86967532574824d02260e, used in the original July 2026 walkthrough. Backend selection and layouts vary across versions. This article follows the FlashAttention cache path at that source revision.

Mapping one token to a physical slot

Suppose a request has three logical KV blocks, with 16 token positions per block. The allocator has assigned physical blocks 7, 2, and 9:

Logical block      Physical block
      0                  7
      1                  2
      2                  9

block_table = [7, 2, 9]

These are illustrative numbers, not a fixed vLLM block size or allocation policy.

Token position 20 belongs to logical block 20 / 16 = 1, at offset 20 % 16 = 4. The block table maps logical block 1 to physical block 2. Its flattened physical slot is therefore:

physical slot = physical block * block size + offset
              = 2 * 16 + 4
              = 36

If position 20 is being processed in the current step, its entry in slot_mapping is 36. The write kernel receives that destination directly. Later, the attention read path can resolve position 20 through the request's block table and reach physical block 2, offset 4 again.

Metadata	What an entry identifies	Used for
slot_mapping	The physical destination of a token processed in this step	Writing newly computed K/V
block_table	The physical block backing a request's logical KV block	Reading cached K/V

The arrays need different shapes because the operations need different information. A step may write only the newest token for a request while attention reads its whole visible history.

Writing K and V into the cache

In csrc/libtorch_stable/cache_kernels.cu, the relevant kernel is reshape_and_cache_flash_kernel. At this revision, one CUDA thread block handles one input token. The destination lookup starts with:

const int64_t token_idx = blockIdx.x;
const int64_t slot_idx = slot_mapping[token_idx];
if (slot_idx < 0) {
  return;
}
const int64_t block_idx = slot_idx / block_size;
const int64_t block_offset = slot_idx % block_size;

A negative slot means that this input position has no cache destination. Padding for a fixed execution shape is one reason such positions can occur.

The division and remainder undo the flattening from the example. Slot 36 becomes physical block 2 and offset 4. Strides then turn those coordinates into a memory address. Here is a simplified address sketch for the key tensor; strides are measured in elements:

key_src = key + token_idx * key_stride;
key_dst = key_cache
        + block_idx * block_stride
        + block_offset * page_stride;

The remaining offsets select the KV head and feature within that token. The value tensor follows the same destination mapping.

This also explains a common source of confusion: a CUDA block is a cooperating group of threads; a KV block is a storage allocation unit. One CUDA block writing one token does not mean it owns an entire KV block.

Why the cache layout changes the copy path

Writing a token involves copying num_heads * head_size key elements and the same number of value elements. How that work is divided depends on which values are adjacent in memory.

Ignoring the outer physical-block dimension, compare these layouts:

NHD: [token, head, feature]
     One token's heads can form one contiguous span.

HND: [head, token, feature]
     The same token's next head is separated by head_stride.

The kernel tests whether the heads are contiguous and whether conversion uses a shared scale. When head_stride == head_size and kv_scale_stride == 0, it can copy a token's heads as one span. Otherwise, it assigns work by head, with a warp cooperating on that head's copy.

The scale condition matters even when the data is contiguous. A quantized cache with separate scales per head needs to apply the corresponding conversion to each head. Layout alone does not determine the path.

The copy helper can move vector packs rather than individual elements. Adjacent lanes handling adjacent packs give the warp a concentrated address range. That is the useful connection to coalescing: inspect the addresses accessed by the active lanes during one instruction. A contiguous chunk within one lane is not enough to describe the whole warp's access pattern.

Reading logical blocks during attention

The attention kernel receives the query data, sequence metadata, the paged cache, and the block table. To read a logical token position, the addressing relationship is conceptually:

logical block  = token position / block size
offset         = token position % block size
physical block = block_table[logical block]
read K/V at    = physical block, offset, head, feature

This is an addressing sketch, not the production kernel's exact loop. A tiled kernel may resolve pages for a whole data region and reuse address calculations across many loads.

The logical sequence remains ordered even though its physical blocks are scattered. The read kernel follows the block table as it loads K/V tiles into the attention computation. It does not need to first gather the whole request into a contiguous K/V tensor.

Paged storage and tiled attention fit together here. Paging decides where the data lives. FlashAttention organizes how the kernel consumes it: retain a query tile, stream key/value tiles, compute a score tile, and update the output using online softmax. The complete score and probability matrices need not be written to GPU global memory.

The write and read implementations must agree on the cache layout. Swapping in a read kernel that expects a different layout is not a matter of changing a function name. In particular, the older native paged_attention_v1/v2 kernels are not the read side of this FlashAttention path.

Following the path in a debugger or source checkout

Start with the selected attention backend and its cache layout. Then follow one token's slot_mapping entry into the write kernel. On the read side, use the corresponding request's block table to resolve that token's logical block. Both should lead to the same physical cache location.

The small example gives a useful checkpoint: logical token 20 maps to physical block 2, offset 4, while its write slot is 36. Once that relationship is clear, layout strides, vectorized copies, and attention tiles become separate questions you can inspect without losing the data flow.

For the full kernel walkthrough and diagrams, see the expanded article on my blog. The next part builds a verifiable FlashAttention forward in PyTorch and Triton, so the online-softmax state can be examined directly.

Source: vLLM cache kernels at the referenced commit. The Chinese edition contains the original walkthrough. This DEV edition is a shorter adaptation with a new illustrative mapping example.

English adaptation prepared with AI assistance.