Part 2 · vLLM internals

Inside vLLM CUDA Kernels: KV-Cache Writes and Paged Attention

This English edition is adapted from the published Chinese article on Zhihu. It preserves the source-code references, experimental boundaries, and reproducible artifacts while adapting the structure for an international engineering audience.

Part 1 stopped at the Python/CUDA boundary. This article crosses it and follows two connected operations: writing newly computed K/V tensors into the paged KV cache, then reading those physical pages from the FlashAttention path. The write and read kernels are discussed as one data flow rather than as isolated snippets.

You need basic C/C++ familiarity, but not prior CUDA programming experience. The first part introduces only the GPU concepts used later: warps, lanes, memory coalescing, memory hierarchy, and tiles.

Version scope. The source references point to vLLM commit 50ac1c7bab47f14d56d86967532574824d02260e from July 14, 2026. Layout branches, quantization paths, and backend selection will continue to change, so the article cites filenames and functions rather than depending on unstable line numbers. The native CUDA paged_attention_v1/v2 comparison is explicitly historical: those kernels were removed from main by PR #47361 on July 2, 2026. For that comparison, this article uses the last pre-removal tree at commit 258f8de91f99b40a4dfb234e55a09e9493c6ece8.

Picking Up at the Python/CUDA Boundary

The previous article stopped at two Python-to-CUDA calls:

# Write KV: Write the newly calculated K and V into the paged KV cache
torch.ops._C_cache_ops.reshape_and_cache_flash(key, value, ..., slot_mapping)

# Read K/V and compute attention
flash_attn_varlen_func(q, ..., block_table=block_table)

These calls mark the end of the Python-level walkthrough, but they do not form one universal path. vLLM supports multiple attention backends, and each read kernel must be paired with the cache layout produced by its write kernel.

Part 1 followed the FlashAttention route. We will therefore read reshape_and_cache_flash and flash_attn_varlen_func as a pair. The removed paged_attention_v1/v2 path appears later only as a historical comparison. It provides an explicit partition-and-reduce implementation and consumes a different KV layout; it is not presented as a backend available at the target commit.

Three questions organize the article: why layout and quantization scale produce two K/V copy paths; how block_table maps a logical K/V tile to physical cache pages; and how online softmax avoids materializing the complete score matrix in GPU memory.

The two sides are examined at different levels. The vLLM cache-write kernel is read directly. The production FlashAttention read kernel is traced at the metadata and data-flow level because its hardware-specialized implementation lives behind generated CuTe/CUTLASS templates rather than one compact CUDA function. The removed native paged_attention_v1/v2 kernels provide the readable historical comparison later in Part 3.

Part 1: Build the GPU Mental Model

This section introduces only the GPU concepts used later. Readers familiar with threads, warps, and the memory hierarchy can skip ahead to tiles and the end-to-end overview.

A CUDA kernel launch creates a grid of thread blocks. Threads within a block can cooperate through shared memory and block-wide synchronization. Hardware issues threads in groups called warps; on NVIDIA GPUs, a warp contains 32 lanes. SIMT execution means those lanes follow one instruction stream, although lane activity and instruction scheduling have architecture-specific details.

Branch divergence occurs when lanes in one warp take different paths. The warp must execute the required paths with non-participating lanes masked off. The cost depends on path length and the active-lane pattern; it is not necessarily a clean 2× penalty. A tail condition such as i < n commonly affects only a boundary warp, while data-dependent branches can diverge across many warps.

Global memory has large capacity and high latency relative to on-chip storage; model weights and the KV cache primarily live there. L2 is shared across the device, while each SM has an L1 cache. A thread block can also use programmer-managed shared memory, and each thread has registers. Kernel tuning balances reuse against occupancy: larger tiles may reuse more data but consume more registers or shared memory, leaving fewer resident warps to hide latency.

Coalesced global-memory access. When a warp executes one global load or store instruction, the memory subsystem coalesces active-lane requests into the transactions needed to cover their addresses. On NVIDIA GPUs with compute capability 6.0 or later, a useful first model is to count the aligned 32-byte segments touched by the warp. Concentrated, aligned addresses usually require fewer transactions; scattered addresses or accesses crossing segment boundaries require more and transfer more unused bytes.

