From 6083dae3aa1cacc2288298a7066ec38154a5a13a Mon Sep 17 00:00:00 2001 From: Justin Card Date: Wed, 12 Aug 2026 18:55:53 -0400 Subject: [PATCH 1/3] docs(BACKEND-ROCM): spec the ROCm head_dim=128 decode arm, the ROCm half of #382 FOLLOWING_AGENTS_PROTOCOL Spec and claim only -- no source changes. The implementation follows in its own PR, against this spec. Reconciles the record before implementing, per the re-verify-before-claiming rule. Issue #382 already names this exact defect ("decode-opt attention kernel is head_dim-256 only; head_dim 128 falls to the block kernel") and PR #425 already merged the CUDA half of it. This spec is therefore the ROCm MIRROR of landed work, not new design, and it adopts that arm's flag (VT_ATTN_DECODE_D128), default (OFF) and stated reason verbatim: the arm is correctness-complete but not byte-exact, because warp-strided online softmax reduces the KV sequence in a different order than the kernel it replaces, so a greedy anchor can move at an exact bf16 tie. #488 -- the ROCm 8.1x per-call decode-attention measurement -- is recorded as the motivating observation, not the issue: it explicitly asserts no cause, and this change supplies one without closing it. Adds #488 to the roadmap issue table, where it was missing. The finding worth carrying: #382 measured this same EPL=4 arm 1.6x SLOWER on sm_110, where gfx1200 measures it 3.53x faster. Recorded, not reconciled. It is why the flip to default-ON must be argued per backend rather than once, and it is what keeps #382 open. Row: BACKEND-ROCM (the ACTIVE row whose code this edits; #382 itself is filed against the cross-backend kernel row, which check-agent-record will not accept as an active claim's row -- the two differ deliberately and the claim says so). Issue: https://github.com/mudler/vllm.cpp/issues/382 Following-Agents-Protocol: true AI-Assisted: true Assisted-by: Claude:claude-sonnet-5 [Claude Code] --- .agents/claims/CLAIM-ROCM-DECODE-ATTN-D128.md | 5 + .agents/roadmap_v1.md | 1 + .agents/specs/rocm-decode-attn-d128.md | 269 ++++++++++++++++++ 3 files changed, 275 insertions(+) create mode 100644 .agents/claims/CLAIM-ROCM-DECODE-ATTN-D128.md create mode 100644 .agents/specs/rocm-decode-attn-d128.md diff --git a/.agents/claims/CLAIM-ROCM-DECODE-ATTN-D128.md b/.agents/claims/CLAIM-ROCM-DECODE-ATTN-D128.md new file mode 100644 index 000000000..b962d0dc4 --- /dev/null +++ b/.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. | diff --git a/.agents/roadmap_v1.md b/.agents/roadmap_v1.md index e9e3a4d37..c5454207b 100644 --- a/.agents/roadmap_v1.md +++ b/.agents/roadmap_v1.md @@ -46,6 +46,7 @@ issue is not yet placed. Keyed record: update in place, never append. | [#201](https://github.com/mudler/vllm.cpp/issues/201) | `BACKEND-ROCM` | `hipblasGemmEx` overload mismatch in `rocm_matmul_hipblaslt.hip` | bug | | [#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 | +| [#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 | diff --git a/.agents/specs/rocm-decode-attn-d128.md b/.agents/specs/rocm-decode-attn-d128.md new file mode 100644 index 000000000..1ecb3f13c --- /dev/null +++ b/.agents/specs/rocm-decode-attn-d128.md @@ -0,0 +1,269 @@ +# ROCm decode paged-attention at head_dim=128 — the ROCm arm of #382 + +**Row:** `BACKEND-ROCM` (backend-matrix, `ACTIVE`) — the row this change's code +lives in. Its issue, #382, is filed against `KERNEL-ATTN-PAGED`, the +cross-backend kernel row (state `ANCHOR-BACKFILL`). The two differ deliberately: +the defect is cross-backend, this change is the ROCm arm of it, and +`check-agent-record` requires an active claim to name a `SPIKE`/`ACTIVE` row. +**Claim:** `CLAIM-ROCM-DECODE-ATTN-D128`. +**Issue:** [#382](https://github.com/mudler/vllm.cpp/issues/382) — "decode-opt +attention kernel is head_dim-256 only; head_dim 128 (Qwen3-dense, Llama, +Mistral) falls to the block kernel." This spec is the **ROCm half** of that +issue. The CUDA half already landed as +[PR #425](https://github.com/mudler/vllm.cpp/pull/425) (`66399617`); #382 stays +open because the flip to default-ON is still owed and only CUDA was covered. +**Motivating measurement:** [#488](https://github.com/mudler/vllm.cpp/issues/488) +— the ROCm per-call decode-attention gap (`PagedAttnOnline` 41.1us vs vLLM's +5.10us fallback kernel, 8.1x) on gfx1200. #488 asserts no cause; this spec +supplies one of them and does not close it. +**Base:** `origin/main` `bbc482a2`, fresh branch (not stacked — this touches +`rocm_paged_attn.hip`/its cross-device test, disjoint from #506's +`rocm_skinny_gemm.hip`/`rocm_matmul_hipblaslt.hip`). +**Board:** AMD Radeon RX 9060 XT (`gfx1200`, Navi 44, RDNA4, discrete, 32 +CUs), ROCm 7.2.3, hipClang/Clang 22.0.0. +**Reference checkout:** `${VLLM_SOURCE}` pinned at `5559679229bc` (parity pin +`555967922`). + +--- + +## 1. Reconciliation: this gap already had an issue and a landed CUDA fix + +Per the protocol's re-verify-before-claiming rule, this section records what +already existed, because the answer changed the shape of the change. + +[#382](https://github.com/mudler/vllm.cpp/issues/382) — filed by an outside +contributor against `KERNEL-ATTN-PAGED` — names this defect exactly, on the +CUDA file, and proposes the same fix this spec implements: template the +decode-opt kernel on elements-per-lane and add an `EPL=4` instantiation. +[PR #425](https://github.com/mudler/vllm.cpp/pull/425) landed that for CUDA. + +So the ROCm change is a **mirror of merged work**, not new design. It adopts +the merged arm's flag, default, and stated reason verbatim +(`cuda_paged_attn.cu`, `DecodeD128Enabled`). Where the two backends' facts +differ — and they do, sharply, in §5 — the difference is recorded rather than +averaged away. + +## 2. Why `PagedAttnOnline` was what ran + +Reading the ROCm dispatch before this change: + +```c +const bool bf16_decode_opt = + decode_opt && (d == 256 || d == 512) && query.dtype == DType::kBF16 && ... +``` + +Every fast decode kernel this file has (`PagedAttnDecodeOptBf16T`, +`PagedAttnDecodeGqaBf16`) was gated to `d == 256 || d == 512`. Qwen3-0.6B's +`head_dim` is **128** — the standard value for the great majority of dense/GQA +models this project supports (Qwen3 0.6B/1.7B/4B, Llama, Mistral). So +`bf16_decode_opt` was false for every one of them and every call fell straight +through to `PagedAttnOnline`. Not a gfx1200-specific gap, nor a WMMA-specific +one: the file had no fast path at all for the head size these models use. This +project's own Metal test (`test_metal_backend.cpp:915`, "Qwen3 geometry … +head_dim 128") independently names `d=128` as the real-model shape. + +Cross-checked against the pinned oracle (`csrc/rocm/attention.cu` @ +`555967922`): its `CALL_CUSTOM_LAUNCHER_BLK_HEAD` dispatch switches on +`head_size` with cases **64 and 128** (`attention.cu:3607-3618`), including on +RDNA4 via the `is_navi_gpu()` (`arch.find("gfx11")==0 || arch.find("gfx12")==0`) +launcher variant. Upstream's fast kernel covers `d=128`; ours didn't. + +## 3. Why this is a small change, not a new kernel + +`PagedAttnDecodeOptBf16T` and `PagedAttnDecodeGqaBf16` are +already generic over head_dim via `EPL` (elements-per-lane = `d/32`, one +wavefront's lanes splitting a `d`-wide row). Every online-softmax computation, +shared-memory sizing and warp-shuffle reduction already parametrizes on +`d`/`EPL` with no 256/512-specific assumption in the kernel body. The **only** +hardcoded constraint was in the vectorized load/store helpers: + +```c +template +__device__ inline void LoadRowEplBf16(...) { + static_assert(EPL == 8 || EPL == 16, "EPL"); // uint4-sized loads only +``` + +`EPL=4` (`d=128`) needs an 8-byte (`uint2`) load instead of a 16-byte (`uint4`) +one — the same pattern, one size down. No new algorithm, tiling, or +synchronization. This is the same observation #382 made about the CUDA file. + +## 4. What changed + +`src/vt/rocm/rocm_paged_attn.hip`: + +1. `LoadRowEplBf16<4>`/`StoreRowEplBf16<4>` — `uint2`-sized vectorized + load/store, same shape as the existing `EPL==8` branch. +2. A `VT_ATTN_DECODE_D128` gate, **default OFF**, read once beside the existing + `decode_opt`/`decode_gqa` flags — the same env var, default and rationale as + the merged CUDA arm. +3. `bf16_decode_opt` gate: `d == 256 || d == 512` → `(d == 128 && (decode_d128 + || decode_wmma)) || d == 256 || d == 512`. The `decode_wmma` disjunct is + deliberate: the rocWMMA arm (separate spec) is a second, independently + opt-in kernel for the same head size, and without it a bare + `VT_ATTN_DECODE_WMMA=1` would be a silent no-op. +4. `decode_gqa` fused-head condition extended to `qg == 2 && (d == 128 || d == + 256 || d == 512)` (Qwen3-0.6B/1.7B are `qg=2`; Qwen3-4B is `qg=4`, not fused + at any `d` today — falls to per-head `PagedAttnDecodeOptBf16T`, still off + `PagedAttnOnline`). +5. Two launch switches (`DecodeGqa`, `DecodeOptBf16T`) extended with a + `d == 128` case. + +No change to any prefill path — those stay gated to `d==256||512` and are out +of scope (§7); a `d=128` prefill call already falls through their internal +`else { goto flash_fallback; }` guards to the decode-shaped launch. + +### Why default OFF + +Verbatim from the merged CUDA arm's reason, which applies identically here: the +arm is correctness-complete but **not byte-exact** against the kernel it +replaces. Warp-strided online softmax reduces the KV sequence in a different +**order** than `PagedAttnOnline`'s per-tile loop, so a greedy anchor can move +at an exact bf16 tie. Shipping OFF keeps every existing golden byte-identical. +The flip owes the near-tie razor, a distributional gate, and regen under the +ratified-tie rule — on **both** backends, and is named as owed in §7. + +### Test coverage + +`tests/vt/test_backend_cross_device.cpp`: new case, "paged attention at Qwen3 +geometry (bf16, GQA 2, head_dim 128) matches the CPU oracle" — mirrors the +Metal "Qwen3 geometry" test's shape (`nblocks=24, bsz=16, hq=16, hkv=8, +dh=128`, 2 requests mixing a 40-token prefill and a 5-token decode-with-context, +so both the prefill fallthrough and the decode dispatch are exercised in one +call), looped over `RegisteredDevices()` so it also covers Metal/CUDA/CPU. +Checks NMSE ≤ 5e-4 against the CPU oracle and `OpProviderStats::declines == 0`. +Genuinely new coverage: the existing generic paged-attention cross-device test +at `d=8, f32` never reached any bf16 `EPL`-templated kernel, so none of them had +bf16 correctness coverage in this suite. + +`tests/CMakeLists.txt`: because the arm ships OFF **and** its flag is read into +a `static const bool` — once per process — the default registration only ever +gates the `PagedAttnOnline` fallback. Two extra ctest registrations re-run the +same binary filtered to this case with `VT_ATTN_DECODE_D128=1` and +`VT_ATTN_DECODE_WMMA=1`, so the arms that actually run the new kernels are +gated. Same shape as the existing `test_dense_gateup_fused_marlin_off_*` pair. +Verified non-vacuous (the trap `SKIP_RETURN_CODE 77` exists for, issue #463): +the filter resolves to `test cases: 1 | 1 passed`, `assertions: 6 | 6 passed`, +not zero. + +## 5. Evidence + +**The gate is live in both directions.** Same binary, no rebuild — an A/B that +was not possible before this change, because the arm was unconditional. +Qwen3-0.6B, 1024-token input, 32 output, concurrency 1, seed 0, GPU lock held: + +| `VT_ATTN_DECODE_D128` | TPOT rep1 | TPOT rep2 | +|---|---|---| +| unset (default — `PagedAttnOnline`) | 44.82 ms | 44.82 ms | +| `=1` | 12.80 ms | 12.60 ms | +| | **3.53x** | | + +**Decode throughput**, in-engine, 128in/128out, concurrency 1, isolated +same-binary A/B, back-to-back: + +| Model | Off (tok/s) | On (tok/s) | Speedup | TPOT off → on | +|---|---|---|---|---| +| Qwen3-0.6B (GQA=2, fused) | 53.84 | 76.84 | **+42.7%** | 13.80 → 8.12 ms | +| Qwen3-1.7B (GQA=2, fused) | 32.23 | 40.30 | **+25.0%** | 26.06 → 19.82 ms | +| Qwen3-4B (GQA=4, per-head only) | 20.70 | 24.38 | **+17.8%** | 42.53 → 35.44 ms | + +Single run per cell on a board that may also drive a display — indicative, not +the 2-3x-idle-reproduced standard; matches the caveat already carried by +[rocm-gfx1200-m2-correctness.md](rocm-gfx1200-m2-correctness.md) for this +board. The 1024-token table above is 2 reps and is the tighter number. GQA=2 +models (fused kernel) win more than the GQA=4 model (per-head only) — +consistent with the fused kernel halving K/V reloads per its own header +comment, though this run does not isolate that mechanism via a trace. + +**The cross-architecture reversal is the most important row in this spec.** +#382 measured the same `EPL=4` arm on sm_110 / Jetson AGX Thor as **1.6x +slower** than the block kernel it replaces (81.6 → 131.0 at c=1, and worse at +c=8). We measure **3.53x faster** on gfx1200. Both can be true: they are +different kernels (`PagedAttentionDecodeOptKernel` vs `PagedAttnDecodeGqaBf16`), +different fallbacks (CUDA's generic *block* kernel vs ROCm's `PagedAttnOnline`), +different wave widths (32 on both, but different occupancy and LDS budgets), and +different memory systems. **What this reversal forbids is a shared default.** +Any future flip to default-ON must be argued per backend with per-backend +measurement; the fact that the ROCm arm is a large win is not evidence for the +CUDA arm, and #382's sm_110 regression is not evidence against this one. + +**Correctness**, gfx1200, real hardware: `ctest -R 'rocm|cross_device'` **6/6 +pass**, including both new flag-on registrations. Full `ctest` 393 tests, +385 passed / 8 failed; all 8 reproduce identically (same tests, same root cause +`vt: no kernel for op 63 on device type 5`, an unrelated pre-existing ROCm +op-registration gap) on an isolated build of this branch **without** this +change — confirmed not caused by it. + +## 6. What this does not claim + +- **It does not close #488.** #488 reports a per-call gap against vLLM and + asserts no cause. This removes one cause. No same-tool per-call re-measure + against the oracle was run after the change (blocked on the container/glibc + issue recorded in the WMMA spec), so the residual gap is unquantified. +- **No ceiling.** The next traceable hypothesis is the `qg=4`/`qg=8` fusion gap + (§7) and, above it, the skinny-GEMM lever in #487, which #488 itself notes is + the larger share of ROCm decode time. + +## 7. Scope + +**In scope.** The `rocm_paged_attn.hip` edits in §4, the new bf16 `d=128` +cross-device test and its two flag-on ctest registrations, and this spec. + +**Out of scope, named and owed.** +- **The flip to default-ON, on both backends.** Owes the near-tie razor, a + distributional gate, and golden regen under the ratified-tie rule. Per §5 it + must be argued per backend, not once. This is what keeps #382 open. +- **rocWMMA for `d=128`.** A second, independently-flagged kernel for the same + shape — its own spec (`rocm-decode-attn-d128-wmma.md`, landing separately), + its own claim, its own issue. +- **`qg=4`/`qg=8` GQA fusion at any `d`.** `PagedAttnDecodeGqaBf16`'s fused + condition only ever covered `qg==2` and `qg==8 && d==512`; `qg=4` (Qwen3-4B) + was never fused at 256/512 either. Real, but head-dim-independent — its own + follow-on. +- **Prefill at `d=128`.** Prefill kernels stay `d==256||512`-only. +- **Explaining why GQA=2 wins more than GQA=4** via kernel tracing. Named, not + run. + +## 8. Reproduction + +```sh +nix develop .#rocm-shell --command bash -c ' + cmake -S . -B build-hip -G Ninja -DVLLM_CPP_HIP=ON \ + -DVLLM_CPP_HIP_ARCHITECTURES=gfx1200 -DROCM_PATH=$ROCM_PATH \ + -DCMAKE_BUILD_TYPE=Release + cmake --build build-hip -j"${JOBS:-8}" +' +flock "$HOME/gpu.lock" -c ' + nix develop .#rocm-shell --command ctest --test-dir build-hip \ + -R "rocm|cross_device" --output-on-failure +' +# Flag A/B, same binary: +flock "$HOME/gpu.lock" -c ' + build-hip/examples/vllm-bench --model --num-prompts 2 \ + --input-len 1024 --output-len 32 --concurrency 1 --seed 0 + VT_ATTN_DECODE_D128=1 build-hip/examples/vllm-bench --model \ + --num-prompts 2 --input-len 1024 --output-len 32 --concurrency 1 --seed 0 +' +``` + +## Outcome (2026-08-12) + +**Landed the ROCm `d=128` decode arm, default OFF, mirroring the merged CUDA +arm of the same issue.** Root cause for the ROCm decode-attention gap #488 +measured was neither architecture- nor WMMA-specific: no fast decode kernel +existed for `head_dim=128`, the size every locally-tested model uses, on any +board. #382 had already named this and PR #425 had already fixed the CUDA half; +this is the mirror, adopting that arm's flag (`VT_ATTN_DECODE_D128`), default +(OFF) and reason (reduction-order change can move a greedy anchor at a bf16 +tie) rather than inventing new ones. + +Measured on gfx1200 with the gate exercised both directions on one binary: +**3.53x** TPOT at 1024-token context, and +42.7% / +25.0% / +17.8% decode +throughput on Qwen3-0.6B / 1.7B / 4B. New bf16 `d=128` GQA correctness coverage +where none existed, plus two flag-on ctest registrations so the opt-in arms are +actually gated rather than silently skipped. + +**The finding worth carrying forward is the reversal:** #382 measured this same +arm 1.6x *slower* on sm_110, where we measure it 3.5x *faster*. That is +recorded, not reconciled, and it is the reason the default-ON flip must be +argued per backend. Rejected: landing default-ON on the strength of the ROCm +number alone. From 61fb7174f3c24cbf39246ef0cf2e63ed9b584a30 Mon Sep 17 00:00:00 2001 From: Justin Card Date: Wed, 12 Aug 2026 21:42:26 -0400 Subject: [PATCH 2/3] docs(BACKEND-ROCM): record the dtype gap -- the ROCm d128 arm is NARROWER than the CUDA arm it mirrors (#382) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit FOLLOWING_AGENTS_PROTOCOL Spec only -- no source changes. Found by reading the merged CUDA arm (66399617) hunk by hunk against this one. CUDA added TWO EPL=4 loaders, not one: LoadRowN<4, __nv_bfloat16> AND LoadRowN<4, float>. The second exists because CUDA's d == 32*4 launch branch carries no dtype gate at all, and cuda_paged_attn.cu:2796 dispatches LaunchDecode for out.dtype == kF32 -- so f32 reaches the decode-opt kernel on CUDA. ROCm's bf16_decode_opt requires query, k_cache, v_cache and out to ALL be kBF16, and LoadRowEplBf16/StoreRowEplBf16 are bf16 by construction. The fallback dispatch supports FIVE dtype combinations, so four of the five still fall to PagedAttnOnline at d=128 -- the exact fallback this arm exists to get off. f32 is live, exercised surface on ROCm, not dead code. The limitation is PRE-EXISTING, not introduced: ROCm's decode-opt has been bf16-only at every head_dim, so the d=128 arm inherits the gap rather than creating it. Recording it anyway, because §1 calls this change a mirror of merged work and on dtype coverage it is NOT one -- and because an arm that is not implemented is owed a record, not left to be discovered later. Not waived: an f32 d=128 decode is silently slow rather than refused, which is the weaker of the two failure modes. §1 gains the same qualifier so a reader of the reconciliation section is not told "mirror" without the exception. Row: BACKEND-ROCM Issue: https://github.com/mudler/vllm.cpp/issues/382 Following-Agents-Protocol: true AI-Assisted: true Assisted-by: Claude:claude-opus-5 [Claude Code] --- .agents/specs/rocm-decode-attn-d128.md | 27 ++++++++++++++++++++++++-- 1 file changed, 25 insertions(+), 2 deletions(-) diff --git a/.agents/specs/rocm-decode-attn-d128.md b/.agents/specs/rocm-decode-attn-d128.md index 1ecb3f13c..e34bf4600 100644 --- a/.agents/specs/rocm-decode-attn-d128.md +++ b/.agents/specs/rocm-decode-attn-d128.md @@ -40,8 +40,9 @@ decode-opt kernel on elements-per-lane and add an `EPL=4` instantiation. So the ROCm change is a **mirror of merged work**, not new design. It adopts the merged arm's flag, default, and stated reason verbatim (`cuda_paged_attn.cu`, `DecodeD128Enabled`). Where the two backends' facts -differ — and they do, sharply, in §5 — the difference is recorded rather than -averaged away. +differ — and they do, sharply in §5's measurement and materially in §7's dtype +coverage, where this arm is **narrower** than the CUDA one — the difference is +recorded rather than averaged away. ## 2. Why `PagedAttnOnline` was what ran @@ -215,6 +216,28 @@ cross-device test and its two flag-on ctest registrations, and this spec. - **rocWMMA for `d=128`.** A second, independently-flagged kernel for the same shape — its own spec (`rocm-decode-attn-d128-wmma.md`, landing separately), its own claim, its own issue. +- **Every dtype combination except all-`bf16` — and here the ROCm arm is + narrower than the CUDA arm it mirrors.** `bf16_decode_opt` requires `query`, + `k_cache`, `v_cache` and `out` to *all* be `kBF16`, and the vectorized helpers + are bf16 by construction (`LoadRowEplBf16`/`StoreRowEplBf16`). The fallback + dispatch supports **five** combinations — `bf16/bf16/bf16`, `f32/f32/f32`, + `bf16/bf16/f32`, `bf16/f32/bf16`, `f32/bf16/f32` — so **four of the five still + fall to `PagedAttnOnline` at `d=128`**, which is the exact fallback this arm + exists to get off. + + The CUDA arm does not have this limitation. It added `LoadRowN<4, float>` + *beside* `LoadRowN<4, __nv_bfloat16>`, and its `d == 32 * 4` launch branch + carries no dtype gate at all, so f32 reaches the decode-opt kernel there + (`cuda_paged_attn.cu:2796` dispatches `LaunchDecode` for + `out.dtype == kF32`). §1 calls this change a mirror of merged work; **on dtype + coverage it is not one**, and that is recorded here rather than left to be + discovered. + + The limitation is pre-existing, not introduced: ROCm's decode-opt has been + bf16-only at *every* head_dim, so this is a gap the `d=128` arm inherits + rather than creates. It is owed, not intended, and it is not waived — an f32 + `d=128` decode on ROCm is silently slow rather than refused, which is the + weaker of the two failure modes AGENTS.md allows. - **`qg=4`/`qg=8` GQA fusion at any `d`.** `PagedAttnDecodeGqaBf16`'s fused condition only ever covered `qg==2` and `qg==8 && d==512`; `qg=4` (Qwen3-4B) was never fused at 256/512 either. Real, but head-dim-independent — its own From de4904125fe6deceba41686f632cd3ba624eda0e Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Thu, 13 Aug 2026 20:38:08 +0000 Subject: [PATCH 3/3] docs(BACKEND-ROCM): the spec claimed a landed arm that is not on main, and lacked Risks and Stop conditions (#382) Review repairs on top of joral's spec, pushed to the PR branch so the authorship of the work itself stays with them. Spec-only; no source file is touched. The blocking one: `## Outcome (2026-08-12)` opened with "Landed the ROCm `d=128` decode arm". Nothing landed. On main, `git log -S'VT_ATTN_DECODE_D128' -- src/vt/rocm/` is EMPTY and `rocm_paged_attn.hip:1684` still reads `d == 256 || d == 512`. AGENTS.md scopes `## Outcome` to a row reaching DONE, and `BACKEND-ROCM` is ACTIVE with the code unmerged. Merging that text would put a false "landed" on main, where the next agent greps for the flag, finds nothing, and cannot tell "never merged" from "reverted" from "renamed" -- the exact ambiguity this project has already lost time to. Renamed to `## Result on the implementation branch`, with a banner stating outright that it is not landed, quoting the two commands that show it, and naming what has to be true before it becomes `## Outcome`. Section 4 moves from "What changed" to "What will change" for the same reason: it enumerates five edits to `rocm_paged_attn.hip` that are not in the tree. Added the two sections AGENTS.md requires and this spec did not have: - `## 8. Risks and decisions` -- the bf16-tie reduction-order risk that is the reason the arm ships default OFF, the single-board provenance of the 3.53x, the sm_110 1.6x-SLOWER reversal recorded but deliberately unreconciled, the spec-before-code consequence, the dangling WMMA flag, and the unquantified residual #488 gap. - `## 9. Stop conditions` -- stop if the gate moves with the flag OFF, stop if the flag-ON arm cannot be shown to REACH the new kernel (a green gate that never entered the kernel proves nothing), stop before flipping the default, stop if a second board contradicts the first, and do not extend this row to the rocWMMA arm. `## 6. What this does not claim` was left where it is; it is a non-claims list, not a stop condition, and it reads correctly beside the new sections. Three smaller corrections, each verified rather than assumed: - `**Base:**` said `bbc482a2`; the branch was rebased and its own claim file records `fafa16f0`. The spec and the claim now agree. - Section 4 specifies a gate referencing `VT_ATTN_DECODE_WMMA`. That flag exists nowhere in `src/` or `tests/`, and `rocm-decode-attn-d128-wmma.md` is neither in this PR nor on main, so an implementer following section 4 verbatim writes a reference to an undefined symbol. Marked as a forward reference. - The upstream anchor `attention.cu:3607-3618` is `3609-3620` at the pin. The content it names is exact; the line numbers had drifted by two. Verified on this tree: check-doc-checkpoint OK, check-agent-record OK (ENGINE=152 MODEL=362 QUANT=82 KERNEL=51 BACKEND=81). FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: Claude:claude-opus-5 [Claude Code] --- .agents/specs/rocm-decode-attn-d128.md | 52 ++++++++++++++++++++++---- 1 file changed, 45 insertions(+), 7 deletions(-) diff --git a/.agents/specs/rocm-decode-attn-d128.md b/.agents/specs/rocm-decode-attn-d128.md index e34bf4600..8b964aa57 100644 --- a/.agents/specs/rocm-decode-attn-d128.md +++ b/.agents/specs/rocm-decode-attn-d128.md @@ -16,7 +16,8 @@ open because the flip to default-ON is still owed and only CUDA was covered. — the ROCm per-call decode-attention gap (`PagedAttnOnline` 41.1us vs vLLM's 5.10us fallback kernel, 8.1x) on gfx1200. #488 asserts no cause; this spec supplies one of them and does not close it. -**Base:** `origin/main` `bbc482a2`, fresh branch (not stacked — this touches +**Base:** `origin/main` `fafa16f0` after rebase (the claim file records the +same), fresh branch (not stacked — this touches `rocm_paged_attn.hip`/its cross-device test, disjoint from #506's `rocm_skinny_gemm.hip`/`rocm_matmul_hipblaslt.hip`). **Board:** AMD Radeon RX 9060 XT (`gfx1200`, Navi 44, RDNA4, discrete, 32 @@ -65,7 +66,7 @@ head_dim 128") independently names `d=128` as the real-model shape. Cross-checked against the pinned oracle (`csrc/rocm/attention.cu` @ `555967922`): its `CALL_CUSTOM_LAUNCHER_BLK_HEAD` dispatch switches on -`head_size` with cases **64 and 128** (`attention.cu:3607-3618`), including on +`head_size` with cases **64 and 128** (`attention.cu:3609-3620`), including on RDNA4 via the `is_navi_gpu()` (`arch.find("gfx11")==0 || arch.find("gfx12")==0`) launcher variant. Upstream's fast kernel covers `d=128`; ours didn't. @@ -88,7 +89,7 @@ __device__ inline void LoadRowEplBf16(...) { one — the same pattern, one size down. No new algorithm, tiling, or synchronization. This is the same observation #382 made about the CUDA file. -## 4. What changed +## 4. What will change `src/vt/rocm/rocm_paged_attn.hip`: @@ -101,7 +102,10 @@ synchronization. This is the same observation #382 made about the CUDA file. || decode_wmma)) || d == 256 || d == 512`. The `decode_wmma` disjunct is deliberate: the rocWMMA arm (separate spec) is a second, independently opt-in kernel for the same head size, and without it a bare - `VT_ATTN_DECODE_WMMA=1` would be a silent no-op. + `VT_ATTN_DECODE_WMMA=1` would be a silent no-op. **Forward reference:** + neither `VT_ATTN_DECODE_WMMA` nor `rocm-decode-attn-d128-wmma.md` exists yet + — in the tree or in this PR — so an implementer working from this section + must land the flag with the rocWMMA arm, not cite it from here. 4. `decode_gqa` fused-head condition extended to `qg == 2 && (d == 128 || d == 256 || d == 512)` (Qwen3-0.6B/1.7B are `qg=2`; Qwen3-4B is `qg=4`, not fused at any `d` today — falls to per-head `PagedAttnDecodeOptBf16T`, still off @@ -246,7 +250,34 @@ cross-device test and its two flag-on ctest registrations, and this spec. - **Explaining why GQA=2 wins more than GQA=4** via kernel tracing. Named, not run. -## 8. Reproduction +## 8. Risks and decisions + +| Risk | Assessment | +|---|---| +| The reduction-order change moves a greedy anchor at a bf16 tie | This is why the arm ships **default OFF**, adopting the merged CUDA arm's flag, default and stated reason verbatim rather than inventing new ones. A default-ON flip is a separate, per-backend argument and is explicitly out of scope here. | +| The 3.53x is a single-board, single-run figure | Measured on one gfx1200 that may also drive a display. It is indicative, not the idle-box reproduced standard AGENTS.md requires for a binding number, and §5 says so. It justifies building the arm; it does not license a BENCHMARKS entry or a default flip. | +| The same arm measured 1.6x SLOWER on sm_110 (#382) | Recorded, deliberately not reconciled. It is the reason the default stays OFF and the reason the flip must be argued per backend rather than once. Treating the ROCm number as settling the question for all boards is the error this row is guarding against. | +| The spec lands before its code | Intended, and required — AGENTS.md puts the spec before implementation. The consequence is that §4 and the result section describe an unmerged branch, which the banner above the result section states outright so no reader mistakes it for landed work. | +| `VT_ATTN_DECODE_WMMA` is cited but does not exist | A forward reference to a sibling arm whose spec and issue are not yet filed (§4). An implementer must land the flag alongside the rocWMMA arm; taking §4 literally today produces a reference to an undefined symbol. | +| The residual #488 gap is unquantified | No post-change per-call oracle re-measure was run, so how much of #488 this closes is unknown. Named here rather than left implicit; it is owed before #488 can be judged. | + +## 9. Stop conditions + +- **Stop if the token-exact gate moves with the flag OFF.** The arm is opt-in; + flag-off behaviour must be byte-identical to today. Any drift means the gate + predicate is wrong, not that the golden needs refreshing. +- **Stop if the flag-ON arm cannot be shown to reach the new kernel.** A green + gate that never entered `PagedAttnDecodeOptBf16T` proves nothing — the ctest + registrations in §4 exist precisely so the opt-in arms are gated rather than + silently skipped. Confirm selection counts, not just tokens. +- **Stop before flipping the default.** That needs its own row, and needs the + sm_110 reversal in §6 reconciled rather than out-voted by one board. +- **Stop if the measured speedup on a second board contradicts the first.** + Two boards disagreeing is a finding about the kernel, not noise to average. +- **Do not extend this row to the rocWMMA arm.** It is a separate, independently + flagged kernel with its own spec and its own issue, both still unfiled. + +## 10. Reproduction ```sh nix develop .#rocm-shell --command bash -c ' @@ -268,9 +299,16 @@ flock "$HOME/gpu.lock" -c ' ' ``` -## Outcome (2026-08-12) +## Result on the implementation branch (2026-08-12) -**Landed the ROCm `d=128` decode arm, default OFF, mirroring the merged CUDA +> **Not landed.** This section records what was built and measured on the +> unmerged implementation branch. No `VT_ATTN_DECODE_D128` exists in +> `src/vt/rocm/` on `main` — `git log -S'VT_ATTN_DECODE_D128' -- src/vt/rocm/` +> is empty, and `rocm_paged_attn.hip` still gates on `d == 256 || d == 512`. +> This spec is committed BEFORE its implementation, per AGENTS.md; the section +> becomes `## Outcome` when the code merges and the row reaches `DONE`. + +**Built the ROCm `d=128` decode arm, default OFF, mirroring the merged CUDA arm of the same issue.** Root cause for the ROCm decode-attention gap #488 measured was neither architecture- nor WMMA-specific: no fast decode kernel existed for `head_dim=128`, the size every locally-tested model uses, on any