Asynchronous KV Cache Prefetching

1. TL;DR

During the PagedAttention decode kernel, each warp iterates over the sequence's KV blocks one at a time. When a warp finishes block $b$ and reaches for block $b{+}\text{NUM\_WARPS}$, the address of that next block is only resolved as the iteration begins, so the warp stalls on a long-latency global load before its compute can proceed (an exposed long-scoreboard stall). This paper adds a single PTX instruction — cp.async.bulk.prefetch.L2.global — that, while the current block is still being multiplied, issues an asynchronous bulk prefetch of the next block's key (and later value) tensor toward L2. Because the request is issued early, those bytes are moving toward L2 ahead of use, so when the warp reaches that block its long-latency load is hidden behind compute rather than serialized in front of it. (Caveat, measured: the baseline KV is already ~85% L2-resident from natural reuse — see Sections 6-7 — so the prefetch is hiding the exposed load latency, not removing traffic or changing what sits in L2.)

2. Motivation — Why Prefetch the KV-Cache At All?

2.1 Why now: decode is memory-bound and getting worse

Autoregressive decoding generates one token per step and, at each step, must read the entire key/value history of every active sequence. With batch size $B$, sequence length $L$, number of KV heads $H_{kv}$, and head dimension $d$ in two bytes, the per-step KV read is roughly $2 \cdot B \cdot L \cdot H_{kv} \cdot d \cdot 2$ bytes, while the matmul work per step is tiny — a single query vector against the cached keys. The arithmetic intensity of decode is therefore far below the GPU's compute-to-bandwidth ratio, so the attention kernel spends most of its time waiting on HBM. As context windows grow into the tens of thousands of tokens and agentic workloads accumulate ever-longer histories, the fraction of decode time spent moving KV bytes only increases.

2.2 The bad scenario: a warp stalls every block boundary

PagedAttention stores the KV-cache in fixed-size physical blocks scattered across HBM, indexed through a per-sequence block table. The decode kernel walks these blocks in a strided loop — warp $w$ handles blocks $w, w{+}\text{NUM\_WARPS}, w{+}2\,\text{NUM\_WARPS}, \dots$ Each loop iteration must (1) translate the logical block to a physical block via the block table, (2) issue a global load for that block's keys, and (3) do the dot products. Because the load address is only known once the iteration begins, the kernel cannot issue that load early enough for the hardware to hide its long latency; the warp issues the load and then sits idle for hundreds of cycles waiting on it — what Nsight calls a long-scoreboard stall — before compute can resume. The stall is the exposed latency of a load that was issued too late, not a claim that the bytes had to travel from HBM (in steady state much of the KV is already L2-resident); the cost is that the warp waits in front of compute instead of behind it. Repeated across every block of every sequence, these stalls dominate the kernel's runtime.

2.3 Baseline pain: what stock vLLM already does, and where it stops

