spec(BACKEND-ROCM): the ROCm head_dim=128 decode arm, the ROCm half of #382

.agents/claims/CLAIM-ROCM-DECODE-ATTN-D128.md

@@ -0,0 +1,5 @@
+# CLAIM-ROCM-DECODE-ATTN-D128
+
+| Claim | Row IDs | Agent | Worktree / remote dir | Branch | Owned scope | State | Last update |
+|---|---|---|---|---|---|---|---|
+| `CLAIM-ROCM-DECODE-ATTN-D128` | `BACKEND-ROCM` (`ACTIVE`) | Claude Code (sonnet-5), helper role | worktree `rdna3-kernel-porting-b9ec47`, real gfx1200 hardware (AMD Radeon RX 9060 XT, RDNA4, 32 CU), `$GPU_LOCK` respected | `row/ROCM-DECODE-ATTN-D128-SPEC` (this spec; the implementation follows on `row/ROCM-DECODE-ATTN-D128-IMPL`, stacked), base `main` `fafa16f0`; issue [#382](https://github.com/mudler/vllm.cpp/issues/382) (the ROCm half; the CUDA half landed as [PR #425](https://github.com/mudler/vllm.cpp/pull/425), `66399617`), motivated by [#488](https://github.com/mudler/vllm.cpp/issues/488). NOTE: #382 is filed against the cross-backend kernel row (state `ANCHOR-BACKFILL`), while this claim's Row ID is the `ACTIVE` backend row whose code it edits — `check-agent-record` requires an active claim to name a `SPIKE`/`ACTIVE` row, so the two deliberately differ | Owns ONLY: the `LoadRowEplBf16`/`StoreRowEplBf16` `EPL=4` case, the `VT_ATTN_DECODE_D128` gate (default OFF, same flag/default/reason as the merged CUDA arm), the `bf16_decode_opt`/`decode_gqa` gate extensions and the two `d==128` launch-dispatch branches in `src/vt/rocm/rocm_paged_attn.hip`; the new "Qwen3 geometry (bf16, GQA 2, head_dim 128)" case in `tests/vt/test_backend_cross_device.cpp` and its two flag-on ctest registrations in `tests/CMakeLists.txt`; `.agents/specs/rocm-decode-attn-d128.md` and this claim file. **NON-COLLISION:** disjoint from `CLAIM-ROCM-SKINNY-GEMM-GFX1200` (different files: `rocm_skinny_gemm.hip`/`rocm_matmul_hipblaslt.hip` vs `rocm_paged_attn.hip`), not stacked on any other branch. EXCLUDED: **the flip to default-ON on either backend** (owes the near-tie razor + distributional gate + golden regen, and per the spec §5 cross-arch reversal must be argued per backend — this is what keeps #382 open), rocWMMA for `d=128` (separate claim, separate spec, separate issue), `qg=4`/`qg=8` GQA fusion at any `d` (pre-existing, board-independent gap), any `d=128` prefill path, and the 8 pre-existing unrelated `ctest` failures (`vt: no kernel for op 63 on device type 5`) | `ACTIVE` | 2026-08-12 — **reconciled against the existing record before landing**, per the re-verify-before-claiming rule: #382 already named this exact defect and PR #425 had already merged the CUDA half, so this became a mirror of merged work rather than new design, and was re-gated from default-ON to **default OFF behind `VT_ATTN_DECODE_D128`** — the merged arm's own flag, default and stated reason (warp-strided online softmax reduces the KV sequence in a different order, so a greedy anchor can move at an exact bf16 tie; OFF keeps every golden byte-identical). gfx1200-verified: `ctest -R 'rocm\|cross_device'` **6/6** including two new flag-on registrations (verified non-vacuous: 1 case, 6 assertions, not zero); full `ctest` 385/393 with the 8 failures independently confirmed pre-existing. Gate exercised **both directions on one binary** — Qwen3-0.6B @1024 ctx TPOT 44.82/44.82 ms OFF vs 12.80/12.60 ms ON = **3.53x**; decode throughput +42.7% / +25.0% / +17.8% on 0.6B / 1.7B / 4B. **Carried finding:** #382 measured this same `EPL=4` arm **1.6x slower** on sm_110 where we measure it 3.5x faster — recorded, not reconciled; it is why the default-ON flip must be argued per backend. Rebased from `bbc482a2` onto `main` `fafa16f0` (167 commits), which required reformatting `Assisted-by` for the `check-commit-trailers` gate that landed in between, and de-linking §7's forward reference to the rocWMMA spec — that spec now lands on its own branch, so a markdown link to it fails `check-agent-record` as a dangling link. Spec content otherwise byte-identical. Re-gated on the new base, gfx1200: build 783/783, `ctest -R 'rocm\|cross_device'` 6/6, the new case non-vacuous under both flags (1 case, 6 assertions), full `ctest` with 8 pre-existing `kSharedExpertGate` (`OpId(63)`) failures owed to unmerged PR #509. `agent-preflight` fails 11, set-identical to a clean `fafa16f0` baseline. Spec PR open; implementation PR follows. |

.agents/roadmap_v1.md

@@ -48,6 +48,7 @@ issue is not yet placed. Keyed record: update in place, never append.
 | [#269](https://github.com/mudler/vllm.cpp/issues/269) | `BACKEND-ROCM` | ROCm gfx1200: Gemma-3 is 48/48 exact vs two vLLM-ROCm oracles; Qwen3-0.6B exposes a deterministic cross-version near-tie, not a backend defect | verification |
 | [#332](https://github.com/mudler/vllm.cpp/issues/332) | `BACKEND-ROCM` | ROCm: no decode-graph capture — the hipGraph seam is unimplemented, costing ~3x decode throughput vs vLLM on gfx1200 | perf |
 | [#444](https://github.com/mudler/vllm.cpp/issues/444) | `BACKEND-ROCM` | ROCm: main does not build for gfx1200/gfx1201 — `rocm_paged_attn.hip` includes rocwmma unconditionally on arch, not availability | bug |
+| [#488](https://github.com/mudler/vllm.cpp/issues/488) | `BACKEND-ROCM` | ROCm gfx1200: PagedAttnOnline is 8.1x slower per call than vLLM's paged-attention kernel (41.1us vs 5.10us) | perf |
 | [#125](https://github.com/mudler/vllm.cpp/issues/125) | `BACKEND-VULKAN` | Vulkan on AMD Strix Halo (gfx1151) does not load | bug |
 | [#203](https://github.com/mudler/vllm.cpp/issues/203) | `BACKEND-VULKAN` | Vulkan on unified memory holds TWO copies of the weights: 27B peaks at 100.8 GiB RSS and OOM-reboots a Spark | bug |
 | [#310](https://github.com/mudler/vllm.cpp/issues/310) | `BACKEND-VULKAN` | docs/FEATURES.md understates Vulkan: says decode 4.24 vs 4.35 where the binding figure is 4.36 vs 4.35 | bug |