Skip to content
Closed
Show file tree
Hide file tree
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
2 changes: 1 addition & 1 deletion CONTRIBUTING.md
Original file line number Diff line number Diff line change
Expand Up @@ -40,7 +40,7 @@ those commits.
**Foreign HIP** (leapdragon, a17t, anyone who is not this lab) keeps
**their** Author **and** Committer. Same rule as extras
[PR #1](https://github.com/opengfx1030/vllm-rdna/pull/1) (leapdragon
`rdna_ar`, now dest extras @ `a4060647` squash / tip `2a5e893`; pick unique Aron Hsiao commits,
`rdna_ar`, now dest extras @ `a4060647` squash / tip `700753d9`; pick unique Aron Hsiao commits,
not the GitHub squash). Do not rewrite those as Blivion or Cursor.

| Rule | Meaning |
Expand Down
2 changes: 1 addition & 1 deletion docs/ARCHITECTURE.md
Original file line number Diff line number Diff line change
Expand Up @@ -200,7 +200,7 @@ One V1 id per tile (`HIPPIHX_V1_OP_*`). Serve wraps as
`torch.ops.hippihx.<op>` — never a second Triton path in this library.
Python mirror: `hippihx.v1` (`V1OpId`, `ABI_REVISION`). Caps include
optional activation `dtype` (revision **2**). Qualnames `attention.*`
(revision **3**; ids unchanged). Dest extras tip `2a5e89368272` still
(revision **3**; ids unchanged). Dest extras tip `700753d9add5` still
has no `torch.ops.hippihx.*` rewire — see [`BACKPORT.md`](BACKPORT.md).
Do not edit `opengfx1030/vllm-rdna` from this tree.

Expand Down
33 changes: 17 additions & 16 deletions docs/BACKPORT.md
Original file line number Diff line number Diff line change
@@ -1,7 +1,7 @@
# extras → hippihx backport review