Coalescing is therefore a property of the warp’s address set for one instruction, not of one thread in isolation. Mapping adjacent lanes to adjacent elements is the usual easy-to-coalesce pattern, but lane order is not itself the rule: a permutation that still touches the same few segments may require the same number of transactions. Conversely, one thread reading a contiguous span does not prove that the warp is coalesced. These transactions are requests into the GPU memory hierarchy. They may be served by L1/L2 or continue to HBM/GDDR after a miss, so a memory transaction is not the same thing as a final DRAM command.

Figure B1: CUDA execution hierarchy, GPU memory hierarchy, and the address pattern behind coalesced global-memory access.

A Tile Is a Computation Unit, Not a CUDA Block or KV Block

A matrix is usually too large to keep on chip in full. Kernels therefore operate on smaller submatrices called tiles. For attention:

Q tile: [Br, D]
K tile: [Bc, D]
scores tile = Q tile × K tileᵀ: [Br, Bc]

A group of threads keeps a Q tile resident while streaming K/V tiles through it. Each iteration retains only the current tile data and a small amount of running state. Tile dimensions such as Br and Bc depend on the implementation, dtype, and GPU; no single value is universal.

Three uses of “block” must remain distinct: a tile is an algorithmic data region, a CUDA block or CTA is a cooperating group of threads, and a KV-cache block/page is a storage-allocation unit. A CTA may process one or more tiles, and one tile may read one or more cache blocks.

A Minimal FlashAttention Primer

Standard attention still computes scores = QKᵀ, P = softmax(scores), and O = PV. Materializing full scores and P tensors in HBM creates substantial memory traffic. FlashAttention remains exact; it changes the data flow. Q/K/V are processed in tiles, the current score tile stays in registers or shared memory, online softmax updates the output accumulator immediately, and only the final result is written back.

Paged KV and FlashAttention solve different problems. FlashAttention determines how attention is evaluated tile by tile; block_table determines which physical cache blocks supply a logical K/V range. Together they let the kernel scan a logically ordered sequence stored in non-contiguous physical blocks.

Warp partitioning, pipelines, and specialized instructions differ across FlashAttention generations and GPU architectures. This article keeps only the stable backbone: resident Q tiles, streamed K/V tiles, online softmax, and P·V accumulation.

The overview figure now has enough context to be useful. Its upper section shows new Q/K/V tensors and metadata prepared by Python, the middle shows the paged KV block pool in HBM, and the lower section shows FlashAttention’s on-chip tile pipeline. Green arrows are writes, purple arrows are logical-to-physical lookups, and blue arrows are reads and attention computation.

End-to-end paged-KV data flow. slot_mapping writes newly computed K/V to the physical block pool. During attention, query_start_loc and seq_lens delimit each request, while block_table maps its logical KV blocks to physical blocks. The fused kernel keeps the current Q/K/V tiles on-chip while performing QK, online softmax, and P·V accumulation.

We will revisit each concept at the point where the code uses it. Start with the write kernel.

Part 2: reshape_and_cache Writes the Paged KV Cache

Open csrc/libtorch_stable/cache_kernels.cu and find reshape_and_cache_flash_kernel—the _flash path identified in Part 1. Its signature contains many parameters:

