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/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/.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 8b964aa57..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 @@ -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:** @@ -114,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 @@ -125,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 @@ -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,29 +202,105 @@ 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 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. @@ -255,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 @@ -301,12 +388,24 @@ 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, +> 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 @@ -319,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 f73993197..f661a31da 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]); @@ -311,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, @@ -436,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, @@ -1638,7 +1656,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 +1698,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 +1725,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 +1907,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 +1936,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,