Lock: [`opengfx1030/vllm-rdna`](https://github.com/opengfx1030/vllm-rdna)
`rdna_extras` @ `2a5e89368272` (2026-09-24 21:52 UTC). Dest **default
`rdna_extras` @ `700753d9add5` (2026-09-25 05:48 UTC). Dest **default
branch is `rdna_extras`**, not `main` (`c00091e02670` upstream vLLM — no
dest HIP). **No** `torch.ops.hippihx.*`. Observe only: no `.cu` dump, no
tok/s, no extras maxdiff.
Expand Down Expand Up @@ -32,28 +32,29 @@ until a body migrates and extras binds it.
| `3bddd3c9a5d1` | Flash-Next production launcher (`serve_gfx1030_flashnext.sh`) | extras scripts. PIECEWISE + seq cap 6. leapdragon `VLLM_RDNA_AR` unset (still **0**). Zoo still no Finegrained. |
| `50120e13b468` | HC `_contig()` capture-safe per-shape cache | extras product Python. Gate still **off**. Remaining same-shape clobber. No V1 op. |
| `5c4ab9891910` | dest-reverted gfx1030 wvSplitK decode port | extras dense GEMM. Kernel asserts on gfx1030 under capture. Keeps `gemv_f16_rdna2` for decode `M<=8`. **No zoo tile.** Do not reintroduce `gemv_f16`. |
| `e45dd5cb2de8` | QSA `CircularBufferManager` empty-ring prefix hits | extras v1 core. Observe `attention/qsa_indexer`. Empty compression-group ring is valid. Do not copy tok/s. |
| `e45dd5cb2de8` | QSA `CircularBufferManager` empty-ring prefix hits | extras v1 core. Observe `attention/qsa_indexer`. Empty compression-group ring is valid. |
| `c59a23f625e0` | recovered extras probes / benches / HC Triton WIP | extras ops. HC Triton + graph-keepalive diagnostics **not** on the serve path. Do not dump. |
| `31003ff` / `ae5c7edc` / `37345ee9` / `aaa4775a` / `0dd38115` / `657bdba8` / `9c602943` / `609c9c0d` | Flash-Next launcher serve knobs | extras scripts. Vision-on, then `609c9c0d` FULL_AND_PIECEWISE (ROCm executes as PIECEWISE). Stay extras. Do not copy tok/s. |
| `31003ff` / `ae5c7edc` / `37345ee9` / `aaa4775a` / `0dd38115` / `657bdba8` / `9c602943` / `609c9c0d` | Flash-Next launcher serve knobs | extras scripts. Vision-on, then `609c9c0d` FULL_AND_PIECEWISE (ROCm executes as PIECEWISE). Stay extras. |
| `741e5bc31ae5` | mamba spec-decode tables index by `req_idx` | extras worker (upstream vLLM #55506 port; Author Karl0007). Persistent per-request-slot tables. V1 `req_idx == batch_idx`. Stay extras. |
| `3b59ee16e553` | extras PR **#13** squash (T44b `rdna_ar` VRAM flags + wedge) | Later `comm.pcie`. Still opt-in (`VLLM_RDNA_AR=0`). Dest extras `MAX_KB` default **64**. Zoo lock **512**. Do **not** pick squash (Cursor rewrite). |
| `dbb1e7764aba` | extras PR **#14** merge (V620 Triton MoE JSON / ROCR amdsmi / PLE fp8) | Stay extras. HIP MoE ignores the JSON. Cursor rewrite — do not pick. |
| `4425834a26ac` | ROCm platform/worker startup (amdsmi import, hip fallback, `torch.cuda.init`) | Stay extras. Serve/platform only. No HIP body. |
| `3d6df9ed617a` | extras `exl3_dot2_*` `__HIP__RDNA__` guard for docker multi-arch | Stay extras. Zoo still one `--offload-arch` per fatbin. gfx1150 / gfx12xx not dest. Do not dump. |
| `ed94e3f3d299` | Qwen4Exp MTP proposer allowlist + skinny `w2_zp` | Stay extras. MTP still not dest. No HIP body. |
| `b33f9b66eb2b` | `get_device_name` torch fallback when amdsmi has no handles | Stay extras. Platform only. |
| `f3dd65fa7063` | extras PR **#15** merge (QSA live-context bound + folded **#12** PLE/MTP/graph-redirect) | Stay extras. Python/Triton/serve. No HIP body. Foreign George + Codex. Do not pick merge. Do not copy tok/s. |
| `e131562936f9` | extras PR **#17** merge (resident W4A16 MoE / Mamba #55450 retire) | Stay extras. ATen HIP (`moe_resident_decode.cu`). `VLLM_RDNA_MOE_RESIDENT*` default **off**. Dest `VLLM_RDNA_AR` still **0**. Foreign George + Codex. Do not pick merge. Do not dump. Do not copy tok/s. |
| `48c56efbe80b` | pin `AttentionConfig.backend = RDNA_ATTN` when `VLLM_USE_RDNA2_FA=1` | Stay extras. Python platform (`vllm/platforms/rocm.py`). API-server env; workers inherit config. Dest `VLLM_RDNA_AR` still **0**. Unique Blivion. No HIP. Do not copy tok/s. |
| `68a635ed8d78` | extras **#20** capture + dest-integrated **#21** FULL+PIECEWISE | Stay extras. Python/serve. Keeps FULL decode graphs; piecewise for mixed/prefill. Foreign Waldecir + Cursor. Do not pick merges. Do not copy tok/s. Dest `VLLM_RDNA_AR` still **0**. Zoo does not own graph mode. |
| `2a5e89368272` | extras PR **#19** merge (MTP unquantized-weight detect) | Stay extras. Python only (`qwen3_5_mtp.py`). MTP still not dest. Foreign a17t + opencode. Do not pick merge. Do not copy tok/s. Dest `VLLM_RDNA_AR` still **0**. |
| `f3dd65fa7063` | extras PR **#15** merge (QSA live-context bound + folded **#12** PLE/MTP/graph-redirect) | Stay extras. Python/Triton/serve. Foreign George + Codex. Do not pick. |
| `e131562936f9` | extras PR **#17** merge (resident W4A16 MoE / Mamba #55450 retire) | Stay extras. ATen HIP. `VLLM_RDNA_MOE_RESIDENT*` default **off**. Foreign George + Codex. Do not pick. |
| `48c56efbe80b` | pin `AttentionConfig.backend = RDNA_ATTN` when `VLLM_USE_RDNA2_FA=1` | Stay extras. Python platform. Unique Blivion. |
| `68a635ed8d78` | extras **#20** capture + dest-integrated **#21** FULL+PIECEWISE | Stay extras. Keeps FULL decode graphs; piecewise mixed/prefill. Foreign Waldecir + Cursor. Zoo does not own graph mode. |
| `2a5e89368272` | extras PR **#19** merge (MTP unquantized-weight detect) | Stay extras. Python only. MTP still not dest. Foreign a17t + opencode. Do not pick. |
| `700753d9add5` | MTP draft unbreak (amdsmi shutdown swallow + local-argmax) | Stay extras. Python platform + MTP model. MTP still not dest. Unique Blivion. |

Other dest-fixed ISA (keep in tile locks, not a dump): GDN prefill `o`
`i_t_local`; causal_conv FIR pre-shift then shift; ConfigH
(`K_STEP=64`) reverted. Live dest bugs: [Dest extras defects](#dest-extras-defects).

Open extras PRs **#2** (GLM Later), **#18**. Draft **#9/#10**
sdot — skip. Closed **#3/#16** unmerged, **#21** dest-integrated
Open extras PRs **#2** (GLM Later), **#18**, **#22**, **#24**. Draft **#9/#10**
sdot, **#23** PCIe P2P KV — skip. Closed **#3/#16** unmerged, **#21** dest-integrated
unmerged, **#11** dest-reverted wvSplitK (**no zoo tile**), **#12**
superseded by dest **#15**. Merged **#13/#14/#15/#17/#19/#20** dest
*presence*, observe, do not pick.
Expand All @@ -70,15 +71,15 @@ superseded by dest **#15**. Merged **#13/#14/#15/#17/#19/#20** dest
| leapdragon `rdna_ar` + PIX helpers + dest T44b (`3b59ee16`) | **Later** `comm.pcie`. Still opt-in. Dest extras `MAX_KB` **64**; zoo **512**. Do not pick squash. |
| PR **#2** GLM-5.3 KDA/DSA | **Later** `kda_scan` / `dsa_nope` / `qsa_indexer`. Do not name tiles `glm5_*`. |
| GDN arenas, `rdna2_graph_keepalive.cuh`, breakable cudagraphs, Hybrid W4 gfx10, Flash-Next HC/QSA/PLE/M-RoPE HIP, skinny GEMM / `gemv_f16_rdna2`, PLE schema, W4 MoE oracle, `new_zeros`/`zeros_like`, V1 FULL→PIECEWISE, vLLM custom AR, `bench_report.py`, seq cap 6, Flash-Next launcher knobs, HC `_contig()` cache, QSA Triton bounds, mamba spec-decode, T44b wedge, V620 MoE JSON / amdsmi / PLE, ROCm platform init, extras EXL3 docker arch-guard, Qwen4Exp MTP, recovered extras probes, resident MoE, FA RDNA_ATTN pin, `qwen4_exp/**` | **Stay extras.** Product gates default off (S6 revert). No new V1 op until dest locks a class. |
| PR **#3** a17t / explore **#9/#10** sdot / closed **#12** / closed **#16** / PR **#18** | **Skip.** Second W4 family, not dest, serve-only. **#16** unmerged. **#18** GPTQ `BLOCK_KN_SIZE` 256. |
| PR **#3** a17t / explore **#9/#10** sdot / closed **#12** / closed **#16** / PR **#18** / PR **#22** / PR **#23** / PR **#24** | **Skip.** Second W4 family, not dest, serve-only. **#16** unmerged. **#18** GPTQ `BLOCK_KN_SIZE` 256. **#22** two-shot `rdna_ar` (Cursor; library still **0**; zoo **512**; do not dump). **#23** PCIe P2P KV (Cursor; HIP IPC SDMA; zoo `comm.pcie` is AR, not KV disagg). **#24** RAM KV offload (George/Codex; no HIP). |
| PR **#5** / **#8** Flash-Next, extras **#11** wvSplitK, extras **#13** T44b, extras **#14**, extras **#15**, extras **#17**, extras **#19**, extras **#20**, extras **#21** | **Closed.** Dest-landed T44b / #14 / #15 / #17 / #19 / #20 / dest-integrated #21 are observe-only. Dest reverted wvSplitK (`5c4ab989`). Do not re-merge. Do not grow `gemv_f16`. |

Bodies stay in extras because every dest HIP file includes `torch/all.h`.
A dump is not a migrate. CONTRIBUTING: do not edit extras from this repo.

## Dest file → tile

Observed at `2a5e89368272`. Numbers are extras observations, not dest
Observed at `700753d9add5`. Numbers are extras observations, not dest
locks. Fill tile READMEs.

| extras path | hippihx tile | Notes |
Expand Down Expand Up @@ -108,7 +109,7 @@ extras consume.
journals FPP13 16k c=8 green via serve arenas — still extras.
2. Tile README LDS / `__launch_bounds__` / wave / DOT unit are filled.
3. hippihx ships one HIP entry extras can bind. **Started:** V1
`plan`/`run`. `run` is `NOT_READY`. extras @ `2a5e89368272` has no
`plan`/`run`. `run` is `NOT_READY`. extras @ `700753d9add5` has no
`torch.ops.hippihx.*`.
4. extras is rewired **in extras**. The extras copy is then deleted.

Expand All @@ -117,10 +118,10 @@ Until then: observe, lock numbers, keep stubs. Tracker:

## Dest extras defects

Dest tip `2a5e89368272`. Live dest bugs / rolled-back paths — hippihx
Dest tip `700753d9add5`. Live dest bugs / rolled-back paths — hippihx
must not reproduce them.

| Defect | Where | Status @ `2a5e893` | Zoo lock |
| Defect | Where | Status @ `700753d9` | Zoo lock |
|---|---|---|---|
| `llvm.amdgcn.fdot2.bf16.bf16` ISel abort | Hybrid W4 gfx10 + bf16. Dest HIP has **no** `fdot2.bf16`. | **Serve-mitigated** (`59237b3`). | Never `fdot2.bf16`. DOT + GDN HIP = **fp16 act**. |
| Scale-baked W4 zero-point | `qdq_4_rdna2.cuh` `prep_zero_scale_fp16`: `0xE400 \| zero` then `scale * (-1024 - zero)` in `half`. | **Still live** | Integer `q - zero`, then `* scale`. |
Expand Down Expand Up @@ -183,4 +184,4 @@ This review is documentation, not a kernel migrate.
| PR #3 a17t unique W4 (`d53572644`, later Simon Siebert) | **Not taken** | Closed unmerged. If dest ever locks that family, pick **their** commits |
| Explore PRs **#9/#10** sdot | **Not taken** | Not dest |
| extras dest-presence serve (**#13** T44b / **#14** / **#15** / **#17** / **#19** / **#20** / dest-integrated **#21** / mamba `741e5bc3` / closed **#12**) | observe | Python/Triton/serve + dest-landed ATen HIP. No HIP migrate. Do not pick Cursor/George/Codex/Karl/Waldecir/a17t rewrites. Unique Aron Hsiao `rdna_ar` still pickable. |
| extras not dest (**#16** closed / **#18**) | **Not taken** | GPTQ `BLOCK_KN_SIZE` 256. |
| extras not dest (**#16** closed / **#18** / **#22** / **#23** / **#24**) | **Not taken** | GPTQ `BLOCK_KN_SIZE` 256. Two-shot `rdna_ar`. PCIe P2P KV. RAM KV offload. |
2 changes: 1 addition & 1 deletion docs/CONSUME.md
Original file line number Diff line number Diff line change
Expand Up @@ -28,7 +28,7 @@ fa_fdot2.run(binding)
Register **one** `torch.ops.hippihx.<op>` that calls `hippihx_v1_run`.
No Triton→HIP double-fire. Bind keys on **arch + wave**, not GFX name
alone. extras V1 consume is the HIP fatbin. Dest extras tip
`2a5e89368272` still has no `torch.ops.hippihx` bind. FlyDSL extras
`700753d9add5` still has no `torch.ops.hippihx` bind. FlyDSL extras
consume waits on `FLYDSL_V1_CONSUME` (graph-safe JIT) — see
[`FLYDSL.md`](FLYDSL.md).

Expand Down
44 changes: 20 additions & 24 deletions docs/EXTRAS.md
Original file line number Diff line number Diff line change
Expand Up @@ -7,12 +7,13 @@ one V1 op. Review: [`BACKPORT.md`](BACKPORT.md).
## Unvalidated extras inventory

Snapshot of [`opengfx1030/vllm-rdna`](https://github.com/opengfx1030/vllm-rdna)
`rdna_extras` @ `2a5e89368272` (2026-09-24 21:52 UTC). Dest default
`rdna_extras` @ `700753d9add5` (2026-09-25 05:48 UTC). Dest default
branch is **`rdna_extras`**. `main` is upstream vLLM `c00091e02670`.
Merged extras [PR #1](https://github.com/opengfx1030/vllm-rdna/pull/1)
squash is `a4060647cfbb`. Open **#2**, **#18**. Closed **#3/#16**
unmerged; **#21** dest-integrated unmerged; **#12** superseded by
**#15**. Merged **#13**, **#14**, **#15**, **#17**, **#19**, **#20**.
squash is `a4060647cfbb`. Open **#2**, **#18**, **#22**, **#24**. Draft **#23**.
Closed **#3/#16** unmerged; **#21** dest-integrated unmerged; **#12**
superseded by **#15**. Merged **#13**, **#14**, **#15**, **#17**, **#19**,
**#20**.
**No** `torch.ops.hippihx.*`.

**Unvalidated.** Not dest. Not silicon-signed. No tok/s. hippihx still
Expand Down Expand Up @@ -50,6 +51,9 @@ Status: **extras** = live on dest tip · **Later** = side branch ·
| leapdragon push AR | `rdna_allreduce.{cu,cuh}` (merged PR **#1**) | `comm/pcie` | Dest extras, default **off**. Occupancy pin closed. **unvalidated** |
| a17t extra AWQ GEMM / GEMV | `awq_gemm_rdna2.cu` etc. (closed PR **#3**) | — | **skip** — closed unmerged, second W4 family |
| GPTQ exllama `BLOCK_KN_SIZE` 256 | `q_gemm.cu` (PR **#18**) | — | **skip** — not dest DOT W4 |
| two-shot `rdna_ar` | `rdna_allreduce.{cu,cuh}` (PR **#22**) | — | **skip** — not dest. Library `VLLM_RDNA_AR` still `"0"`. Zoo `AR_MAX_KB` **512**. Do not dump |
| PCIe P2P KV disagg | `PcieP2pConnector` (draft PR **#23**) | — | **skip** — not dest. HIP IPC SDMA. Zoo `comm.pcie` is AR, not KV. Do not dump |
| RAM KV offload | offloading connector (PR **#24**) | — | **skip** — not dest. George/Codex. Serve only. No HIP. Do not pick |
| Explore W4A8 sdot4 / W4A4 sdot8 | extras PRs **#9/#10** | — | **skip** — not dest |
| Resident W4A16 MoE skinny decode | `moe_resident_decode.cu` (dest **#17** @ `e1315629`) | `moe/routed` (watch) | Stay extras. Opt-in `VLLM_RDNA_MOE_RESIDENT*`. ATen HIP. Do not dump. Closed **#16** unmerged. **unvalidated** |

Expand Down Expand Up @@ -128,23 +132,15 @@ stay extras (no D2H under capture in the zoo).

### Not taken / leave in extras

Capture plumbing (`torch.zeros` / `new_zeros` / `zeros_like`, persist
keepalive, GDN arenas, `eager_break_during_capture`), product serve
(`qwen4_exp/**`, TunableOp, PLE offload), skinny GEMM, ConfigH, V1
FULL→PIECEWISE, vLLM custom AR, `bench_report.py`, Qwen4Exp HIP gates /
wrappers / HC compute / `_contig()` cache, seq cap 6, Flash-Next
launcher knobs, dest-reverted wvSplitK, QSA Triton bounds, recovered
extras ops, mamba spec-decode `req_idx`, dest T44b `rdna_ar` (still
opt-in; do not pick Cursor squash), dest PR **#14** (HIP MoE ignores
the JSON; do not pick), dest **#17** (resident MoE ATen HIP; do not
dump; do not pick), dest PR **#20** / dest-integrated PR **#21** (FULL+PIECEWISE
serve; zoo does not own graph mode; do not pick; do not copy tok/s),
dest PR **#19** (MTP unquantized-weight detect; MTP still not dest;
do not pick a17t/opencode), a17t PR **#3** (closed unmerged), PR **#18**
(GPTQ `BLOCK_KN_SIZE` 256),
closed **#5/#11/#13/#14/#15/#17/#19/#20/#21**, explore **#9/#10**, closed PR **#12**
(superseded by **#15**), closed PR **#16**, ROCm platform init, extras
EXL3 docker arch-guard (one `--offload-arch` per fatbin; gfx1150 /
gfx12xx not dest), Qwen4Exp MTP (not dest), amdsmi `get_device_name`
fallback, dest FA `RDNA_ATTN` pin (`48c56ef`). Produce (`-cb 3inst`,
AWQ pack) stays outside hippihx.
Stay extras: capture plumbing, product serve, skinny GEMM, ConfigH,
graph mode, vLLM custom AR, `bench_report.py`, dest T44b `rdna_ar`
(still opt-in), dest PR **#13** / PR **#14** / PR **#15** / PR **#17** /
PR **#19** / PR **#20** (observe, do not pick), dest-integrated PR **#21**,
dest FA `RDNA_ATTN` pin, resident MoE ATen HIP, Qwen4Exp MTP (not dest),
ROCm platform init, extras EXL3 docker arch-guard. Skip: a17t PR **#3**,
PR **#18** GPTQ `BLOCK_KN_SIZE` 256, PR **#22** two-shot `rdna_ar`
(Cursor; library `VLLM_RDNA_AR` still `"0"`; zoo `AR_MAX_KB` **512**;
do not dump), draft PR **#23** PCIe P2P KV (Cursor; HIP IPC SDMA; do not
dump), PR **#24** RAM KV offload (George/Codex; no HIP; do not pick),
explore **#9/#10**, closed PR **#12**, closed PR **#16**,
closed **#5/#11**. Produce stays outside hippihx.
9 changes: 7 additions & 2 deletions tests/test_no_b12x_cuda.py
Original file line number Diff line number Diff line change
Expand Up @@ -112,7 +112,7 @@ def test_readme_unvalidated_inventory() -> None:
assert "e131562" in extras
assert "48c56ef" in extras
assert "68a635" in extras
assert "2a5e893" in extras
assert "700753d9" in extras
assert "e45dd5c" in extras
assert "d1b200" in extras
assert "1046782" in extras
Expand Down Expand Up @@ -146,6 +146,7 @@ def test_readme_unvalidated_inventory() -> None:
assert "48c56ef" in bp
assert "68a635" in bp
assert "2a5e893" in bp
assert "700753d9" in bp
assert "dbb1e77" in bp
assert "e45dd5c" in bp
assert "d1b200" in bp
Expand Down Expand Up @@ -181,13 +182,17 @@ def test_readme_unvalidated_inventory() -> None:
assert "PR **#19**" in extras or "PR **#19**" in bp
assert "PR **#20**" in extras or "PR **#20**" in bp
assert "PR **#21**" in extras or "PR **#21**" in bp
assert "PR **#22**" in extras or "PR **#22**" in bp
assert "PR **#23**" in extras or "PR **#23**" in bp
assert "PR **#24**" in extras or "PR **#24**" in bp
assert "VLLM_RDNA_MOE_RESIDENT" in extras
assert "moe_resident" in extras
contrib = (ROOT / "CONTRIBUTING.md").read_text(encoding="utf-8")
assert "2a5e893" in contrib
assert "700753d9" in contrib
assert "a4060647" in contrib
# CONTRIBUTING is tip-only. Historical dest SHAs live in BACKPORT.
for token in (
"2a5e893",
"68a635",
"48c56ef",
"e131562",
Expand Down
4 changes: 2 additions & 2 deletions tiles/attention/fa_fdot2/README.md
Original file line number Diff line number Diff line change
Expand Up @@ -33,11 +33,11 @@ body yet — when you do, cherry-pick BlivionIaG, do not re-author
| `__launch_bounds__` | extras @ `a4060647`: decode D=128 `__launch_bounds__(128)`; decode D=256 + GQA-aware D=256 `__launch_bounds__(256)`; prefill `THREADS_PREFILL=128` | **not dest-locked** — occupancy pin closed |
| GQA D=256 decode | extras @ `6c5ff94`: idle waves still shuffle; skip online-softmax when `m_new == -inf` (NaN guard); zero `O_partial` on empty seq. `GQA_MAX_G=8`, `GQA_BC=32`, `GQA_DSK=256+8` | observed; not dest-locked |
| fp16 flash KV writer | extras: `reshape_and_cache_flash_rdna2`, `__launch_bounds__(128, 4)`, used for non-native KV write, fp16 only | extras HIP; do not dump ATen wrapper. Persist workspaces stay extras. |
| GQA-subgroup prefill | extras @ `ecfec4e412ad`: `fa_rdna2_prefill_paged_varlen_gqa`, `HEADS_PER_CTA=2` / `BR=8`, dual-acc fdot2. Env `VLLM_FA_RDNA2_GQA_MODE` default `subgroup`. True-GQA (`HEADS_PER_CTA=6`) dropped. Dest @ `d1b200b1`: GQA prefill O accumulator is register-resident (was smem). Occupancy pin closed. Do not dump `.cu`. Do not copy tok/s. | observed; occupancy pin closed. Do not copy tok/s. |
| GQA-subgroup prefill | extras @ `ecfec4e412ad`: `fa_rdna2_prefill_paged_varlen_gqa`, `HEADS_PER_CTA=2` / `BR=8`, dual-acc fdot2. Env `VLLM_FA_RDNA2_GQA_MODE` default `subgroup`. True-GQA (`HEADS_PER_CTA=6`) dropped. Dest @ `d1b200b1`: GQA prefill O accumulator is register-resident (was smem). Occupancy pin closed. Do not dump `.cu`. | observed; occupancy pin closed. |

Dest extras @ `48c56ef` pins `AttentionConfig.backend = RDNA_ATTN` from
the API-server `VLLM_USE_RDNA2_FA` env so workers actually select
FA-RDNA2. Stay extras. No HIP. Do not copy tok/s.
FA-RDNA2. Stay extras. No HIP.

Scratch is sized by `plan`. C consume id: `HIPPIHX_V1_OP_ATTN_FA_FDOT2`
(`include/hippihx/v1.h`). Serve wraps as one `torch.ops.hippihx.*` — no
Expand Down