template <typename scalar_t, typename cache_t, Fp8KVCacheDataType kv_dt>
__global__ void reshape_and_cache_flash_kernel(
    const scalar_t* __restrict__ key,    // [num_tokens, num_heads, head_size]
    const scalar_t* __restrict__ value,  // [num_tokens, num_heads, head_size]
    cache_t* __restrict__ key_cache,     // NHD or HND layout
    cache_t* __restrict__ value_cache,
    const int64_t* __restrict__ slot_mapping,  // [num_tokens]
    const int64_t block_stride, const int64_t page_stride,
    const int64_t head_stride, const int64_t key_stride,
    const int64_t value_stride, const int num_heads, const int head_size,
    const int block_size, const float* k_scale, const float* v_scale,
    const int kv_scale_stride) {

Several parameters define the contract:

The template parameters let one kernel support multiple precision paths. scalar_t is the input K/V type, cache_t is the cache-storage type, and kv_dt selects FP8-related conversion behavior. For a first reading, the addressing logic can be followed independently of the conversion details.

key and value contain newly computed K/V with logical shape [num_tokens, num_heads, head_size]. __restrict__ tells the compiler that the pointer ranges do not alias, enabling more aggressive optimization. slot_mapping has one destination slot per token.

The key_cache comment names two layouts, NHD and HND. N is token position (later decomposed into cache block plus in-block offset), H is head, and D is head_size. The order determines which values are adjacent after fixing a token. The *_stride parameters translate multidimensional coordinates into a linear address, just as the row stride does in A[row * stride + col]. Computing strides outside the kernel lets one implementation support both layouts. The reason for two copy paths becomes clear once we compare those addresses.

The contract is now clear: the kernel receives newly computed K/V, slot_mapping, and layout strides, then writes key_cache and value_cache. The body answers three questions: how work is assigned to threads, how the destination address is formed, and how values are copied.

Figure B2: the reshape_and_cache_flash write path. One CUDA block handles one token. NHD with a shared scale can copy all heads as one contiguous span; HND or per-head scales use a warp-per-head path.

One CUDA Block Handles One Token

const int64_t token_idx = blockIdx.x;

This establishes the work assignment: one CUDA block handles one token. Part 1 flattened requests into [num_tokens]; here the grid launches num_tokens blocks, with blockIdx.x selecting the token.

One token contains num_heads × head_size K elements and the same number of V elements. Assigning a thread block to that token exposes enough parallel copy work while keeping one block responsible for one destination slot.

Decomposing a Slot into Block and Offset

const int64_t slot_idx = slot_mapping[token_idx];
if (slot_idx < 0) {     // Negative means this padded/unused token has no cache destination
  return;
}
const int64_t block_idx = slot_idx / block_size;
const int64_t block_offset = slot_idx % block_size;

The lookup returns the flattened physical cache slot for this token. A negative value marks an unused position with no cache destination—for example, padding introduced for a fixed CUDA-graph shape—so the block returns without writing.

The quotient identifies the physical cache block and the remainder identifies the token offset within it. This is the inverse of slot = block_id × block_size + in_block_offset from Part 1.

The cache is laid out by physical block, so the flattened slot must be decoded into “which block” and “which token position within that block” before layout strides can form an address.

Strides then form the source address in the newly computed K/V tensor and the destination address in the paged cache:

// Source: the starting position of this token in the input key/value
const scalar_t* key_src = key + token_idx * key_stride;
// Target: the starting position of this token in the KV cache
cache_t* key_dst = key_cache + block_idx * block_stride + block_offset * page_stride;

block_idx * block_stride selects the cache block; block_offset * page_stride selects the token position within that block. The remaining head and feature offsets depend on NHD versus HND layout.

Why Layout and Quantization Scale Produce Two Copy Paths

After the addressing is completed, the kernel needs to answer two questions:

1. Is the layout contiguous? For one token, are all heads adjacent in memory?
2. Is the conversion uniform? When writing a quantized cache, do those heads share one scale?

Start with the layout. Omitting the outer cache-block dimension makes the relative positions of H and N easier to see:

NHD: [token, head, D]
Fixed token: D elements of head 0 | D elements of head 1 | head 2 ...
Address: token_base + head * D

HND: [head, token, D]
Fixed token: D elements of head 0 |<-- head_stride -->| D elements of head 1...
Address: block_base + head * head_stride + token_offset * D

In NHD, the num_heads × head_size values for one token are contiguous when head_stride == head_size, so the kernel can treat them as one long span. In HND, tokens within one head are adjacent, but the same token’s next head is a head_stride away; the whole row cannot be copied as one contiguous span.

Scale is the second constraint. kv_scale_stride == 0 means every head shares one scale; otherwise head h uses scale[h * kv_scale_stride]. Even when the NHD data is contiguous, different per-head conversion parameters prevent one uniform conversion over the whole row. Assigning work by head lets a warp reuse one head’s scale and keeps both the conversion and addressing explicit.

The branch therefore depends on both layout contiguity and scale granularity:

const bool is_contiguous_heads = (head_stride == head_size);

if (is_contiguous_heads && kv_scale_stride == 0) {
  // NHD + shared scale: all heads for this token form one contiguous span
  vectorize_with_alignment(
      key_src, key_dst, num_heads * head_size,
      threadIdx.x, blockDim.x, k_op);
} else {
  // HND, or a separate scale per head: use the warp-per-head path below
}

With contiguous NHD data and one shared scale, the whole thread block can copy the token’s head data as one span. Otherwise, the kernel uses the per-head path.

Warp, lane, and warp_id

CUDA groups threads into 32-lane warps. A lane is a thread’s index within its warp, from 0 to 31; warp_id identifies the warp within the block:

const int lane = threadIdx.x & 31;        // Equivalent to threadIdx.x % 32
const int warp_id = threadIdx.x >> 5;     // Equivalent to threadIdx.x / 32
const int warps_per_block = blockDim.x >> 5;

For a 256-thread block, threadIdx.x == 70 belongs to warp 2, lane 6. The outer loop assigns heads to warps; the inner helper assigns vector packs within one head to lanes.

The kernel uses the following work assignment for HND or per-head quantization scales:

for (int head = warp_id; head < num_heads; head += warps_per_block) {
  const scalar_t* k_src_h = key_src + head * head_size;
  cache_t* k_dst_h = key_dst + head * head_stride;
  ...
  // The warp's 32 lanes cooperate on the vector copy for this head.
  vectorize_with_alignment<VEC_SIZE>(k_src_h, k_dst_h, head_size, lane, 32, k_op);
}

The outer loop distributes heads across warps. A 256-thread block contains eight warps: warp 0 handles heads 0, 8, 16, and 24; warp 1 handles heads 1, 9, 17, and 25; and so on. All 32 lanes in a warp process the same head and share its destination base and quantization scale. Their offsets within that head differ.

Why adjacent lanes should move adjacent packs

The inner vectorize_with_alignment assigns work by vector pack, not by individual scalar. On its aligned fast path, the core loop is equivalent to:

int num_vec = len / VEC_SIZE;
for (int i = tid; i < num_vec; i += stride) {
  src_pack = vector_load(src + i * VEC_SIZE);
  dst_pack = op(src_pack);
  vector_store(dst + i * VEC_SIZE, dst_pack);
}

The HND path passes tid = lane and stride = 32, so the first loop iteration assigns pack i to lane i. For half or bfloat16 with VEC_SIZE=8, each pack contains 8 × 2 B = 16 B:

lane 0 -> pack 0 -> elements 0–7  -> bytes 0–15
lane 1 -> pack 1 -> elements 8–15 -> bytes 16–31
lane 2 -> pack 2 -> elements 16–23 -> bytes 32–47
...
lane 31 -> pack 31 -> bytes 496–511 (if this round has at least 32 packs)

There are two kinds of locality here. Each lane loads one contiguous 16-byte pack, and adjacent lanes load adjacent packs. With 32 valid packs, the warp covers a contiguous 32 × 16 B = 512 B range. With fewer packs, a prefix of the lanes participates, but their addresses are still contiguous. The hardware can therefore service the warp with a small number of aligned memory transactions. Coalesced access does not mean that all 32 lanes are combined into one transaction. The exact number of sectors or transactions depends on the GPU architecture, cache-line/sector granularity, access width, and alignment.

When can one instruction touch only one 32-byte sector?

Compare the vectorized helper with this scalar work assignment:

for (int i = lane; i < head_size; i += 32) {
  dst[i] = op(src[i]);
}

Across loop iterations, each lane handles a strided sequence:

lane 0: elements 0, 32, 64, 96
lane 1: elements 1, 33, 65, 97
...

This scalar pattern is also coalesced when viewed iteration by iteration. In the first loop iteration, lanes 0–31 read elements 0–31 concurrently; in the next, they read 32–63. Lane 0 eventually handles 0, 32, 64, and 96, but those accesses occur in different warp instructions. Each instruction still covers adjacent addresses.

“One transaction for the whole warp” is meaningful only after specifying access width, alignment, and the transaction model. If 32 lanes each access one byte from an aligned 32-byte range, the instruction touches one 32-byte sector. If each lane accesses a half, the range is 64 bytes and touches at least two such sectors; floats cover 128 bytes and at least four. Misalignment can add another sector. Loads and stores must also be counted separately—for example, a half load followed by an FP8 store has different widths on the two sides. Older descriptions sometimes call one aligned 128-byte warp access “one coalesced transaction”; under the 32-byte-segment model used here, that phrase should not be read as one physical transaction.

The scalar version can also be coalesced, so vectorization does not make the required bytes disappear. It primarily reduces instruction, loop, and address calculation overhead. For head_size=128 with half-precision elements:

Scalar work assignment: 32 lanes × 2B = 64B per round, 4 rounds in total
Vector pack: 16 valid lanes × 16B = 256B moved in one round

Both methods move 256 bytes and can generate contiguous warp accesses. The pack version lets each active lane move eight halves with a vector load/store, requiring fewer loop iterations and instructions. The helper checks alignment before taking the vector path and handles unaligned prefixes or tails separately. Its goal is to preserve coalescing while doing more useful work per lane instruction, not to collapse several required sectors into one.

By contrast, if lane 0 reads bytes 0–15, lane 1 reads 64–79, and lane 2 reads 128–143, each lane’s access is locally contiguous but the warp touches many separated segments. Coalescing depends on the address set of the entire warp for one instruction, not on one lane in isolation.

Later iterations add stride=32, preserving the adjacent-lane/adjacent-pack relationship for the next group of packs. The helper handles any unaligned prefix or tail outside the vectorized core.

Figure B3: vectorize_with_alignment assigns adjacent 16-byte packs to adjacent lanes. A scalar stride can also be coalesced but needs more instructions; a scattered P0/P4/P8 assignment touches more memory segments.

k_op applies the write-side conversion. It is effectively a copy when source and cache formats match; an FP8 cache path also applies scaling and type conversion. Vector width determines both instruction count and the mapping from lanes to elements.

What This Part Established

Stripping away the template and FP8 details, the logic of this kernel is:

• One block is responsible for one token (token_idx = blockIdx.x).
• Read the physical destination from slot_mapping[token_idx]; skip negative entries that represent unused positions.
• Split the slot into block_idx and block_offset, then use tensor strides to form the source and destination addresses.
• Contiguous NHD data with one scale uses the whole-span fast path; HND data or per-head scales use the warp-per-head path.
• vectorize_with_alignment assigns adjacent vector packs to adjacent lanes and handles misaligned prefixes and tails.

This is the kernel behind Part 1’s reshape_and_cache_flash(..., slot_mapping) call. One CUDA block handles one token, slot_mapping selects the destination, and layout plus scale granularity select either the whole-span fast path or the warp-per-head path.

Part 3: Reading Paged K/V through the FlashAttention Path

The overview established the complete data path, and Part 2 examined its write half. We now follow the corresponding production call into the read half: resolve logical pages, load K/V tiles, and accumulate attention output.

Zooming into the Read Path

Two independent designs meet at this boundary:

• Paged KV solves placement: a request sees a logically ordered history, but physical K/V blocks may be scattered; block_table translates logical blocks to physical blocks.
• FlashAttention solves tiled evaluation: QK, softmax, and P·V are fused so the full score matrix is not materialized in global memory.

Their combination is what this article calls FlashAttention over paged KV. FlashAttention scans the logical KV sequence; for each K/V tile, block_table resolves the physical pages before the kernel loads data, computes QK, and updates online softmax. It neither gathers paged K/V into a temporary contiguous tensor nor calls vLLM’s native paged_attention_v2 kernel.

Keep two units separate: a page or KV block is a storage and allocation unit, whereas a tile is a unit of computation. A K/V tile may lie inside one physical block or span several blocks. The kernel advances by tiles while block_table resolves the underlying pages.

The actual call in the previous article appears as follows on the Python side:

flash_attn_varlen_func(
    q=query,
    k=key_cache,
    v=value_cache,
    cu_seqlens_q=query_start_loc,
    seqused_k=seq_lens,
    block_table=block_table,
    ...
)

flash_attn_varlen_func dispatches according to GPU and configuration. Tile sizes, warp partitioning, and pipelines vary by implementation, but the stable data flow can be summarized in five steps:

① Use metadata to select the request, Q tile, and visible KV range
② Advance through logical KV and resolve physical blocks through block_table
③ Load the current K/V tile and compute its QK score tile
④ Update online softmax and the P·V accumulator
⑤ Write the result; a split-KV path first combines partial states stably

Figure B4: FlashAttention over paged KV. Metadata supplies request boundaries, visible KV lengths, and logical-to-physical block mappings; stages ②–④ repeat until the visible KV range has been scanned.

The next sections expand three points along this path.

How Metadata Connects a Q Tile to Paged K/V

FlashAttention does not receive a padded [batch, seq, ...] tensor here. vLLM packs request tokens, then supplies three metadata tensors: query_start_loc marks request ranges in the packed query, seq_lens gives each request’s visible KV length, and block_table maps logical KV blocks to physical blocks.

The KV cache uses one of two logical shapes:

NHD: [num_blocks, block_size, num_kv_heads, head_size]
HND: [num_blocks, num_kv_heads, block_size, head_size]

K and V use the NHD or HND paged layout written by reshape_and_cache_flash. query_start_loc identifies the query’s request, seq_lens bounds its visible history, and block_table resolves that history’s physical cache blocks. One variable-length kernel can therefore handle requests with different lengths and non-contiguous block placement.

For one CTA, the conceptual loop is:

for each query tile:
  initialize this tile's softmax state

  for each logical KV tile:
    use block_table to locate the physical KV blocks it covers
    load the K/V tile directly from the paged cache
    compute this QK tile, update softmax, and accumulate the output

  write back the normalized output tile

Address translation can be summarized as physical_block = block_table[seq_idx, logical_block]. As the kernel advances through logical KV positions, it repeatedly resolves the underlying blocks; it does not first gather the complete history into a contiguous temporary tensor.

FlashAttention implementations assign data movement, matrix multiplication, and tiles per CTA differently. A particular release’s threadIdx.x mapping should therefore not be treated as a permanent rule. The invariant is that a CTA owns a Q tile, resolves and loads K/V tiles while scanning the K dimension, then performs QK, softmax, and P·V. Under GQA/MQA, several query heads reuse one KV head; the selected backend decides the exact mapping.

How Each K/V Tile Contributes to the Output

The attention equations remain unchanged:

scores = q @ k.T * scale
probs  = softmax(scores)
out    = probs @ v

An unfused implementation may write full scores to global memory and read it back for softmax and probs @ v. FlashAttention computes and consumes one score tile within the fused kernel. Stages ③ and ④ in Figure B4 therefore share one data-flow lifetime rather than materializing a complete intermediate matrix.

For each query row, the kernel maintains a running maximum m, normalization sum l, and unnormalized output accumulator acc:

m_new   = max(m, rowmax(scores_tile))
alpha   = exp(m - m_new)
p       = exp(scores_tile - m_new)
l_new   = alpha * l + rowsum(p)
acc_new = alpha * acc + p @ V_tile

The final output is acc / l. alpha rescales prior state when the running maximum changes, making the recurrence algebraically equivalent to a stable softmax over the complete row, aside from normal floating-point rounding.

“Online” means that m/l/acc is updated as each tile is consumed; earlier score tiles are not retained. For a short query with a long KV history, parallelism over queries and heads may be insufficient. Some implementations therefore split the KV range, compute partial output and log-sum-exp state, and combine the partial states with stable max/LSE rescaling.

Split-KV changes the work partition, not the softmax identity. Without splits, one state is updated sequentially across tiles. With splits, CTAs produce local states that are rescaled to a common maximum before merging.

Why This Path Is Not the Legacy paged_attention_v2

FlashAttention split-KV and paged_attention_v2 use the same stable-softmax principle to merge partial ranges, but they are different kernel families:

• The number of splits, tile scheduling, and combine paths for FlashAttention are determined by the chosen implementation and runtime configuration.
• The legacy paged_attention_v2 uses an explicit partition kernel followed by a separate reduction kernel in csrc/libtorch_stable/attention/ of the pre-removal tree.

paged_attention_v2 does not mean FlashAttention-2. The “v2” names vLLM’s native partition-and-reduce variant, whose control flow is comparatively easy to inspect:

paged_attention_v2_kernel
Each block is responsible for seq × head × partition
Save local max_logits/exp_sums/tmp_out
        ↓
paged_attention_v2_reduce_kernel
Recalibrate each partition and merge tmp_out

This path is paired with reshape_and_cache and its dedicated cache layout; it is not the read side of reshape_and_cache_flash.

paged_attention_v1 shares most of its compute body with v2 but launches work differently:

The comparison is most useful after following the production FlashAttention path. paged_attention_v2 then provides a more explicit example of splitting the K/V range and stably merging partial softmax results; v1 completes the comparison by showing the cost of scanning the whole sequence in one block.

What This Part Established

• The production path examined here is reshape_and_cache_flash → flash_attn_varlen_func, not the legacy paged_attention_v2.
• FlashAttention over paged KV combines logical-to-physical block addressing with tiled fusion and online softmax.
• query_start_loc, seq_lens, and block_table provide packed-query boundaries, visible KV lengths, and physical block mappings. The kernel need not gather the complete K/V history first.
• QK, online softmax, and P·V share a fused tile data flow without materializing the full score matrix in global memory.
• Split-KV and native v2 partitioning share stable merge mathematics but use different layouts, scheduling, and kernels.
• The legacy paged_attention_v2 is useful as a readable comparison path; v1 needs only a brief work-partitioning comparison.

What to Carry Forward

Three design choices carry beyond these particular kernels.

First, cache layout and consuming kernel form one contract. The NHD/HND cache written by reshape_and_cache_flash is consumed by the corresponding FlashAttention path; native PagedAttention uses another layout. A layout cannot be evaluated independently of the kernel it serves.

Second, FlashAttention consumes QK, online softmax, and P·V within the lifetime of a tile, avoiding a round trip of the full score matrix through global memory. When a long K/V range is split for parallelism, local maximum and log-sum-exp state allow the partial results to be combined into the same global softmax.

Third, block_table brings page-table-style indirection to the KV cache. A logically contiguous history may be physically scattered across the block pool. The scheduler maintains that mapping; the attention kernel resolves it while loading each K/V tile.

Together, the first two articles now trace the path from llm.generate() to paged addressing, QK tiles, online softmax, and P·V accumulation inside the FlashAttention backend. Going deeper means entering generation- and architecture-specific territory: warp specialization, asynchronous pipelines, and Tensor Core instruction selection.

What this article doesn’t cover

To keep the main path readable, this article leaves out:

• FP8/BF16 precision branches and in-kernel quantization conversions; the diagrams treat values as unquantized for clarity.
• Generation-specific tile parameters, warp specialization, asynchronous copy pipelines, and backend selection. The article retains only the common data flow.
• How a FlashAttention split-KV scheduler chooses split counts and how its combine kernel is implemented at thread level.
• Thread-level reduction details inside paged_attention_v2_reduce_kernel; the kernel appears only as a comparison path.
• Warp-level reductions, shared-memory bank conflicts, and lower-level memory optimizations.

This walkthrough targets vLLM commit 50ac1c7bab47f14d56d86967532574824d02260e, primarily csrc/libtorch_stable/cache_kernels.cu, csrc/libtorch_stable/quantization/vectorization_utils.cuh, and vllm/v1/attention/backends/flash_attn.py. The removed native kernels from csrc/libtorch_stable/attention/ at commit 258f8de9 are used only for the v1/v2 comparison. For current coalescing and transaction details, see the NVIDIA CUDA C++ Best Practices Guide. FlashAttention templates vary by version and architecture, so the article retains only the common execution skeleton. Editable Excalidraw sources for Figures B0–B4 are published with the site.