Stock vLLM's PagedAttention is already a heavily-tuned kernel: it vectorizes the loads, uses shared memory for the running softmax, and overlaps thread-group work. What it does not do is look ahead in the block table. The hardware prefetcher cannot help, because the next block's address is a pointer-chase through the block table rather than a linear stride in virtual memory — the very indirection that makes paging flexible also defeats automatic prefetch. So the baseline issues each block's load only at the start of its iteration and pays that load's full latency in the open, even though the address of the next block is, in fact, computable several iterations early from the same block table the kernel already holds in registers. (The latency is exposed because the load is issued late, not because the bytes are necessarily cold in HBM — in steady state much of the KV is already L2-resident; what's missing is the early issue, not the data.)

2.4 Strawman approaches and why they fall short

2.5 The obstacle and the positioning

The obstacle the paper must clear is precision-without-disruption: it needs to issue the next block's load/prefetch early enough to hide exposed long-scoreboard latency before use, without consuming shared memory, without blocking the issuing warp, and without rewriting the kernel. Hopper (sm_90+) supplies exactly the right primitive — cp.async.bulk.prefetch.L2.global — a non-blocking bulk-prefetch hint that requests data be moved from global memory toward L2 and returns immediately, touching neither registers nor shared memory of the issuing warp. (It is a scalar L2-prefetch instruction in the same async-bulk-copy family as TMA, but unlike a full TMA tensor copy it needs no tensor-map descriptor and writes nothing into shared memory.) The paper's positioning is therefore narrow and honest: it is not a new attention algorithm; it is a minimal, additive latency-hiding patch that exploits a Hopper hardware feature to prefetch a tensor whose address the software already knows.

2.6 Contributions

1. A one-iteration-ahead block lookahead computed from the existing block table (block_idx + NUM_WARPS), reusing data the kernel already touches.
2. An asynchronous L2 prefetch for both the key and the value cache of that next block, issued by a single lane while the rest of the warp computes the current block.
3. A drop-in deployment story: a single header file replaces vLLM's attention_kernels.cuh; no scheduler, model, or API changes. Measured ~2.15× kernel and up to ~1.97× e2e on H20.

3. Core Mechanism

3.1 The overlap, visualized

The whole idea is to convert serial "load-then-compute" into overlapped "compute-while-loading". In the baseline, each block iteration begins with an exposed late global-load (long-scoreboard) stall (the warp waits in front of compute), then computes. In the patched kernel, lane 0 fires a prefetch for the next block at the start of the current iteration; the bulk-prefetch engine issues that next block's load toward L2 in the background while the warp computes the current block. By the time the next iteration needs that block, its load is already in flight, so the long-latency global load is hidden behind compute rather than serialized in front of it. (Measured caveat: the baseline KV is already ~85% L2-resident, so the observable signature is a small drop in memory stalls, not a large L2-hit jump — see Sections 6-7.)

Baseline: exposed late global-load (long-scoreboard) latency

Patched: prefetch overlapped with compute

The orange dashed cells (baseline stalls) shrink once the next block's bytes have been pulled toward L2 ahead of use. The bottom lane is the non-blocking bulk-prefetch path through the cache subsystem, independent of the warp's compute, which is what makes the overlap free of register and shared-memory cost.

3.2 The lookahead distance: why NUM_WARPS, not 1

Because the block loop is strided by NUM_WARPS — warp $w$ processes $w, w{+}\text{NUM\_WARPS}, \dots$ — the block a given warp will touch next is exactly block_idx + NUM_WARPS, not block_idx + 1. The patch therefore looks ahead by the stride, which keeps each warp prefetching precisely the block it itself will consume. This is read straight from the block table the warp already holds, so the lookahead address costs one extra table read and a comparison.

// attention_kernels.cuh:258 — next physical block from the block table
const int64_t next_physical_block_number = (block_idx + NUM_WARPS) < end_block_idx ?
              block_table[block_idx + NUM_WARPS] : 0;

3.3 The key-cache prefetch (the one PTX line)

Inside the per-token vector loop, exactly one lane (lane 0, first vector) issues the prefetch, guarded so the last block — whose next_physical_block_number is the sentinel 0 — is never prefetched. The address is masked to a 16-byte boundary (the low nibble cleared by & 0xFFFFFFFFFFFFFFF0) as the bulk-prefetch instruction requires an aligned global address, and the size is a fixed 4096 bytes — one paged block's worth of key data.

// attention_kernels.cuh:281-287 — Prefetch Key Cache
if (j == 0 && lane == 0) {
    const cache_t* next_k_ptr = k_cache + next_physical_block_number * kv_block_stride +
                                kv_head_idx * kv_head_stride;
    if (next_physical_block_number != 0)
    asm volatile ("cp.async.bulk.prefetch.L2.global [%0], 4096;"
        ::"l"((unsigned long)next_k_ptr & 0xFFFFFFFFFFFFFFF0UL));
}

3.4 The value-cache prefetch (the symmetric second line)

The value cache is read in a separate loop (the second pass that computes the weighted sum of values). The patch repeats the same pattern there: recompute next_physical_block_number for the value loop's iteration, and on the first row of lane 0 issue an identical bulk prefetch against the value pointer. The structure is intentionally a mirror of the key prefetch, which is why the diff is so small.

// attention_kernels.cuh:425-431 — Prefetch Value Cache
if (i == 0 && lane == 0) {
    const cache_t* next_v_ptr = v_cache + next_physical_block_number * kv_block_stride +
                                kv_head_idx * kv_head_stride;
    if (next_physical_block_number != 0)
        asm volatile ("cp.async.bulk.prefetch.L2.global [%0], 4096;"
            ::"l"((unsigned long)next_v_ptr & 0xFFFFFFFFFFFFFFF0UL));
}

3.5 The cost model

The benefit per block iteration is the hidden latency $t_{hide} = \min(t_{load},\, t_{compute})$, where $t_{load} \approx t_{scoreboard}$ is the exposed global-load / cache-hierarchy latency of the late load — the prefetch can hide at most as much exposed global-load latency as there is compute to overlap with. The cost is one block-table read, one comparison, one address mask, and the launch of one asynchronous prefetch per block, all on a single lane. Net speedup appears when $t_{hide}$ accumulated over all blocks dominates the tiny issue overhead, which is exactly the memory-bound regime decode lives in. This is why the paper sees its largest wins on long sequences and decode-heavy workloads.

3.6 Three correctness details hidden in the two lines

Although the patch is tiny, each token in the PTX is load-bearing, and getting any of them wrong would either crash the kernel or silently corrupt the prefetch. Three details are worth spelling out because they are exactly the kind of thing a reproduction must preserve byte-for-byte.

4. Repo Analysis

The repository is striking in its minimalism. It contains no build system, no benchmark scripts, no Python — just a license, a short README, and one CUDA header. The entire intellectual contribution is shipped as a drop-in replacement for one vLLM file. We cloned it to the path below and confirmed its contents are exactly as described.

4.1 The diff against stock vLLM v0.10.2

The header is a near-verbatim copy of vLLM's own csrc/attention/attention_kernels.cuh with the prefetch additions woven in. The additive change is small — two next-block computations, two PTX prefetch sites (one for K at line 286, one for V at line 430), each behind a last-block guard. The table below maps the paper's mechanism onto exact code locations we read in the cloned header.

5. Claims vs Implementation

Each paper/README claim, checked against the cloned header. Performance claims are the paper's H20 numbers; our own measured H100 numbers are in Sections 6-7.

6. Our Reproduction & Measurement MEASURED

6.1 How much of this gain is the mechanism we bound?

The measured answer is concrete: ~0% of this paper's speedup is the passive inter-kernel activation residency we bound at <0.65% BOUND. The reasoning follows directly from the code and is confirmed by the counters. The carryover here is (1) software-issued — an explicit PTX instruction, not data left behind incidentally; (2) intra-kernel — the same kernel prefetches a later iteration of itself, so no second kernel is involved at all; and (3) on the KV-cache — a tensor that is read every decode step, not a transient activation. Each of those three properties places it outside the definition of the quantity our study measures. The measured signature matches at the kernel level: the cold-replay per-kernel delta (ncu, 1.01-1.02x) tracks a memory-stall reduction (long-scoreboard 9.2%->8.5%) with HBM read flat — not any cross-kernel cache hand-off. The headline 1.80x is the warm-loop e2e figure; the cold-replay-vs-warm-loop gap is reported separately (Section 7), not folded into the stall drop.

6.2 Method: what we build and run, on which GPU

• GPU: PACE H100 (sm_90, primary), optionally H200 — the closest available analogue to the paper's H20, since cp.async.bulk.prefetch requires Hopper. A100/L40S/RTX6000 cannot compile the kernel and are excluded.
• Two isolated trees: baseline = stock v0.10.2 header, treatment = the repo header, each editable-installed into its OWN conda env (vllm-prefetch-base / -patch) to avoid the two-tree editable-install drift where the wrong .so loads.
• Same everything else: both trees built from the v0.10.2 source (baseline self-reports 0.10.2; the patched build self-reports 0.10.3.dev), identical flags, VLLM_ATTENTION_BACKEND=XFORMERS, VLLM_USE_V1=0 so the legacy PagedAttention decode kernel (where the patch lives) is actually on the hot path. The only source difference is the prefetch .cuh.

6.3 The Nsight metrics that isolate L2-reuse from the other levers

We use the project's proven application-replay recipe (--replay-mode application --cache-control all --clock-control none --launch-count 40 --launch-skip 200, kernel-filtered to regex:paged_attention) to get cold-cache semantics without kernel-replay's host-RAM snapshot, which OOM'd this project at 67 GB on an H200.

6.4 E2E throughput (FACT) MEASURED

Random dataset, input-len 512, output-len 512, tok/s base -> patch. This reproduces the paper's e2e regime (up to ~1.97x on Hopper): the win is big at low batch / decode-heavy and vanishes by batch 32 (compute-bound). We ran H100; the paper's headline was H20.

6.5 Per-kernel attribution (FACT) MEASURED

Nsight Compute on paged_attention_v1 (cold-cache application replay, --cache-control all, input-len 1024 output-len 256, mean over ~40 launches). The isolated kernel speedup is only ~1.01-1.02x — far below the e2e 1.80x — because the baseline is already ~85% L2-resident and cold-cache replay under-represents the warm-loop async overlap. The mechanism is visible as a memory-stall drop, not an L2-residency jump.

7. L2 Attribution — Reading the Mechanism Off the Counters MEASURED

Profiling both builds on the same workload decides, from the hardware counters alone, which lever produces the speedup. There are four candidate explanations; the measured metric pattern distinguishes them. The H100 counters confirm lever (A): the per-kernel time delta tracks the long-scoreboard mem-stall drop (9.2%->8.5%), L2 read-hit barely moves (already ~85% resident), and HBM read is flat (2.01 MB) — so it is active latency-hiding of the prefetched KV, not residency or a traffic cut.

The decisive test was correlation, and the counters bear it out: the per-kernel time reduction tracks the drop in long-scoreboard stall percentage, so the speedup is latency-hiding of the prefetched KV — lever (A). HBM bytes staying flat (2.01 MB) prove the same KV reads were issued earlier / overlapped — the traffic was not removed, only its latency hidden; and by construction the patch adds no kernel launch and no shared-memory allocation, ruling out levers (C) and (B). One honest subtlety: the L2 read-hit rate barely rises (+0.2-0.4 pp) because the baseline KV is already ~85% L2-resident — so the prefetch's fingerprint is the stall drop, not a residency jump. This is the same separation logic our measurement study applies to bound passive residency — here it identifies an entirely different, active lever.

8. Related Work & Significance

8.1 Where this sits among attention kernels

Modern inference engines mostly run FlashAttention- and FlashInfer-style kernels that fuse the whole attention computation and already overlap global loads with compute through software pipelining and asynchronous copies into shared memory. This paper does not compete with those; it retrofits the same overlap idea onto vLLM's older XFORMERS PagedAttention path, which lacks block-table lookahead. Its contribution is therefore best read as an existence proof: even on a non-fused, paged kernel, a single bulk-prefetch instruction recovers most of the latency-hiding that a fully fused kernel gets for free. The mechanism — anticipate the next memory access and stage it one level up the hierarchy before it is needed — is the oldest idea in computer architecture, here expressed in two lines of PTX against the L2 cache.

8.2 The V0/V1 path caveat

A practical subtlety we must respect during reproduction: vLLM's default execution path has moved to the V1 engine, and the kernel this patch modifies — paged_attention_v1/v2 in csrc/attention/attention_kernels.cuh — is dispatched by the legacy XFORMERS backend on the V0 path. We therefore set VLLM_USE_V1=0 alongside VLLM_ATTENTION_BACKEND=XFORMERS so the patched kernel is genuinely on the hot path; otherwise the profiler's regex:paged_attention filter would capture nothing and the speedup would not appear. The page's reproduction section verifies the kernel is actually exercised before any number is trusted.

8.3 Why it matters for the survey

For a survey of KV-cache management, this paper is a clean datapoint on the kernel-level end of the spectrum: most cache-management work operates at the scheduler or block-allocation layer, but the bytes ultimately have to traverse the memory hierarchy, and this is what that traversal looks like when optimized by hand. For our measurement study specifically, it is the indispensable contrast that gives the <0.65% bound its meaning. A bound is only informative if there exist cases on the other side of it; this paper is precisely such a case — a real, large L2-mediated win — and showing that it lives in a different category (active, intra-kernel, KV-cache) is what tells the reader the bound is narrow and exact, not a sign that L2 never matters.

9. Separability From Our <0.65% Bound

This is why the paper belongs in our study as a sharpening case rather than a counterexample. Our bound is about one specific quantity: the value of letting a prior kernel's output activation sit in L2 long enough that the next kernel reads it from L2 instead of HBM, with no software doing anything to make that happen. The prefetch patch differs on every axis that defines that quantity.

10. How to Reproduce

The repo ships no build or benchmark scripts, so reproduction means: build stock and patched v0.10.2, then drive both with vLLM's own throughput benchmark and Nsight Compute. The two sbatch scripts below ran on PACE H100 (job 9948442); the build and the H100 run are complete and the numbers in Sections 6-7 are the result. (vLLM has since moved benchmark_throughput.py to `vllm bench throughput`.)

10.1 Install (per the repo README)

# 1. clone vLLM v0.10.2 (twice: baseline + patched)
git clone --depth=1 --branch v0.10.2 https://github.com/vllm-project/vllm.git baseline
git clone --depth=1 --branch v0.10.2 https://github.com/vllm-project/vllm.git patched

# 2. drop the repo header into the patched tree only
cp attention_kernels.cuh patched/csrc/attention/attention_kernels.cuh

# 3. build EACH tree into its OWN conda env (avoids two-tree .so drift)
export TORCH_CUDA_ARCH_LIST="9.0a"   # 'a' variant exposes the bulk-prefetch instruction
# module load cuda/12.9.1  (NOT 12.6.1 — ptxas SIGSEGVs on FA3 Hopper kernels)
pip install --no-cache-dir --no-build-isolation -e .
pip install xformers

10.2 Run e2e + attribution

export VLLM_ATTENTION_BACKEND=XFORMERS
export VLLM_USE_V1=0   # exercise the legacy PagedAttention decode kernel

# end-to-end throughput (decode-heavy: long output_len)
python benchmarks/benchmark_throughput.py --model Llama-3.1-8B \
  --dtype float16 --input-len 512 --output-len 512 --num-prompts 8 --output-json out.json

# kernel attribution (cold-cache, application-replay; NOT kernel-replay -> host-RAM OOM)
ncu --replay-mode application --cache-control all --clock-control none \
    --target-processes all -k "regex:paged_attention" \
    --launch-count 40 --launch-skip 200 --csv --page raw \
    --metrics gpu__time_duration.sum,dram__bytes_read.sum,\
lts__t_sector_op_read_hit_rate.pct,\
smsp__warp_issue_stalled_long_scoreboard_per_warp_active.pct,\
sm__throughput.avg.pct_of_peak_sustained_elapsed \
    python benchmarks/benchmark_throughput.py --model Llama-3.1-8B ...