From d6f98865427189bdec5769d1ec5e4be91858b07d Mon Sep 17 00:00:00 2001 From: Justin Card Date: Wed, 12 Aug 2026 21:58:07 -0400 Subject: [PATCH 1/3] feat(rocm): head_dim=128 decode arm -- the ROCm half of #382, default OFF FOLLOWING_AGENTS_PROTOCOL Implements .agents/specs/rocm-decode-attn-d128.md, committed ahead of this change. The ROCm mirror of PR #425's CUDA arm for issue #382. Every fast decode kernel in rocm_paged_attn.hip was gated to d == 256 || d == 512, so head_dim 128 -- Qwen3-dense, Llama and Mistral -- fell straight through to PagedAttnOnline. LoadRowEplBf16/StoreRowEplBf16 gain an EPL=4 (uint2) case, and both decode dispatch switches gain a d == 128 branch, so that geometry reaches PagedAttnDecodeGqaBf16 (qg=2, fused) or PagedAttnDecodeOptBf16T (per-head) instead. Gated VT_ATTN_DECODE_D128, DEFAULT OFF -- the same env var, default and reason as the merged CUDA arm, because the warp-strided online softmax reduces the KV sequence in a different ORDER than PagedAttnOnline, so a greedy anchor can move at an exact bf16 tie and OFF keeps every golden byte-identical. No allowlist entry is owed: the CUDA arm already added VT_ATTN_DECODE_D128. Testing: a new bf16 Qwen3-geometry (GQA 2, head_dim 128) cross-device case, where the suite previously had no bf16 coverage of any EPL-templated kernel. Because the arm ships OFF and its flag is read into a static const bool -- once per process -- the default registration can only ever gate the fallback, so tests/CMakeLists.txt registers a second invocation of the same binary with the flag set. Verified non-vacuous against the issue #463 trap: 1 case, 6 assertions, not zero. Scope note: this commit is the SCALAR arm only. The rocWMMA d=128 decode spike that was originally developed alongside it is split into its own spec and its own PR -- it is a spike with its own (unfiled) issue, and bundling the two put 474 lines of two unrelated kernels in one review. The `|| decode_wmma` disjunct in the bf16_decode_opt gate belongs to that change, not this one: it only means anything once a second opt-in kernel for this head size exists. Row: BACKEND-ROCM Issue: https://github.com/mudler/vllm.cpp/issues/382 Following-Agents-Protocol: true AI-Assisted: true Assisted-by: Claude:claude-sonnet-5 [Claude Code] Assisted-by: Claude:claude-opus-5 [Claude Code] --- .agents/claims/CLAIM-ROCM-DECODE-ATTN-D128.md | 2 +- src/vt/rocm/rocm_paged_attn.hip | 86 +++++++++++++---- tests/CMakeLists.txt | 10 ++ tests/vt/test_backend_cross_device.cpp | 93 +++++++++++++++++++ 4 files changed, 171 insertions(+), 20 deletions(-) diff --git a/.agents/claims/CLAIM-ROCM-DECODE-ATTN-D128.md b/.agents/claims/CLAIM-ROCM-DECODE-ATTN-D128.md index b962d0dc4..554d7b7bf 100644 --- a/.agents/claims/CLAIM-ROCM-DECODE-ATTN-D128.md +++ b/.agents/claims/CLAIM-ROCM-DECODE-ATTN-D128.md @@ -2,4 +2,4 @@ | 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. | +| `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-IMPL` (the implementation; its spec landed from `row/ROCM-DECODE-ATTN-D128-SPEC` as [PR #564](https://github.com/mudler/vllm.cpp/pull/564), squashed to `373aa125`), rebased off the now-merged spec commits onto `main` `2784dd7b`; 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 `VT_ATTN_DECODE_D128` flag-on ctest registration in `tests/CMakeLists.txt` (the second, `VT_ATTN_DECODE_WMMA`, moved to the rocWMMA branch with the arm it gates); `.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 pre-existing unrelated `ctest` failures (`vt: no kernel for op SharedExpertGate` on ROCm, plus a missing `shellcheck`, an mmap-RSS assertion and a JSON type error) | `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. **2026-08-14 — spec LANDED as PR #564 (`373aa125`); this claim now tracks the implementation.** Rebased off the two now-squashed spec commits onto `main` `2784dd7b`; the commit is source-only (3 files) and carries no forward reference to the rocWMMA flag, so it stands alone. Re-gated on that base, gfx1200, `$GPU_LOCK` held: build 1220/1220; `ctest -R 'rocm\|cross_device'` **5/5** (5 not 6 — the `VT_ATTN_DECODE_WMMA` registration left with its arm); flag A/B on ONE binary re-measured **3.47x** (Qwen3-0.6B @1024 ctx, 45.15/45.13 ms OFF vs 12.97/13.04 ms ON), holding the 3.53x from the old base across 76 commits of drift. Full `ctest` 448/455 with **7** failures, and those 7 are now PROVEN pre-existing rather than argued: a clean `main` `2784dd7b` worktree, built from source with none of this code, fails the identical set (only `test_op_parity`'s index shifts 403→404, from the added registration). `agent-preflight` fails 9, a strict SUBSET of that same baseline's 10 (differing only by `role-undeclared`). Note `origin/main` (the `joral` fork) is 75 commits behind `upstream/main`, so preflight's range gates grade 76 commits of which 75 are other people's — `check-commit-trailers` and `check-doc-checkpoint` both pass against `upstream/main`, the base the spec actually merged to. **Fresh evidence 2026-08-14, SUPERSEDING the `+42.7% / +25.0% / +17.8%` figures above** — those came from a 128-token-context stash-based A/B; every number here is 1024-token synthetic prompt, 128 generated, greedy, seed 0, one binary, `$GPU_LOCK` held, 2 reps per cell agreeing within ~1%. Four-model TPOT OFF→ON: Qwen3-0.6B 42.53→11.78 ms (**3.61x**, `qg=2` fused), Qwen3-1.7B 52.85→21.93 ms (**2.41x**, `qg=2` fused), Qwen3-4B 81.89→39.22 ms (**2.09x**, `qg=4` per-head — no GQA fusion at any `d`, so this isolates the `EPL` widening from the fusion), Qwen3.5-0.8B 23.76→23.55 ms (**1.01x**). The last is the **NEGATIVE CONTROL** and it earned its keep: its `head_dim` is 256, so the `d == 128` gate provably cannot reach it, yet its first OFF rep landed a 33% outlier at 31.14 ms — a blind 2-rep average would have reported a bogus ~1.2x "win" for a model the flag cannot affect. Re-run 3x it gives 23.86/23.75/23.68 against ON's 23.52/23.57. End-to-end output throughput rises less than TPOT on the same runs (0.6B 2.48x, 1.7B 2.05x, 4B 2.02x) because these carry a 1024-token prefill the flag does not touch; TPOT isolates decode, throughput dilutes it. **Qwen3-1.7B concurrency sweep** (`--num-prompts` = 2x concurrency), throughput tok/s OFF→ON (ratio): c1 12.89→24.66 (1.91x), c2 23.27→47.45 (2.04x), c4 39.10→86.86 (2.22x), c8 58.97→147.35 (**2.50x**), c16 78.43→227.08 (**2.90x**); TPOT ratio over the same points 2.40x→3.18x. **The advantage GROWS with concurrency rather than compressing** — the opposite of the prediction made before the run, which reasoned that a tiny c1 grid flatters the fast kernel. The dominant effect is the reverse: from c8 to c16 the fallback scales only 1.33x against the arm's 1.54x, and scaling efficiency at c16 relative to perfect-linear-from-c1 is **38% OFF against 58% ON**. `PagedAttnOnline` is therefore the batch-scaling bottleneck, not merely slow per call, and the win is largest in the regime a server actually runs in. The c1 row reproduces the independent four-model sweep to within ~1% (52.85/21.93 vs 53.40/22.26), a cross-check on run-to-run stability. **Caveats:** single board; `--input-len` builds synthetic tokens, so all of the above is a decode-path A/B and not a serving benchmark. Implementation PR not yet opened. | diff --git a/src/vt/rocm/rocm_paged_attn.hip b/src/vt/rocm/rocm_paged_attn.hip index f73993197..e0e658bfe 100644 --- a/src/vt/rocm/rocm_paged_attn.hip +++ b/src/vt/rocm/rocm_paged_attn.hip @@ -262,32 +262,48 @@ __global__ void PagedAttnOnline(TO* out, const TQ* query, const TKV* k_cache, constexpr int kDecWarps = 8; constexpr int kDecEpl = 8; // 32 lanes * 8 = 256 head_dim -// EPL elems/lane: 8 → d=256, 16 → d=512 (32 lanes). +// EPL elems/lane: 4 → d=128 (Qwen3/Llama-class GQA), 8 → d=256, 16 → d=512 +// (32 lanes). d=128 (#488, ROCM-DECODE-ATTN-D128): the kernel body below is +// already generic in EPL; only the 4/8/16-element vectorized load/store and +// the host dispatch needed a new case, not a new kernel. template __device__ inline void LoadRowEplBf16(const __hip_bfloat16* p, int64_t base, int lane, float r[EPL]) { - static_assert(EPL == 8 || EPL == 16, "EPL"); - const uint4* u = reinterpret_cast(p + base); - if constexpr (EPL == 8) { - const uint4 w = u[lane]; + static_assert(EPL == 4 || EPL == 8 || EPL == 16, "EPL"); + if constexpr (EPL == 4) { + const uint2* u = reinterpret_cast(p + base); + const uint2 w = u[lane]; const __hip_bfloat16* h = reinterpret_cast(&w); #pragma unroll - for (int i = 0; i < 8; ++i) r[i] = __bfloat162float(h[i]); + for (int i = 0; i < 4; ++i) r[i] = __bfloat162float(h[i]); } else { - const uint4 w0 = u[lane]; - const uint4 w1 = u[lane + 32]; - const __hip_bfloat16* h0 = reinterpret_cast(&w0); - const __hip_bfloat16* h1 = reinterpret_cast(&w1); + const uint4* u = reinterpret_cast(p + base); + if constexpr (EPL == 8) { + const uint4 w = u[lane]; + const __hip_bfloat16* h = reinterpret_cast(&w); #pragma unroll - for (int i = 0; i < 8; ++i) r[i] = __bfloat162float(h0[i]); + for (int i = 0; i < 8; ++i) r[i] = __bfloat162float(h[i]); + } else { + const uint4 w0 = u[lane]; + const uint4 w1 = u[lane + 32]; + const __hip_bfloat16* h0 = reinterpret_cast(&w0); + const __hip_bfloat16* h1 = reinterpret_cast(&w1); +#pragma unroll + for (int i = 0; i < 8; ++i) r[i] = __bfloat162float(h0[i]); #pragma unroll - for (int i = 0; i < 8; ++i) r[i + 8] = __bfloat162float(h1[i]); + for (int i = 0; i < 8; ++i) r[i + 8] = __bfloat162float(h1[i]); + } } } template __device__ inline void StoreRowEplBf16(__hip_bfloat16* p, int64_t base, int lane, const float r[EPL]) { - static_assert(EPL == 8 || EPL == 16, "EPL"); - if constexpr (EPL == 8) { + static_assert(EPL == 4 || EPL == 8 || EPL == 16, "EPL"); + if constexpr (EPL == 4) { + __hip_bfloat16 tmp[4]; +#pragma unroll + for (int i = 0; i < 4; ++i) tmp[i] = __float2bfloat16(r[i]); + reinterpret_cast(p + base)[lane] = *reinterpret_cast(tmp); + } else if constexpr (EPL == 8) { __hip_bfloat16 tmp[8]; #pragma unroll for (int i = 0; i < 8; ++i) tmp[i] = __float2bfloat16(r[i]); @@ -1638,7 +1654,6 @@ __global__ void PagedAttnPrefillSharedKWmma(__hip_bfloat16* out, const __hip_bfl #endif // VT_ROCWMMA_OK } - } // namespace void PagedAttentionKernelRocm(Queue& q, Tensor& out, const Tensor& query, const Tensor& k_cache, @@ -1681,8 +1696,26 @@ void PagedAttentionKernelRocm(Queue& q, Tensor& out, const Tensor& query, const return true; // default ON when decode_opt }(); + // head_dim-128 arm (#382) -- the Qwen3-dense / Llama / Mistral head size. + // Same kernels as d==256/512, EPL=4 instead of 8/16; before this, bf16 decode + // at d==128 fell all the way to the generic PagedAttnOnline (#488 measured + // that fallback at 41.1us/call against vLLM's 5.10us on gfx1200). + // + // DEFAULT OFF, opt in with VT_ATTN_DECODE_D128=1 -- the SAME flag, default and + // reason as the merged CUDA arm (cuda_paged_attn.cu, DecodeD128Enabled, #382 / + // PR #425). It 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 + distributional gate + regen under the ratified- + // tie rule, on both backends, and is a separate change. + static const bool decode_d128 = [] { + const char* e = std::getenv("VT_ATTN_DECODE_D128"); + return e != nullptr && e[0] == '1'; + }(); const bool bf16_decode_opt = - decode_opt && (d == 256 || d == 512) && query.dtype == DType::kBF16 && + decode_opt && ((d == 128 && decode_d128) || d == 256 || d == 512) && + query.dtype == DType::kBF16 && k_cache.dtype == DType::kBF16 && v_cache.dtype == DType::kBF16 && out.dtype == DType::kBF16 && kc_hd % 8 == 0; @@ -1690,7 +1723,8 @@ void PagedAttentionKernelRocm(Queue& q, Tensor& out, const Tensor& query, const const int64_t qg = hq / num_kv_heads; // GQA fuse: sliding QG=2. Global QG=8: tile as DecodeGqa QG=2 × z=4 // (proven kernel; halves K/V reloads vs DecodeOpt's 8×). QG=4/8 CTAs lost A/B. - if (decode_gqa && ((qg == 2 && (d == 256 || d == 512)) || (qg == 8 && d == 512))) { + if (decode_gqa && + ((qg == 2 && (d == 128 || d == 256 || d == 512)) || (qg == 8 && d == 512))) { // Prefill SharedK (scoreless). Default ON for long single-req prefill (lab win // ~1.09× @11k vs DecodeGqa-z4). Set VT_ATTN_PREFILL_FLASH_SHAREDK=0 to disable. // SHAREDK_WMMA=1: rocWMMA QK + online V (lab; keep only if beats scalar SharedK). @@ -1871,7 +1905,14 @@ void PagedAttentionKernelRocm(Queue& q, Tensor& out, const Tensor& query, const const unsigned z_groups = (qg == 8 && d == 512) ? 4u : 1u; dim3 grid(static_cast(total_q), static_cast(num_kv_heads), z_groups); dim3 block(static_cast(nwarps * 32)); - if (d == 256) { + if (d == 128) { + PagedAttnDecodeGqaBf16<2, 4, 8><<>>( + out.Ptr<__hip_bfloat16>(), query.Ptr<__hip_bfloat16>(), k_cache.Ptr<__hip_bfloat16>(), + v_cache.Ptr<__hip_bfloat16>(), block_table.Ptr(), seq_lens.Ptr(), + query_start_loc.Ptr(), num_reqs, hq, num_kv_heads, d, block_size, bt_row, + bt_col, kc_blk, kc_pg, kc_hd, vc_blk, vc_pg, vc_hd, args.scale, args.logits_soft_cap, + args.causal, window_left, window_right); + } else if (d == 256) { PagedAttnDecodeGqaBf16<2, 8, 8><<>>( out.Ptr<__hip_bfloat16>(), query.Ptr<__hip_bfloat16>(), k_cache.Ptr<__hip_bfloat16>(), v_cache.Ptr<__hip_bfloat16>(), block_table.Ptr(), seq_lens.Ptr(), @@ -1893,7 +1934,14 @@ void PagedAttentionKernelRocm(Queue& q, Tensor& out, const Tensor& query, const sizeof(float) * (static_cast(kDecWarps) * static_cast(d) + 2 * static_cast(kDecWarps)); dim3 grid(static_cast(total_q), static_cast(hq)); - if (d == 256) { + if (d == 128) { + PagedAttnDecodeOptBf16T<4><<>>( + out.Ptr<__hip_bfloat16>(), query.Ptr<__hip_bfloat16>(), k_cache.Ptr<__hip_bfloat16>(), + v_cache.Ptr<__hip_bfloat16>(), block_table.Ptr(), seq_lens.Ptr(), + query_start_loc.Ptr(), num_reqs, hq, num_kv_heads, d, block_size, bt_row, bt_col, + kc_blk, kc_pg, kc_hd, vc_blk, vc_pg, vc_hd, args.scale, args.logits_soft_cap, args.causal, + window_left, window_right); + } else if (d == 256) { PagedAttnDecodeOptBf16T<8><<>>( out.Ptr<__hip_bfloat16>(), query.Ptr<__hip_bfloat16>(), k_cache.Ptr<__hip_bfloat16>(), v_cache.Ptr<__hip_bfloat16>(), block_table.Ptr(), seq_lens.Ptr(), diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt index 82c035f1f..b85c0509b 100644 --- a/tests/CMakeLists.txt +++ b/tests/CMakeLists.txt @@ -1437,6 +1437,16 @@ add_test(NAME test_vulkan_loader COMMAND test_vulkan_loader) # so it is meaningful on a CUDA box, on the Metal M4, and inert (CPU-only bit- # exactness checks) on a plain CPU build. vllm_cpp_add_test(test_backend_cross_device vt/test_backend_cross_device.cpp) +# The head_dim-128 decode arm (#382) ships default OFF, and its flag is read into +# a `static const bool` -- once per process. The default registration above +# therefore only ever gates the PagedAttnOnline fallback; the arm that runs the new +# kernel needs its own INVOCATION with the flag set. Same shape as the +# test_dense_gateup_fused_marlin_off_* lever pair above. The filter avoids commas: +# doctest reads them as filter separators and the case name contains three. +add_test(NAME test_backend_cross_device_vt_attn_decode_d128 + COMMAND test_backend_cross_device "--test-case=*head_dim 128*") +set_tests_properties(test_backend_cross_device_vt_attn_decode_d128 PROPERTIES + ENVIRONMENT "VT_ATTN_DECODE_D128=1" SKIP_RETURN_CODE 77) # The ROCm GCN-arch capability parse (BACKEND-ROCM W0). UNCONDITIONAL by design: # include/vt/rocm/rocm_arch.h is deliberately free of HIP headers so the one part # of that skeleton carrying a DECISION is gated on a CPU-only CI runner with no diff --git a/tests/vt/test_backend_cross_device.cpp b/tests/vt/test_backend_cross_device.cpp index cd3820c63..fc5d0e61b 100644 --- a/tests/vt/test_backend_cross_device.cpp +++ b/tests/vt/test_backend_cross_device.cpp @@ -1129,6 +1129,99 @@ std::vector Bf16Bits(const std::vector& src) { return out; } +TEST_CASE("paged attention at Qwen3 geometry (bf16, GQA 2, head_dim 128) matches the CPU oracle") { + // #488 / ROCM-DECODE-ATTN-D128: bf16 decode at head_dim==128 (Qwen3/Llama- + // class GQA) fell all the way to the generic PagedAttnOnline on ROCm -- + // every "fast" decode kernel was gated to d==256/512 only. Mirrors the + // Metal "Qwen3 geometry" test's shape (nblocks/bsz/hq/hkv/dh, mixed + // prefill+decode across 2 requests) so a bf16, GQA=2, d=128 case exists + // for every registered device, not just Metal. + constexpr int64_t kNBlocks = 24, kBsz = 16, kHq = 16, kHkv = 8, kDh = 128; + constexpr int64_t kNumReqs = 2; + const std::vector qsl{0, 40, 45}; // req0: 40 new (prefill); req1: 5 new + const std::vector slens{40, 71}; // req1 carries 66 context tokens + const int64_t t_total = qsl.back(); + constexpr int64_t kMaxBlocks = 6; + std::vector btab(static_cast(kNumReqs * kMaxBlocks)); + for (int64_t r = 0; r < kNumReqs; ++r) { + for (int64_t c = 0; c < kMaxBlocks; ++c) { + btab[static_cast(r * kMaxBlocks + c)] = static_cast(r * kMaxBlocks + c); + } + } + + const size_t cache_elems = static_cast(kNBlocks * kBsz * kHkv * kDh); + const std::vector qf = RandomVec(static_cast(t_total * kHq * kDh), 811, -1.5f, 1.5f); + const std::vector kf = RandomVec(cache_elems, 812, -1.5f, 1.5f); + const std::vector vf = RandomVec(cache_elems, 813, -1.5f, 1.5f); + const std::vector qb = Bf16Bits(qf), kb = Bf16Bits(kf), vb = Bf16Bits(vf); + + vt::PagedAttentionArgs args; + args.scale = 1.0f / std::sqrt(static_cast(kDh)); + args.causal = true; + args.query_start_loc_host = qsl.data(); + args.max_seq_len = 71; + + std::vector ref(qb.size(), 0); + { + vt::Backend& cpu = vt::GetBackend(DeviceType::kCPU); + Queue cq = cpu.CreateQueue(); + const Device cd{DeviceType::kCPU, 0}; + std::vector cq_v = qb, ckc = kb, cvc = vb; + std::vector cbt = btab, csl = slens, cqsl = qsl; + Tensor tq = Tensor::Contiguous(cq_v.data(), DType::kBF16, cd, {t_total, kHq, kDh}); + Tensor tkc = Tensor::Contiguous(ckc.data(), DType::kBF16, cd, {kNBlocks, kBsz, kHkv, kDh}); + Tensor tvc = Tensor::Contiguous(cvc.data(), DType::kBF16, cd, {kNBlocks, kBsz, kHkv, kDh}); + Tensor tbt = Tensor::Contiguous(cbt.data(), DType::kI32, cd, {kNumReqs, kMaxBlocks}); + Tensor tsl = Tensor::Contiguous(csl.data(), DType::kI32, cd, {kNumReqs}); + Tensor tqsl = Tensor::Contiguous(cqsl.data(), DType::kI32, cd, {kNumReqs + 1}); + Tensor to = Tensor::Contiguous(ref.data(), DType::kBF16, cd, {t_total, kHq, kDh}); + vt::PagedAttention(cq, to, tq, tkc, tvc, tbt, tsl, tqsl, args); + cpu.DestroyQueue(cq); + } + std::vector reff(ref.size()); + for (size_t i = 0; i < ref.size(); ++i) reff[i] = vt::BF16ToF32(ref[i]); + + for (DeviceType dt : RegisteredDevices()) { + if (!OpAvailable(vt::OpId::kPagedAttention, dt)) continue; + CAPTURE(DeviceName(dt)); + vt::Backend& dev = vt::GetBackend(dt); + Queue q = dev.CreateQueue(); + const Device d{dt, 0}; + + DevBufBytes dq(dev, q, qb.size() * 2), dkc(dev, q, kb.size() * 2), dvc(dev, q, vb.size() * 2), + dout(dev, q, qb.size() * 2); + dq.Upload(qb.data()); + dkc.Upload(kb.data()); + dvc.Upload(vb.data()); + DevBufI32 dbt(dev, q, btab.size()), dsl(dev, q, slens.size()), dqsl(dev, q, qsl.size()); + dbt.Upload(btab); + dsl.Upload(slens); + dqsl.Upload(qsl); + dev.Synchronize(q); + + Tensor tq = Tensor::Contiguous(dq.ptr(), DType::kBF16, d, {t_total, kHq, kDh}); + Tensor tkc = Tensor::Contiguous(dkc.ptr(), DType::kBF16, d, {kNBlocks, kBsz, kHkv, kDh}); + Tensor tvc = Tensor::Contiguous(dvc.ptr(), DType::kBF16, d, {kNBlocks, kBsz, kHkv, kDh}); + Tensor tbt = Tensor::Contiguous(dbt.ptr(), DType::kI32, d, {kNumReqs, kMaxBlocks}); + Tensor tsl = Tensor::Contiguous(dsl.ptr(), DType::kI32, d, {kNumReqs}); + Tensor tqsl = Tensor::Contiguous(dqsl.ptr(), DType::kI32, d, {kNumReqs + 1}); + Tensor to = Tensor::Contiguous(dout.ptr(), DType::kBF16, d, {t_total, kHq, kDh}); + + vt::ResetOpProviderStats(vt::OpId::kPagedAttention, dt); + vt::PagedAttention(q, to, tq, tkc, tvc, tbt, tsl, tqsl, args); + dev.Synchronize(q); + CHECK(vt::GetOpProviderStats(vt::OpId::kPagedAttention, dt).declines == 0); + + std::vector got(qb.size()); + dout.Download(got.data()); + std::vector gotf(got.size()); + for (size_t i = 0; i < got.size(); ++i) gotf[i] = vt::BF16ToF32(got[i]); + CHECK(Nmse(reff, gotf) <= kNmseTol); + + dev.DestroyQueue(q); + } +} + // Rank-3 padded-row view [T, H, D] over a [T, row_stride] f32 buffer — the // merged-qkvz slice shape the GDN/attention glue ops consume in the model. Tensor T3PaddedF32(void* p, Device d, int64_t t, int64_t h, int64_t w, From b839fe6f91a570908d40525d4061b01263a8b926 Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Sat, 15 Aug 2026 11:28:13 +0000 Subject: [PATCH 2/3] docs(BACKEND-ROCM): reconcile the d=128 spec with what actually lands (#382) Review repair on top of joral's commits, which are untouched. The spec merged as #564 with a banner reading "**Not landed.** ... No `VT_ATTN_DECODE_D128` exists in `src/vt/rocm/` on `main` -- `git log -S'VT_ATTN_DECODE_D128' -- src/vt/rocm/` is empty". Landing this PR makes all three clauses false and hands a reader a git command that returns the opposite of what the spec claims. That is the MIRROR IMAGE of the finding that blocked #564: that review stopped a spec claiming work had landed when it had not. This one would have landed the code and left the spec claiming it had not. Same defect, opposite sign, and the same consequence -- the record disagreeing with the tree, which AGENTS.md treats as the thing git exists to prevent. Four corrections, each verified against the landed diff rather than assumed: 1. The banner now says the arm lands here, and states why the section stays `## Result` rather than becoming `## Outcome`: BACKEND-ROCM remains ACTIVE, and `## Outcome` is scoped to a row reaching DONE. 2. Section 4's gate was specified as `(d == 128 && (decode_d128 || decode_wmma)) || 256 || 512`. As landed the `decode_wmma` disjunct is ABSENT, because that flag does not exist anywhere in `src/` or `tests/`. Recorded as the deliberate deviation it is, next to the forward-reference paragraph that already anticipates it. 3. "Two extra ctest registrations" -> ONE. The planned second, `VT_ATTN_DECODE_WMMA=1`, could not land for the same reason. 4. "6/6 pass, including both new flag-on registrations" -> 5/5 with the one that exists, with the superseded figure named rather than silently swapped. WHAT THIS DOES NOT FIX, and now says so in two places. The flag-ON arm has no proof it REACHES the new kernel. `OpProviderStats` counts at provider granularity, so the case's `declines == 0` is identical with the flag set and unset; and because `RegisteredDevices()` excludes kCPU, the case runs 1 test case and 0 assertions on every non-ROCm machine while exiting 0. So the registration added to make the arm "actually gated rather than silently skipped" is green on nothing everywhere this project owns hardware. Section 9's stop condition 2 -- "stop if the flag-ON arm cannot be shown to reach the new kernel; confirm selection counts, not just tokens" -- is therefore still OPEN, and the spec now says that instead of implying the ctest registrations discharged it. Closing it needs a kernel-selection counter in `rocm_paged_attn.hip` asserted to differ between the two registrations. 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 | 49 +++++++++++++++++++------- 1 file changed, 36 insertions(+), 13 deletions(-) diff --git a/.agents/specs/rocm-decode-attn-d128.md b/.agents/specs/rocm-decode-attn-d128.md index 8b964aa57..14d2bd5ac 100644 --- a/.agents/specs/rocm-decode-attn-d128.md +++ b/.agents/specs/rocm-decode-attn-d128.md @@ -99,7 +99,9 @@ synchronization. This is the same observation #382 made about the CUDA file. `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 + || decode_wmma)) || d == 256 || d == 512`. **As landed the gate omits the + `decode_wmma` disjunct**, because that flag does not exist — see the forward + reference below. 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. **Forward reference:** @@ -142,10 +144,19 @@ 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. +gates the `PagedAttnOnline` fallback. **As landed there is ONE extra ctest +registration**, `VT_ATTN_DECODE_D128=1`; the planned second, +`VT_ATTN_DECODE_WMMA=1`, does not exist because that flag does not. Same shape +as the existing `test_dense_gateup_fused_marlin_off_*` pair. + +That registration does NOT by itself prove the new kernel ran. It re-runs the +same case with the env set, and the case's only backend assertion is +`declines == 0`, which `OpProviderStats` reports at PROVIDER granularity — +identical with the flag set and unset. On any non-ROCm machine the case runs +1 test case and **0 assertions** and exits 0, so the registration is green on +nothing everywhere this project has hardware. Closing §9's stop condition 2 +needs a kernel-selection counter in `rocm_paged_attn.hip` asserted to DIFFER +between the two registrations. 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. @@ -191,8 +202,10 @@ 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, +**Correctness**, gfx1200, real hardware: `ctest -R 'rocm|cross_device'` **5/5 +pass** as landed, including the one new flag-on registration. (An earlier draft +of this section said 6/6 "including both new flag-on registrations", from the +two-registration plan above that did not land.) 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 @@ -301,12 +314,22 @@ flock "$HOME/gpu.lock" -c ' ## Result on the implementation branch (2026-08-12) -> **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`. +> **Landed by PR #767**, which carries this correction. The banner this +> paragraph replaced said "Not landed" and offered +> `git log -S'VT_ATTN_DECODE_D128' -- src/vt/rocm/` as proof — a command that +> returns the opposite once the code is in, which is how a record starts +> disagreeing with the tree. +> +> The section stays `## Result` rather than becoming `## Outcome`: `BACKEND-ROCM` +> remains `ACTIVE`, and `## Outcome` is scoped to a row reaching `DONE`. The +> arm ships **default OFF**, so nothing here is a shipped-behaviour claim. +> +> **Still owed, and NOT discharged by this landing:** the flag-ON arm has no +> proof it REACHES the new kernel. `OpProviderStats` counts at provider +> granularity, so `declines == 0` is identical with the flag set and unset, and +> the ctest registration runs 0 assertions on every non-ROCm machine. §9's stop +> condition 2 — "stop if the flag-ON arm cannot be shown to reach the new +> kernel; confirm selection counts, not just tokens" — is therefore still open. **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 From e2f611523f60a784bab65dae55b92ea5e105a9d3 Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Mon, 17 Aug 2026 14:01:29 +0000 Subject: [PATCH 3/3] docs(BACKEND-ROCM): retract the blocker, keep the oracle number, and file the gap (#382) FOLLOWING_AGENTS_PROTOCOL Record repair on top of joral's #767. No kernel behaviour changes: the two `rocm_paged_attn.hip` edits are comments. FIVE THINGS THE RECORD GOT WRONG OR WOULD HAVE LOST. 1. The spec shipped a blocker its own author had RETRACTED. Section 6 said the post-change oracle re-measure was "blocked on the container/glibc issue recorded in the WMMA spec". In the #767 comment of 2026-08-14 joral writes "That diagnosis was wrong and I'm retracting it": our binary runs inside the pinned oracle container, and the failures behind that reading were self-inflicted -- LD_LIBRARY_PATH exported container-wide, which breaks the container's own tools, plus a bind mount that silently yielded nothing and presented as a missing ELF interpreter. Landing it unchanged ships a false blocker, and a false blocker is worse than an empty record because it stops the next person from trying. Section 6 now reads "not run -- NOT blocked", cites the comment, and names what the trace still needs (decode-phase windowing on the oracle side, or it compares our decode against vLLM's model load and graph capture). The WMMA-spec cross-reference is gone. 2. The strongest evidence in the pull request existed only in a GitHub comment, which a squash merge does not carry into the tree. Both sides in the same container, oracle = vLLM 555967922 in production configuration via `vllm bench serve`, Qwen3-0.6B 1024 in / 128 out, concurrency 1, 8 prompts, warmup discarded, 3 reps: ours 42.40 ms unset, 11.67 ms with VT_ATTN_DECODE_D128=1, oracle 6.68 ms -- 6.35x to 1.75x slower. Plus the in-container-vs-native inertness control (42.53/42.79 and 11.78/12.03). Folded into spec section 5 and APPENDED to .agents/benchmark-record.md. The caveats travel with the number rather than behind it. This is per-token LATENCY with asymmetric harnesses -- the oracle over HTTP, ours in-process -- not the same-tool per-call trace AGENTS.md wants before a throughput comparison, and it does not close #488, which asks for a per-CALL kernel comparison. One board, one shape. docs/BENCHMARKS.md's ROCm axis therefore stays PENDING and the new row is marked DIRECTIONAL, not binding; it is added BESIDE the existing row rather than overwriting it, so no entry is evicted. Also recorded: at --num-prompts 2 the oracle gave 6.96 ms and 13.45 ms on consecutive reps, a ~2x spread averaging to a plausible-looking and entirely fictional number. That is why the table uses 8. 3. The spec contradicted itself on a number. 8aedd780 corrected "two flag-on ctest registrations" to ONE in section 4 item 3 and in the Test-coverage section, and missed section 7 and the result banner. Both now say one, with the reason (VT_ATTN_DECODE_WMMA does not exist). 4. The owed list had no heading a checker or a reader can find. Section 7's out-of-scope list is now a literal `## Owed` section, which is the surface AGENTS.md and check-agent-record.py read for an issue that names no owning row. 5. A KNOWN GAP HAD NO ISSUE. Filed as #1134 and appended to .agents/issue-index.md under BACKEND-ROCM. RegisteredDevices() (tests/vt/test_backend_cross_device.cpp:84-96) enumerates {kCUDA, kMETAL, kVULKAN, kXPU, kROCM} and excludes kCPU, so on a CPU-only runner -- which is what CI has -- the new head_dim-128 case reports 1 test case, 0 assertions, exit 0, for BOTH registrations. On ROCm hardware the case's only backend assertion is OpProviderStats::declines == 0, counted at PROVIDER granularity, so it is identical with the flag set and unset. The two compose: no machine in this project distinguishes the flag-ON registration from the flag-OFF one. The spec disclosed this honestly and left section 9 stop condition 2 open, which was right; what was missing is that AGENTS.md requires an issue for a known gap not fixed in flow. Searched before filing -- #463 is the unset-weights-env-var shape and does not describe the declines half, #785 is a kernel that never LAUNCHES behind a dead #if, #900 is the same family on LTX-2.5 -- so this is not a duplicate. The two comment edits are item 6 of the review, non-blocking: the kernel headers at rocm_paged_attn.hip:330 and :455 still enumerated "EPL=8 -> d=256, EPL=16 -> d=512" without the new EPL=4 case, although the top-of-file comment at :264 had been updated. Both now name d=128 and its opt-in flag. Rebased onto origin/main d1e5e9bc, which drops joral's earlier merge commit; the branch was based 39 commits behind. Gates run from this worktree with explicit SHAs. Row: BACKEND-ROCM Issue: https://github.com/mudler/vllm.cpp/issues/382 Issue: https://github.com/mudler/vllm.cpp/issues/1134 Following-Agents-Protocol: true AI-Assisted: true Assisted-by: ClaudeCode:claude-opus-5 [ClaudeCode] --- .agents/benchmark-record.md | 70 +++++++++++++++ .agents/issue-index.md | 1 + .agents/specs/rocm-decode-attn-d128.md | 118 ++++++++++++++++++++----- docs/BENCHMARKS.md | 3 +- docs/STATUS.md | 5 +- src/vt/rocm/rocm_paged_attn.hip | 6 +- 6 files changed, 179 insertions(+), 24 deletions(-) diff --git a/.agents/benchmark-record.md b/.agents/benchmark-record.md index 72283f556..81c3aa0bd 100644 --- a/.agents/benchmark-record.md +++ b/.agents/benchmark-record.md @@ -22505,3 +22505,73 @@ beside its post-fix twin rather than replaced by it, because the pair is what proves the label defect never touched the values: 819.992 s against 819.818584 s. +--- + +## BACKEND-ROCM — the `d=128` decode arm against the pinned oracle: 6.35x -> 1.75x slower on per-token decode, both sides in ONE container; the "container/glibc" blocker was RETRACTED (2026-08-14, `row/ROCM-DECODE-ATTN-D128-IMPL`, gfx1200 / RX 9060 XT / ROCm 7.2.3, PR #767, issues #382 / #488) + +Recorded here from [the #767 comment of +2026-08-14](https://github.com/mudler/vllm.cpp/pull/767#issuecomment-5295395139) +because it was the strongest evidence in that pull request and lived only in a +GitHub thread, which a squash merge does not carry into the tree. Spec: +[`specs/rocm-decode-attn-d128.md`](specs/rocm-decode-attn-d128.md) §5. + +**The blocker this row previously recorded does not exist.** PR #767's body and +the spec's §6 said the post-change oracle re-measure was blocked on a +Nix-glibc-vs-container ABI mismatch. Its author retracted that diagnosis in the +comment above: our binary runs inside the pinned oracle container, and the +earlier failures were self-inflicted — `LD_LIBRARY_PATH` exported +container-wide, which breaks the container's own tools, plus a bind mount that +silently yielded nothing and presented as a missing ELF interpreter. A false +blocker in the record is worse than no record, because it stops the next person +from trying. §6 of the spec now says "not run — NOT blocked". + +### The substitution was proved inert before it was used + +Running our binary against the CONTAINER's ROCm rather than the host's is a +substitution, so it gets a control. Qwen3-0.6B, 1024 in / 128 out, concurrency 1: + +| TPOT | native | in container | +|---|---|---| +| flag unset | 42.53 ms | 42.79 ms | +| `VT_ATTN_DECODE_D128=1` | 11.78 ms | 12.03 ms | + +### Both sides in that container, matched workload + +Qwen3-0.6B, 1024 in / 128 out, concurrency 1, **8 prompts**, warmup discarded, +**3 reps**. Oracle = vLLM `555967922` — the parity pin — in its PRODUCTION +configuration, driven by `vllm bench serve`. + +| | TPOT reps | mean | vs oracle | +|---|---|---|---| +| ours, flag unset (`PagedAttnOnline`) | 42.54 / 42.46 / 42.19 | 42.40 ms | 6.35x slower | +| ours, `VT_ATTN_DECODE_D128=1` | 11.97 / 11.38 / 11.66 | **11.67 ms** | **1.75x slower** | +| vLLM `555967922` | 6.57 / 6.90 / 6.58 | 6.68 ms | — | + +The `d=128` decode arm closes the ROCm decode gap on this shape from **6.35x to +1.75x**. It is the first oracle-relative ROCm decode figure this row has, and +the arm ships **default OFF**, so it is not a shipped-behaviour number. + +### What this is NOT, carried forward verbatim from the author's own caveats + +- **Not the same-tool per-call kernel trace.** This is LATENCY, taken with each + side's own harness. AGENTS.md wants matching traces before a throughput + comparison. `rocprofv3` is present in the container and our binary traces + under it, but the oracle side still needs decode-phase windowing — bucket + dispatches over time, take the final burst — or it compares our decode + against vLLM's model load and graph capture. Reachable, and still owed. +- **Harness asymmetry.** The oracle runs over HTTP via `vllm bench serve`; ours + is in-process. TPOT is the comparable axis. TTFT, E2EL and end-to-end + throughput carry the oracle's HTTP and tokenizer overhead and are + DIRECTIONAL ONLY. +- **Not a #488 closure.** [#488](https://github.com/mudler/vllm.cpp/issues/488) + asks for a PER-CALL kernel comparison; this is PER-TOKEN latency. The ROCm + throughput axis in `docs/BENCHMARKS.md` stays **PENDING/OPEN**. +- Single board, single model shape, one host. + +### The prompt count is load-bearing, and this is why + +At `--num-prompts 2` the oracle returned TPOT **6.96 ms and 13.45 ms on +consecutive reps** — a ~2x spread that averages to a plausible-looking and +entirely fictional number. The figures above use 8 prompts with a discarded +warmup, where both sides hold to ~±0.3 ms. A two-request rate harness is not a +measurement of this axis; it is a coin flip with a mean. diff --git a/.agents/issue-index.md b/.agents/issue-index.md index 9512f191a..da5441913 100644 --- a/.agents/issue-index.md +++ b/.agents/issue-index.md @@ -328,3 +328,4 @@ rather than merged. `scripts/check-agent-record.py` gates both. | [#837](https://github.com/mudler/vllm.cpp/issues/837) | `ROCM-GEMMA4-GETBLAS-DUALSLOT` | ROCm: dual-slot GetBlas TLS + host lifetime so two HIP streams do not share one handle | bug | | [#1117](https://github.com/mudler/vllm.cpp/issues/1117) | `LTX25-A2VID-RECIPE` | `A2VidPipelineTwoStage` (`a2vid_two_stage.py:53` @ `fd4ded7f`) has no recipe row, so `pipeline_kind = a2vid_two_stage` gets the generic table refusal (`src/vllm/model_executor/models/ltx2_pipeline.cpp:1328-1332`) naming the pair rather than the missing machinery. [#922](https://github.com/mudler/vllm.cpp/issues/922) is CLOSED and closed the audio CONDITIONING, not the recipe: a supplied take rides `distilled_two_stage`, which [`ltx25-a2v-audio-input.md`](specs/ltx25-a2v-audio-input.md):438-446 already records as a different trajectory. Measured at `daeff67f2`: `git grep -n '"a2vid' -- src include tests docs examples` returns TWO hits, both upstream anchors inside `Fail`-message assertions (`tests/vllm/multimodal/test_ltx2_video.cpp:4363,:4427`), against a control of 4 for `"one_stage"` in `include/` alone. Four differences from the recipe it rides, each read at the pin: stage 1 is CFG/STG/modality-guided and caller-configured (`:230-240`, fed from `utils/args.py:947-1006`, `--a2v-guidance-scale` defaulting to `video_guider.modality_scale` = 3.0 at `utils/constants.py:54,:64`) where `distilled_two_stage` fixes `allow_guidance_override = false`; stage 1's schedule is scheduler-derived (`:225-227`) against our fixed `DistilledSigmas()`; stage 1 is plain Euler (`:229-258` passes no `stepper`, `utils/blocks.py:526-527`) against our `kEulerAncestral` on 2.5; and the AUDIO guider is the DEFAULT positive-only one (`:237-239`, `ltx-core components/guiders.py:200-210`) rather than the params table's cfg-7.0 row. Two non-schedule facts that must not be guessed: `--audio-path` is `required=True` (`:312-317`), and the distilled LoRA rides stage 2 ALONE (`:114` against `:107`) with `--distilled-lora` `required=True` (`utils/args.py:1140-1153`). Unblocked by `Ltx2GuidedDenoise` landing at `daeff67f2` (#1092/#1102), which [`ltx25-guided-video.md`](specs/ltx25-guided-video.md) `## Owed` names this arm against. Spec [`ltx25-a2vid-recipe.md`](specs/ltx25-a2vid-recipe.md) | enhancement | | [#1118](https://github.com/mudler/vllm.cpp/issues/1118) | `LTX25-A2VID-RECIPE` | LoRA adapters fuse ONCE, at load, into ONE weight set: `src/vllm/multimodal/ltx2_video.cpp:816-820` is the only `dit_options.loras.push_back` in the tree and it runs under `if (!lora_path.empty())`, so every phase of every recipe sees the same fused `im.dit.weights`. Three upstream pipelines build TWO `DiffusionStage`s from the same checkpoint with DIFFERENT adapter sets at `fd4ded7f`: `a2vid_two_stage.py:107` against `:114-119`, `ti2vid_two_stages.py:140` against `:151`, and `ti2vid_two_stages_hq.py:154,:165` at two separate strengths (`:92-101`). `distilled.py:131` builds ONE stage set, which is why `distilled_two_stage`, `dfr` and `retake` have never needed this and no gate has ever asked. Consequence for the arm landing with #1117: that row mirrors `--distilled-lora required=True` (`utils/args.py:1140-1153`) by refusing an `a2vid_two_stage` load with no `lora_path`, so the 3-step stage-2 schedule cannot run on a checkpoint with no distilled adapter — but it CANNOT mirror upstream leaving stage 1 WITHOUT that adapter, so stage 1's guided schedule runs against base + distilled LoRA where upstream runs it against the base alone. That divergence RENDERS, and the PIXELS it renders are not upstream's: it moves the trajectory, so the frames themselves differ, while the frame count, the shapes, the sample rate and the errors are all exactly what they were — nothing in the SHAPE of the result says anything is wrong, which is why it is filed rather than left in a comment. It is not undetectable, and saying so would be the more damaging error: the instrument that WOULD see it is a real-weights comparison against upstream's own render on the same checkpoint, take and seed, upstream's stage 1 on the base weights against ours on base + distilled. Two fix shapes, neither chosen: a second `Ltx2DitWeights` per adapter set (what upstream pays, two `from_checkpoint` calls at `a2vid_two_stage.py:103,:115`) or unfused runtime LoRA selectable per phase; the first doubles resident DiT weights and the second changes the GEMM path. Bounds #1093 and #921, which need the same seam. Listed under `## Owed` in [`ltx25-a2vid-recipe.md`](specs/ltx25-a2vid-recipe.md) | enhancement | +| [#1134](https://github.com/mudler/vllm.cpp/issues/1134) | `BACKEND-ROCM` | The `VT_ATTN_DECODE_D128=1` ctest registration added by #767 cannot show the ROCm `d=128` decode arm REACHED the new kernel, and is empty off ROCm. `RegisteredDevices()` (`tests/vt/test_backend_cross_device.cpp:84-96`) enumerates `{kCUDA, kMETAL, kVULKAN, kXPU, kROCM}` and excludes `kCPU`, so on a CPU-only runner — which is what CI has — the new "Qwen3 geometry (bf16, GQA 2, head_dim 128)" case reports 1 test case, 0 assertions, exit 0, for BOTH registrations. On ROCm hardware the case's only backend assertion is `OpProviderStats::declines == 0`, and `OpProviderStats` counts at PROVIDER granularity, so it is identical with the flag set and unset; the NMSE bound passes on either kernel because the arm is correctness-complete. The two compose: there is no machine in this project on which the flag-ON registration distinguishes itself from the flag-OFF one. Disclosed in [`specs/rocm-decode-attn-d128.md`](specs/rocm-decode-attn-d128.md) §4, its `## Owed` section and its result banner, and §9 stop condition 2 is left OPEN rather than claimed discharged. Closing it needs a kernel-selection counter in `src/vt/rocm/rocm_paged_attn.hip` asserted to DIFFER between the two registrations; the CPU-runner half wants `kCPU` in `RegisteredDevices()` or a non-zero-assertion floor per [#463](https://github.com/mudler/vllm.cpp/issues/463). Not a duplicate of #463 (that is the unset-weights-env-var shape and does not describe the `declines` granularity half), #785 (a kernel that never LAUNCHES behind a dead `#if`, a code defect not a coverage one) or #900 (same family, LTX-2.5 subject) | bug | diff --git a/.agents/specs/rocm-decode-attn-d128.md b/.agents/specs/rocm-decode-attn-d128.md index 14d2bd5ac..b4f1296fb 100644 --- a/.agents/specs/rocm-decode-attn-d128.md +++ b/.agents/specs/rocm-decode-attn-d128.md @@ -41,9 +41,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'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. +differ — and they do, sharply in §5's measurement and materially in the +`## Owed` section'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 @@ -116,8 +116,8 @@ synchronization. This is the same observation #382 made about the CUDA file. `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. +of scope (`## Owed`); a `d=128` prefill call already falls through their +internal `else { goto flash_fallback; }` guards to the decode-shaped launch. ### Why default OFF @@ -127,7 +127,7 @@ 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. +ratified-tie rule — on **both** backends, and is named in `## Owed`. ### Test coverage @@ -211,22 +211,96 @@ two-registration plan above that did not land.) Full `ctest` 393 tests, op-registration gap) on an isolated build of this branch **without** this change — confirmed not caused by it. +### Against the pinned oracle, both sides in the same container (2026-08-14) + +Recorded here from the [#767 comment of +2026-08-14](https://github.com/mudler/vllm.cpp/pull/767#issuecomment-5295395139), +because a number that lives only in a pull-request thread is lost the moment the +thread is squashed. Full entry in +[`.agents/benchmark-record.md`](../benchmark-record.md). + +Running our binary against the container's ROCm rather than the host's is a +substitution, so it was proved inert first. In-container matches native, +Qwen3-0.6B, 1024 in / 128 out, concurrency 1: + +| TPOT | native | in container | +|---|---|---| +| flag unset | 42.53 ms | 42.79 ms | +| `VT_ATTN_DECODE_D128=1` | 11.78 ms | 12.03 ms | + +Both sides then in that same container, matched workload (Qwen3-0.6B, +1024 in / 128 out, concurrency 1, **8 prompts**, warmup discarded, **3 reps**), +oracle = vLLM `555967922` in its production configuration via `vllm bench serve`: + +| | TPOT reps | mean | vs oracle | +|---|---|---|---| +| ours, flag unset | 42.54 / 42.46 / 42.19 | 42.40 ms | 6.35x slower | +| ours, `VT_ATTN_DECODE_D128=1` | 11.97 / 11.38 / 11.66 | **11.67 ms** | **1.75x slower** | +| vLLM `555967922` | 6.57 / 6.90 / 6.58 | 6.68 ms | — | + +**This arm closes the decode gap from 6.35x to 1.75x on this shape.** It is the +first oracle-relative ROCm decode number this row has. + +**What it is not, carried forward from the author's own caveats.** It is +**latency, measured with each side's own harness**, not the same-tool per-call +trace AGENTS.md requires before a throughput claim — the oracle runs over HTTP +via `vllm bench serve` while ours is in-process, so TPOT is the only comparable +axis and TTFT, E2EL and end-to-end throughput carry the oracle's HTTP and +tokenizer overhead and are directional only. It does **not close #488**, which +asks for a per-call kernel comparison (§6). One board, one model shape. The +ROCm throughput axis stays **OPEN**. + +**The prompt count is load-bearing.** At `--num-prompts 2` the oracle returned +TPOT 6.96 ms and 13.45 ms on consecutive reps — a ~2x spread whose average is +plausible-looking and entirely fictional. The table above uses 8 prompts with a +discarded warmup, where both sides hold to ~±0.3 ms. + ## 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. +- **It does not close #488.** #488 reports a per-**call** gap against vLLM and + asserts no cause. This removes one cause, and §5's oracle A/B measures the + residual as a per-**token** latency, which is a different axis. The same-tool + per-call trace #488 asks for was **not run — and it is NOT blocked**; see the + [#767 comment of 2026-08-14](https://github.com/mudler/vllm.cpp/pull/767#issuecomment-5295395139). + An earlier draft of this section attributed the gap to a container/glibc ABI + mismatch. **That diagnosis was retracted by its own author**: our binary runs + inside the pinned oracle container, and the failures behind it were + self-inflicted (`LD_LIBRARY_PATH` exported container-wide, which breaks the + container's own tools, plus a bind mount that silently yielded nothing and + presented as a missing ELF interpreter). What the trace still needs is + decode-phase windowing on the oracle side — bucket dispatches over time and + take the final burst — or it compares our decode against vLLM's model load + and graph capture. `rocprofv3` is present in the container and our binary + traces under it. The work is reachable and owed, not blocked. - **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 + (`## Owed`) 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.** +cross-device test and its **one** flag-on ctest registration, and this spec. +(The planned second registration, `VT_ATTN_DECODE_WMMA=1`, does not land, +because that flag does not exist — §4 item 3 and the Test-coverage section +record the same correction.) + +## Owed + +Out of scope for this change, named here rather than left to be discovered. + +- **Proof that the flag-ON arm REACHES the new kernel** — + [#1134](https://github.com/mudler/vllm.cpp/issues/1134). `RegisteredDevices()` + (`tests/vt/test_backend_cross_device.cpp:84-96`) enumerates + `{kCUDA, kMETAL, kVULKAN, kXPU, kROCM}` and excludes `kCPU`, so on a CPU-only + runner the new case reports 1 test case, **0 assertions**, exit 0 — for both + registrations. And on ROCm, `OpProviderStats::declines` counts at PROVIDER + granularity, so it is identical with the flag set and unset. §9's stop + condition 2 is therefore OPEN, and closing it needs a kernel-selection + counter in `rocm_paged_attn.hip` asserted to DIFFER between the two + registrations. Disclosed in §4 and in the result banner; #1134 is the record + outside this file. +- **The same-tool per-call kernel trace against the oracle** (§6). Reachable, + not blocked; owed before #488 can be judged. - **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. @@ -268,11 +342,11 @@ cross-device test and its two flag-on ctest registrations, and this spec. | 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 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 default flip. The BENCHMARKS row this change adds is the later 3-rep oracle A/B, entered as DIRECTIONAL against an explicitly still-PENDING ROCm axis, not as a binding ratio. | | 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. | +| The residual #488 gap is unquantified | Partly answered, on a different axis. §5's 2026-08-14 oracle A/B measures the residual as PER-TOKEN latency — 6.35x to 1.75x slower than vLLM `555967922` — with both sides in the same container. #488 asks for a PER-CALL kernel comparison, and that is still owed, so #488 is not judged by this. The earlier "blocked on container/glibc" reason for not running it was retracted by its author and is corrected in §6. | ## 9. Stop conditions @@ -329,7 +403,9 @@ flock "$HOME/gpu.lock" -c ' > granularity, so `declines == 0` is identical with the flag set and unset, and > the ctest registration runs 0 assertions on every non-ROCm machine. §9's stop > condition 2 — "stop if the flag-ON arm cannot be shown to reach the new -> kernel; confirm selection counts, not just tokens" — is therefore still open. +> kernel; confirm selection counts, not just tokens" — is therefore still open, +> and [#1134](https://github.com/mudler/vllm.cpp/issues/1134) tracks it outside +> this file. **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 @@ -342,9 +418,11 @@ 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. +throughput on Qwen3-0.6B / 1.7B / 4B, and — measured later, against the pinned +oracle with both sides in one container — **6.35x to 1.75x slower than vLLM +`555967922`** on per-token decode latency (§5). New bf16 `d=128` GQA correctness +coverage where none existed, plus **one** flag-on ctest registration; the +planned second could not land because `VT_ATTN_DECODE_WMMA` does not exist. **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 diff --git a/docs/BENCHMARKS.md b/docs/BENCHMARKS.md index 1665a7414..820d8a02a 100644 --- a/docs/BENCHMARKS.md +++ b/docs/BENCHMARKS.md @@ -503,7 +503,8 @@ built on it rather than keeping the flattering one. | Startup latency (cold to first `/health`) | **36.51 s vs vLLM 0.25.0's 221.51 s = 6.07x** (medians of 3, 27B-NVFP4, GB10). PROVISIONAL: 3 of 6 legs contended, repeat killed by a host reboot. [Detail](../.agents/benchmark-record.md) | Uncontended 3-rep re-run on a quiet box | | Speculation depth (`ROAD-V1-D3-SPEC-K`, #81) | **PENDING.** 27B k=4 acceptance 0.875/0.750/0.618/0.507 vs control 0.000. One-window re-run 2026-08-17: void lifted, 6/6 clear. vLLM leg blocked, the oracle OOM-reboots this box ([spec](../.agents/specs/mtp-k-gt-1.md)) | the vLLM leg (our-ON == vLLM-ON) on the 27B and the whole 35B lane, then the c1/c>1 A/B at matched k + the prose vs code acceptance-vs-depth curve any depth policy needs | | Vulkan vs llama.cpp Vulkan ([`BENCH-VK-LLAMA`](../benchmarks/demo/vulkan_27b_llamacpp.json)) | 25 NATIVE (+8 GDN). **27B prefill 21.5x, a SELF-ratio not a llama.cpp one**; decode **4.36 vs 4.35, MET**, denominator SUPERSEDED (7 clean legs). Smart barriers skip 19.8%/tok, GPU -1.09 ms; e2e 8/12, unresolved. OFF. | `VK-C` coopmat A/B on Thor (`VT_VULKAN_COOPMAT=0` A/Bs it): **11.1x-32.9x** vs our UNTILED scalar kernel, not a competent GEMM. `VK-E`: llama.cpp `-DGGML_VULKAN=ON` at `237ad9b96` on dgx, SUPERSEDED, same GGUF, 3 columns | -| ROCm (`BACKEND-GATE-ROCM-VLLM` / `-SGLANG`) | **PENDING: no binding throughput number.** Runtime-green on 5 gfx archs. Gemma-3 is 48/48 exact vs two vLLM-ROCm oracles; Qwen3.5-0.8B correctness remains open | Same model, quantization, request shape and cache policy vs pinned vLLM-ROCm on one idle AMD host. Add equivalent SGLang; close correctness first ([#41](https://github.com/mudler/vllm.cpp/issues/41)) | +| ROCm (`BACKEND-GATE-ROCM-VLLM` / `-SGLANG`) | **PENDING: no binding throughput number.** The directional `d=128` decode row below is not one. Runtime-green on 5 gfx archs. Gemma-3 is 48/48 exact vs two vLLM-ROCm oracles; Qwen3.5-0.8B correctness remains open | Same model, quantization, request shape and cache policy vs pinned vLLM-ROCm on one idle AMD host. Add equivalent SGLang; close correctness first ([#41](https://github.com/mudler/vllm.cpp/issues/41)) | +| ROCm `d=128` decode arm (`BACKEND-ROCM`, [#382](https://github.com/mudler/vllm.cpp/issues/382)) | **DIRECTIONAL, not binding.** gfx1200, both sides in the pinned oracle container, 1024/128 c1, 8 prompts, 3 reps: TPOT 42.40 -> **11.67 ms** with `VT_ATTN_DECODE_D128=1`; vLLM `555967922` 6.68 ms, so **6.35x -> 1.75x** | Harnesses differ (oracle over HTTP, ours in-process) and no same-tool per-call trace exists, so [#488](https://github.com/mudler/vllm.cpp/issues/488) stays open. Owed: decode-windowed `rocprofv3` both sides | | Tenstorrent Blackhole (`BACKEND-TENSTORRENT`) | **NOT APPLICABLE (speed).** Correctness: OPT-125m STRICT 6/6 e2e on real hardware. Qwen3-0.6B has a device-specific golden and short 4-token warm smoke (~0.28 tok/s), not a completed speed run | Full 16x16 Qwen3 gate, then device-resident tensors + `ttnn::sdpa_decode` before any performance comparison. [Spec](../.agents/specs/tenstorrent-backend.md) | | Mistral-7B-v0.3 on Tenstorrent (`BACKEND-TENSTORRENT-MISTRAL`) | **PENDING (speed).** First data point: 4.26 tok/s warm, batch 1, 32 tok, single run on a P150. Not a gate, not reproduced. No vLLM ratio exists or can (no TT backend). Correctness 16/16 | Reproduce idle with a same-binary A/B before quoting. [Record](../.agents/benchmark-record.md), [spec](../.agents/specs/tenstorrent-mistral.md) | | Prompt logprobs (`SAMPLE-PROMPT-LOGPROBS`, #223) | **NO number measured, claimed or owed.** Correctness-only, CPU. Upstream ships this path explicitly unoptimized (`gpu_model_runner.py:5622-5623`); a step where no request asks is unchanged | Floor if one is ever wanted: vLLM's own `prompt_logprobs=k`, same model and prompt | diff --git a/docs/STATUS.md b/docs/STATUS.md index bbea61041..f09f36303 100644 --- a/docs/STATUS.md +++ b/docs/STATUS.md @@ -503,7 +503,10 @@ on 5 gfx archs; the APU unified-memory fix remains unverified; gfx1200 runs Gemma-3 and Qwen3 all-native, with Gemma-3 strict 48/48 against two vLLM-ROCm oracles and Qwen3 in a measured near-tie regime; Qwen3.5-0.8B GDN runs all-native but its CPU/ROCm divergence remains open; gfx1201 Gemma-4 FP8 MoE is -contributor-measured on 2x R9700 and CPU-link-verified our side; +contributor-measured on 2x R9700 and CPU-link-verified our side; a `head_dim=128` +decode arm lands opt-in behind `VT_ATTN_DECODE_D128`, default OFF, which moves +gfx1200 per-token decode from 6.35x to 1.75x slower than the pinned vLLM oracle +on one shape, a directional figure that leaves the ROCm throughput axis PENDING; [guide](ROCM.md)), inference-time CPU weight offload (`ENG-WEIGHT-OFFLOAD` ACTIVE; the config surface landed W0a (the backend enum, both sub-configs, the validator's two errors and three warnings, and the dot-anchored segment match), diff --git a/src/vt/rocm/rocm_paged_attn.hip b/src/vt/rocm/rocm_paged_attn.hip index e0e658bfe..f661a31da 100644 --- a/src/vt/rocm/rocm_paged_attn.hip +++ b/src/vt/rocm/rocm_paged_attn.hip @@ -327,7 +327,8 @@ __device__ inline void StoreRow8Bf16(__hip_bfloat16* p, int64_t base, int lane, StoreRowEplBf16<8>(p, base, lane, r); } -// BF16 Q/K/V/Out; EPL=8 → d=256, EPL=16 → d=512. No per-key block sync (decode+prefill). +// BF16 Q/K/V/Out; EPL=4 → d=128, EPL=8 → d=256, EPL=16 → d=512. No per-key block sync +// (decode+prefill). The d=128 case is opt-in behind VT_ATTN_DECODE_D128 (default OFF). template __global__ void PagedAttnDecodeOptBf16T(__hip_bfloat16* out, const __hip_bfloat16* query, const __hip_bfloat16* k_cache, const __hip_bfloat16* v_cache, @@ -452,7 +453,8 @@ __global__ void PagedAttnDecodeOptBf16T(__hip_bfloat16* out, const __hip_bfloat1 } // GQA-fused decode: one CTA per KV head; load K/V once, attend QG q-heads. -// Gemma-4: hq=16, kv=8 → QG=2. EPL=8→d=256, EPL=16→d=512. grid=(total_q, num_kv_heads). +// Gemma-4: hq=16, kv=8 → QG=2. EPL=4→d=128, EPL=8→d=256, EPL=16→d=512. +// grid=(total_q, num_kv_heads). The d=128 case is opt-in behind VT_ATTN_DECODE_D128 (default OFF). template __global__ void PagedAttnDecodeGqaBf16(__hip_bfloat16* out, const __hip_bfloat16* query, const __hip_bfloat16* k_cache, const __hip_bfloat16* v_cache,