Skip to content

Stop declaring the paged KV cache mutated on read-only attention ops - #397

Merged
demandal25 merged 2 commits into
amd-integrationfrom
readonly-kv-not-mutated
Oct 21, 2026
Merged

demandal25 merged 2 commits into
amd-integrationfrom
readonly-kv-not-mutated

Conversation

@demandal25

Copy link
Copy Markdown
Collaborator

Summary

With FLASHINFER_USE_TORCH_CUSTOM_OPS=1, torch.compile gave wrong attention, overwrote V, or faulted the GPU when K and V were views: two halves of one [2, pages, ...] tensor, one interleaved tensor under dynamic=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 broke mode="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 default out contiguous 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 declare paged_k_cache/paged_v_cache, and register as {uri}_run_kvro.
  • flashinfer/rocm/prefill.py: the same for the fa2 paged_run op, now {uri}_paged_run_kvro. AITER paged prefill already declared only o and lse. The jit_args variants keep their declarations: they take *args, which schema inference rejects, so they never register.
  • README.md: a known-issue note under torch.compile for append over cache views, and the module-buffer advice for mode="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 contiguous out matching 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 under dynamic=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 one fullgraph graph, over separate caches and over views of one tensor under dynamic=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:

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 a [2, pages, ...] tensor 6.34 6.34 0
Separately allocated 0 0 0
Layer 1 of a two-layer pool GPU memory fault GPU memory fault GPU memory fault

After it, every cell is exact for every op on both arches. Two alternatives were measured and rejected: enable_auto_functionalized_v2 = False fixes only the dynamic=True column, 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_DIR filled 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:

Call gfx942 gfx950
Flag off, AITER decode 27.98 → 27.70 20.46 → 20.28
Flag off, fa2 decode 38.43 → 38.14 32.46 → 32.68
Flag off, tensor-core decode 88.26 → 88.17 42.96 → 43.03
Flag off, fa2 paged prefill 339.60 → 338.88 188.53 → 188.38
Flag on, eager AITER decode 41.95 → 41.33 26.71 → 25.19
Flag on, eager fa2 decode 41.42 → 42.15 32.55 → 32.58
Flag on, eager tensor-core decode 88.57 → 88.24 42.91 → 43.06
Flag on, eager fa2 paged prefill 339.58 → 339.31 188.57 → 188.55
Flag on, compiled AITER decode 69.29 → 70.67 39.33 → 35.64
Flag on, compiled fa2 decode 71.78 → 68.85 36.62 → 35.14
Flag on, compiled tensor-core decode 89.59 → 88.45 45.70 → 43.48
Flag on, compiled fa2 paged prefill 340.07 → 340.08 188.47 → 188.66

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

  • A/B on caa2d8e's ops and shim, the first commit alone, and both commits: kv_cache_views fails 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_order passes throughout. gfx942 and gfx950.
  • Layout × compile-mode matrix above, one process per cell, before and after: the first four layouts on both arches, the layer-1 row on gfx942; the new test covers layer 1 on both.
  • Upgrade: a shared TORCHINDUCTOR_CACHE_DIR filled 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_cascade and test_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 xhigh on origin/amd-integration...HEAD, nine rounds; findings applied in the commits or declined.
  • pre-commit run -a

Copilot AI balanced review requested due to automatic review settings October 21, 2026 13:24

Copilot AI left a comment

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

🔵 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

demandal25 and others added 2 commits October 21, 2026 15:41
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>
Copilot AI balanced review requested due to automatic review settings October 21, 2026 19:48
@demandal25
demandal25 force-pushed the readonly-kv-not-mutated branch from 8370664 to fcb2937 Compare October 21, 2026 19:48

Copilot AI left a comment

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

🔵 Needs a closer look

GPU aliasing and graph-capture lifetime changes warrant hardware-backed human review despite the extensive reported validation.

No open issues 🎉

Review effort: Balanced

@github-actions

Copy link
Copy Markdown
Security scan found no critical or high findings on lines changed by this pull request

View 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.

@demandal25
demandal25 merged commit e1054f3 into amd-integration Oct 21, 2026
14 checks passed
@demandal25
demandal25 deleted the readonly-kv-not-mutated branch October 21, 2026 20:51
@demandal25 demandal25 mentioned this pull request Oct 22, 2026
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>
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants