Repository navigation
Stop declaring the paged KV cache mutated on read-only attention ops - #397
Merged
Merged
Conversation
There was a problem hiding this comment.
🔵 Needs a closer look
Compiler aliasing and HIP graph-capture lifetime changes warrant final human review despite the reported multi-architecture validation.
No open issues 🎉
What changed in this PR
Fixes compiled read-only paged attention by correcting cache mutation declarations and hardening AITER decode.
Changes:
- Removes false K/V mutation declarations and renames ops to invalidate cached graphs.
- Fixes AITER scale storage and layout handling, and makes default decode outputs contiguous.
- Adds regression tests and documents remaining append limitations.
| File | Description |
|---|---|
| tests/rocm/test_torch_compile.py | Tests cache views, append ordering, and graph re-recording. |
| tests/rocm/test_batch_decode_aiter.py | Tests query and cache layouts across decode routes. |
| README.md | Documents append risks and module-buffer advice. |
| flashinfer/rocm/prefill.py | Corrects and renames the read-only paged op. |
| flashinfer/rocm/decode.py | Corrects mutation declarations, renames ops, and allocates contiguous outputs. |
| csrc/rocm/batch_decode_aiter.cu | Uses static device scale storage and validates supported layouts. |
Review effort: Balanced
PA v1 takes k_scale/v_scale pointers (the pinned AITER reads them only
for fp8 caches), so the shim passed a 1.0f tensor it allocated lazily on
the first call and cached in a thread_local. Under
torch.compile(mode="reduce-overhead") that first call can be a
cudagraph-tree warm-up, which allocates from the graph pool; the cached
tensor is then a live allocation the pool does not track, and torch
refuses to continue:
RuntimeError: Detected 1 tensor(s) in the cudagraph pool not tracked
as outputs. All live allocations must be tracked for correctness.
The scale now lives in static device memory (a file-local __device__
float, address via hipGetSymbolAddress, cached per device), which no
allocator owns; a failed lookup reports HIP's error and clears it.
PA v1 also reads V with K's strides, assumes a dense head_dim, and
takes only q's token stride, and the shim checked none of it: a V laid
out differently from K, or a q whose heads are not densely packed (a
head_dim slice of a wider tensor), gave wrong attention with no error,
and a cache on another device reached the kernel as a foreign pointer.
The shim now requires K and V to share page shape and strides (page
counts may differ; only indexed pages are read), q's device and a dense
head_dim, as the fa2 shims require of strides; a q with spread-out
heads is copied to a packed one, a B x H x D copy, rather than
rejected. The unit-scale lookup is keyed by the current device, which
is what hipGetSymbolAddress resolves against.
run() allocates its default out in contiguous format rather than like
q. Every decode kernel writes o densely, so for a head-major q the
q-shaped out made AITER refuse the call and fa2 write its rows in the
wrong order, silently. torch.empty_like(q, memory_format=
contiguous_format) costs 1.63 us against empty_like's 1.46 (torch.empty
with shape, dtype and device measured 3.5 us on gfx942 and showed up as
+2.8 us per flag-on decode call, so it was not used).
Tests: test_torch_compile_aiter_decode_reduce_overhead_rerecords (AITER
decode over a module-buffer cache under reduce-overhead: warm-up as the
process's first AITER call, re-records at new shapes, replay against an
fa2 reference taken before compiling, cache unchanged; its q is a head
slice, so the q copy runs under capture),
test_batch_decode_aiter_kv_and_q_layouts (a V with HND strides raises;
a head-sliced q, a q with only its head_dim strided and a shorter V
pool match fa2) and test_batch_decode_aiter_batched_q_layouts (batch 3;
fa2, tensor-core and AITER decode each read a q sliced from fused QKV
and a head-major q into a contiguous out that matches fa2 on a packed
q) fail on the old code (the RuntimeError above; DID NOT RAISE; the
non-contiguous-output refusal) and pass here, at this commit and with
the next one, on gfx942 and gfx950. Both scale q by 8 so the softmax is
peaked and a misread q moves the output beyond the tolerance.
Flag off, per call, min of 3 interleaved rounds, base -> this commit
and the next, which is unused with the flag off (us):
gfx942 gfx950
AITER decode 27.98 -> 27.70 20.46 -> 20.28
fa2 decode 38.43 -> 38.14 32.46 -> 32.68
tensor-core decode 88.26 -> 88.17 42.96 -> 43.03
fa2 paged prefill 339.60 -> 338.88 188.53 -> 188.38
Neutral: every delta is under 0.8 us and both signs appear.
Co-Authored-By: Claude <noreply@anthropic.com>
fa2 batch decode, AITER decode and the fa2 paged-prefill run op only
read paged_k_cache and paged_v_cache, but declared both mutated. With
FLASHINFER_USE_TORCH_CUSTOM_OPS=1, torch 2.12 mishandles declared-
mutated views: it rebuilds a view from the wrong storage offset and
copies the result back, so compiled attention read the wrong K/V and
overwrote V, or faulted the GPU. Max |compiled - eager|, identical on
gfx942 and gfx950 for all three ops:
K/V passed in dynamic=False default dynamic=True
one interleaved [pages, 2] tensor 0 0 1.05-1.26
(kv[:,0], kv[:,1]) 0 0 0
(kv[0], kv[1]) of [2, pages] 6.34 6.34 0
separately allocated 0 0 0
layer 1 of a two-layer pool memory fault in every mode
With neither cache declared every cell is exact for every op, on both
arches. Rejected alternatives, each measured: disabling
auto_functionalized_v2 fixes only the dynamic=True column; declaring K
alone on AITER decode still faults on layer 1 of a pool.
The ops are renamed (_run_kvro, _paged_run_kvro). Inductor's on-disk
caches do not key on alias schemas, so a cache filled by the old ops
kept replaying the old graph after the fix (6.34 again with a shared
TORCHINDUCTOR_CACHE_DIR, both arches); under the new names it misses
and recompiles (0).
reduce-overhead: the fa2 ops write their workspaces, so CUDA graphs
stay skipped for them as before. AITER decode, which was already
captured for a cache held as a module buffer, is now captured for one
passed as an argument too (the previous commit makes capture safe). As
a module buffer, a 1 GiB cache is neutral (min of 3 rounds:
67.87 -> 67.86 us gfx942, 42.82 -> 42.51 us gfx950); a cache passed as
an argument is copied on every replay, as torch does for any non-static
input (77-80 -> 603 us gfx942, 40-43 -> 506 us gfx950). The README now
tells reduce-overhead users to register the cache as a module buffer.
The jit_args variants keep their declarations: they take *args, which
schema inference rejects, so they never register; a comment says to
drop the caches if they ever do. Append genuinely writes the cache and
keeps both declarations, so it keeps the bug, including the layer-1
fault; the README now documents that and says to allocate K and V
separately.
Tests: test_torch_compile_kv_cache_views runs fa2, tensor-core and
AITER decode and fa2 paged prefill under fullgraph over [2, pages]
views (default compile), an interleaved tensor (dynamic=True) and layer
1 of a pool (both), takes the reference before compiling, checks every
cache tensor bit-identical after, and reports each case. All 16
route x case cells fail with the old declarations and pass here, on
gfx942 and gfx950. test_torch_compile_append_and_attention_order
compiles, under fullgraph, an append of half the positions before and
after fa2 decode, AITER decode and fa2 prefill, over separate caches
and over views of one tensor under dynamic=True; the dropped
declaration no longer orders these, and they pass before and after.
Both turn off Inductor's on-disk caches, so an A/B cannot replay the
other checkout's graph.
Flag on, per call, min of 3 interleaved rounds, base -> this commit
and the previous (us):
gfx942 gfx950
eager AITER decode 41.95 -> 41.33 26.71 -> 25.19
eager fa2 decode 41.42 -> 42.15 32.55 -> 32.58
eager tensor-core decode 88.57 -> 88.24 42.91 -> 43.06
eager fa2 paged prefill 339.58 -> 339.31 188.57 -> 188.55
compiled AITER decode 69.29 -> 70.67 39.33 -> 35.64
compiled fa2 decode 71.78 -> 68.85 36.62 -> 35.14
compiled tensor-core decode 89.59 -> 88.45 45.70 -> 43.48
compiled fa2 paged prefill 340.07 -> 340.08 188.47 -> 188.66
No regression: every slowdown is under 1.4 us and none exceeds 0.2 us
on both arches at once. The compiled gfx950 gains come from a single
round and are not claimed. Flag off the declarations and names are
unused (previous commit's table).
Co-Authored-By: Claude <noreply@anthropic.com>
demandal25
force-pushed
the
readonly-kv-not-mutated
branch
from
October 21, 2026 19:48
8370664 to
fcb2937
Compare
Security scan found no critical or high findings on lines changed by this pull requestView the security scan run for complete reports. “New” means the finding is located in a changed file on an added or modified line; it is not a base-versus-head findings delta. |
6 tasks done
demandal25
added a commit
that referenced
this pull request
Oct 22, 2026
## Summary The README now states what was validated at the release head (the full suite on both architectures, in the supported image), what `torch.compile` does per op, and the one `torch.compile` known issue left after #397, with its workaround. The ROCm example scripts and tutorial notebook are removed: they duplicated the README's quick start and upstream's documentation, and nothing tested them. The README is also cut from 397 to 292 lines: it keeps quick start, usage, the supported configuration, the full support matrix and what upstream has that this does not, and moves the `torch.compile` and environment-variable reference into `docs/rocm/`. ## What changed - **`README.md`**, *Supported hardware and toolchain*: the paragraph saying gfx942 had only run the prefill, arch-caps and AITER-routing suites is replaced by the full-suite result on both arches, anchored to the `v0.6.18+amd.2` tag. - **`docs/rocm/torch-compile.md`** (new; the README keeps a summary and a one-line warning): adds an append row (`append_paged_kv_cache` and `append_paged_mla_kv_cache` compile with the flag on, under `fullgraph` too except `backend="aiter"` append), adds the fp8-quantizing norms to the ops the flag cannot register, and replaces "raises a clear error" with a measured per-op table (flag off, flag on, `fullgraph=True`), a footnote that AITER decode with `return_lse=True` needs `backend="fa2"` under `fullgraph` (#388), the host-side flag-on dispatch cost, and what the flag does not cover: `plan()` stays outside any compiled region, and sampling ops taking a `generator`, `jit_args` custom variants and AITER's opt-in norm/RoPE/activation shims are not registered. With the flag off the error is usually torch's `RuntimeError when making fake tensor call`. - **`docs/rocm/torch-compile.md`**, *Known issue*: merges #397's append note into one rule: give every tensor an op writes its own storage, because PyTorch 2.12 can miscompile a write into a view of a larger tensor under every `dynamic` setting (tables below). That includes append's own single `[pages, 2, ...]` cache, which the wrapper splits into views, so append users are told to pass two separate caches. Four review rounds each found a layout or mode the previous statement of torch's trigger missed, so the README states the rule rather than the trigger. For `reduce-overhead` it keeps only #397's measured advice: register the KV cache as a module buffer, since AITER decode copies an argument cache on every replay. The paragraph this branch carried before #397, which scoped the issue to `dynamic=True`, is dropped. - **`README.md`**, *Basic usage*: the pointer to `examples/` is now a pointer to upstream's `plan()`/`run()` documentation, flagged as tracking upstream's latest release, plus the ROCm list of refused CUDA-only arguments. - **`README.md`**, the rest (third commit): quick start keeps its commands and links `CONTRIBUTING.md` for the container gotchas; the two gfx950 routing paragraphs become a pointer to the per-op notes in `docs/rocm/backends.md`, which already hold them; benchmark caveats are left to `benchmarks/rocm/README.md`; the env-var table keeps the four variables a user typically sets. - **`docs/rocm/environment.md`** (new): the full runtime env-var table and the runtime-helper snippet, moved verbatim. `CONTRIBUTING.md`'s link to the README table now points here. The doc-ownership table in `.claude/skills/pr-workflow/SKILL.md` names both new pages. - **`examples/`**: removes `single_prefill_example.py`, `batch_prefill_example.py`, `batch_decode_example.py`, `amd_flashinfer_rocm_tutorial.ipynb` and `run_jupyter_server.sh`. All five were fork-added (none exists at `v0.6.18`), so the additive-only rule is untouched; upstream's `examples/pytorch/` stays. ## Benchmark results None: documentation and deleted example scripts only, so no kernel or dispatch path changes. ## Test plan Full suite at e1054f3, the release head (this PR is docs and deletions only), in `docker/Dockerfile.rocm`, one allocated card per arch (bus id checked on gfx950): | | gfx942 (MI300X) | gfx950 (MI355X) | |---|---|---| | `-n auto --reruns 2 -m "not slow"` | 33867 passed, 4024 skipped, 0 failed | 33875 passed, 4016 skipped, 0 failed¹ | | `-n 1 --reruns 2 -m slow` | 72 passed | 72 passed | | Wall time | 87 min, warm JIT cache | 89 min, cold JIT cache | ¹ Four git-dependent tests (`test_amd_coverage`'s two manifest tests, `test_redistribution_licences`' sdist and wheel exclusions) failed on gfx950 because that copy of the tree had no `.git`. Both files pass, 125 of 125, in a git checkout of the same commit, and are counted as passed. No other test needed a rerun on either arch. Flag-on dispatch at afb6440, single prefill qo16/kv128, minimum of 3 interleaved rounds on an idle allocated card: | | Flag off | Flag on | Delta | fa2 control delta | |---|---|---|---|---| | gfx942 | 21.61 µs | 42.42 µs | +20.8 µs | +18.7 µs | | gfx950 | 12.65 µs | 23.38 µs | +10.7 µs | +8.2 µs | - [x] `torch.compile` matrix in bf16 at e1054f3 on gfx942 and gfx950, identical cell for cell (fa2 and aiter; single/paged/ragged prefill, decode, MLA, MoE, append, MLA append; flag 0/1; default and `fullgraph=True`), every compiled result compared to eager: all match. The README table is that result. - [x] Writes into views at e1054f3, max \|compiled − eager\|, identical on gfx942 and gfx950: | Written arguments | `dynamic=False` | default | `dynamic=True` | |---|---|---|---| | Append into `kv[0]`, `kv[1]` passed in | 4.594 | 4.594 | – | | `fused_add_rmsnorm(pool[1,0], pool[1,1])` passed in | 8.25 | 8.25 | – | | Append into one `[pages, 2, ...]` cache | 0 | 0 | 5.484 | | Append into separate K and V | 0 | 0 | 0 | One default-compiled function called at T = 32/48/64/80, once a size goes dynamic: | Written arguments | T=32 | 48 | 64 | 80 | |---|---|---|---|---| | `fused_add_rmsnorm(pool[1,0], pool[1,1])`, pool shaped by T | 0 | 12.1 | 12.4 | 13 | | `rmsnorm(out=o[1])`, `o` shaped by T | 0 | 8.94 | 12.5 | 9.44 | | Either, on a buffer allocated once at the largest T | 0 | 0 | 0 | 0 | In-place RoPE on `qkv.split` views and MLA append on `torch.split(buf, [512, 64])` were exact with `dynamic=None`, `False` and `True` on both arches. - [x] No remaining reference to the removed files (`git grep` for the file names, `ipynb`, `jupyter`; the matches left are upstream's). - [x] No Python or C++ changes, so no tests beyond the suite above apply. - [x] `/code-review xhigh` on `origin/amd-integration...HEAD` - [x] `pre-commit run -a` --------- Co-authored-by: Claude <noreply@anthropic.com>
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Summary
With
FLASHINFER_USE_TORCH_CUSTOM_OPS=1,torch.compilegave wrong attention, overwrote V, or faulted the GPU when K and V were views: two halves of one[2, pages, ...]tensor, one interleaved tensor underdynamic=True, or a layer of a multi-layer pool in any mode. The read-only attention ops declared the paged KV cache mutated, which exposes a torch 2.12 bug in how declared-mutated views are rebuilt. They now declare only what they write, under new op names so torch's on-disk compile cache cannot replay the old graphs. AITER decode shim defects found on the way are fixed first: an untracked allocation that brokemode="reduce-overhead", and K/V/q layouts the kernel cannot address being read wrongly without an error.What changed
csrc/rocm/batch_decode_aiter.cu(first commit): the unit k/v scale passed to PA v1 lives in static device memory instead of a lazily allocated tensor, which landed untracked in torch's CUDA-graph pool. K and V must share page shape, strides and q's device and have a dense head_dim; a q with spread-out heads is copied to a packed one. Before, a V laid out differently from K or a head-sliced q gave wrong attention with no error.flashinfer/rocm/decode.py(first commit):run()allocates its defaultoutcontiguous rather than like q. Every decode kernel writes o densely, so for a head-major q AITER refused the call and fa2 wrote its rows in the wrong order without an error.flashinfer/rocm/decode.py: fa2 batch decode and AITER decode no longer declarepaged_k_cache/paged_v_cache, and register as{uri}_run_kvro.flashinfer/rocm/prefill.py: the same for the fa2paged_runop, now{uri}_paged_run_kvro. AITER paged prefill already declared onlyoandlse. Thejit_argsvariants keep their declarations: they take*args, which schema inference rejects, so they never register.README.md: a known-issue note undertorch.compilefor append over cache views, and the module-buffer advice formode="reduce-overhead".tests/rocm/test_batch_decode_aiter.py:test_batch_decode_aiter_kv_and_q_layouts(a V with HND strides raises; a head-sliced q, a head_dim-strided q and a shorter V pool match fa2);test_batch_decode_aiter_batched_q_layouts(batch 3; fa2, tensor-core and AITER decode each read a q sliced from fused QKV and a head-major q into a contiguousoutmatching fa2 on a packed q).tests/rocm/test_torch_compile.py:test_torch_compile_aiter_decode_reduce_overhead_rerecords(first AITER call inside a reduce-overhead warm-up, with a head-sliced q, re-records, replay against an fa2 reference taken before compiling);test_torch_compile_kv_cache_views(fa2, tensor-core and AITER decode, fa2 paged prefill;fullgraph;[2, pages]views, an interleaved tensor underdynamic=True, layer 1 of a pool; reference before compiling, every cache tensor bit-identical after, each case reported);test_torch_compile_append_and_attention_order(append of half the positions before and after each route in onefullgraphgraph, over separate caches and over views of one tensor underdynamic=True). The new compile tests turn off Inductor's on-disk caches.Architecture / design notes
Max |compiled − eager| before this change, identical on gfx942 and gfx950 for fa2 decode, AITER decode and fa2 paged prefill:
dynamic=Falsedynamic=True[pages, 2, ...]tensor(kv[:, 0], kv[:, 1])(kv[0], kv[1])of a[2, pages, ...]tensorAfter it, every cell is exact for every op on both arches. Two alternatives were measured and rejected:
enable_auto_functionalized_v2 = Falsefixes only thedynamic=Truecolumn, and declaring K alone on AITER decode still faults on layer 1 of a pool.Why the ops are renamed. Inductor's FX-graph and AOTAutograd caches do not key on an op's alias schema. With a
TORCHINDUCTOR_CACHE_DIRfilled by the released ops, the fixed code replayed the old graph and was wrong again (6.34, both arches); under the new names it misses the cache and is exact. The default cache directory persists across runs, so without the rename an upgrade would keep the bug for anyone who had compiled with the flag on.mode="reduce-overhead". The fa2 ops write their workspaces, so CUDA graphs stay skipped for them as before. AITER decode is now captured, which the shim fix makes safe. With a 1 GiB cache held as a module buffer it is neutral (67.87 → 67.86 µs gfx942, 42.82 → 42.51 µs gfx950); a cache passed as an argument is copied on every replay, as torch does for any non-static input (77–80 → 603 µs gfx942, 40–43 → 506 µs gfx950). The README now tells reduce-overhead users to register the cache as a module buffer.Append is unchanged. It genuinely writes the cache and keeps both declarations, so it keeps the bug, including the layer-1 fault; the README now documents it and says to allocate K and V separately.
Benchmark results
Per call, minimum of 3 interleaved rounds on one idle card per arch, base caa2d8e → both commits, µs:
No regression: every slowdown is under 1.4 µs, and none exceeds 0.2 µs on both arches at once. The compiled gfx950 gains come from a single round and are not claimed.
Test plan
kv_cache_viewsfails in all 16 route × case cells on the old code and passes with both commits; the reduce-overhead test fails on the old shim (untracked pool tensor) and passes from the first commit on; the two layouts tests fail on the old shim (did not raise; non-contiguous output refused) and pass from the first commit on;append_and_attention_orderpasses throughout. gfx942 and gfx950.TORCHINDUCTOR_CACHE_DIRfilled by the released ops, then read by this change: exact under the new names, both arches.test_torch_compile,test_batch_decode_aiter,test_batch_decode_kernels,test_batch_decode_vllm,test_batch_decode_speculative,test_decode_guards,test_cascadeandtest_deprecated_forward_api, flag off and on: gfx942 2758 / 2758 passed; gfx950 2758 / 2758 passed. Earlier in the branch, 15 decode, prefill and compile files passed on both arches (9050 / 9048)./code-review xhighonorigin/amd-integration...HEAD, nine rounds; findings applied in the commits or declined.pre-commit run -a