docs(rocm): file the gfx1200 M2 near-tie investigation, link issue #269 by joral · Pull Request #273 · mudler/vllm.cpp

joral · 2026-08-10T14:49:47Z

Row

BACKEND-ROCM

Before starting

Issue/PR search and existing claim: checked issue ROCm (AMD GPU) backend #41 (BACKEND-ROCM umbrella)
and its thread; found VikashLoomba's in-flight GDN kernel work for gfx1100
(all ten GDN ops, staged, 68/68 cross-device assertions, awaiting a
claim-ready row) — disjoint op set (classic dense attention vs. GDN
layers) and disjoint board from this investigation, confirmed no overlap.
Opened ROCm gfx1200: Qwen3-0.6B produces wrong greedy output despite all-native execution #269 for this specific defect since it's distinct from the two
already-tracked BACKEND-ROCM bugs, rocm_matmul_hipblaslt.hip:384:13: error: no matching function for call to 'hipblasGemmEx' #201 (hipblasGemmEx overload) and
ROCm: -O0 RmsNorm triggers a CLR HostcallListener teardown deadlock #132 (-O0 teardown deadlock) — neither reproduced on this board/toolchain
(ROCm 7.2.3, -DCMAKE_BUILD_TYPE=Release), and both are since fixed on
main (ce134e1d).
Roadmap or matrix row: .agents/roadmap_v1.md issue table (BACKEND-ROCM
row added here); .agents/backend-matrix.md BACKEND-ROCM row (ACTIVE,
unchanged by this PR). scripts/ready-for-helper.py not run — this isn't a
helper claim on the row, it's a closed investigation writeup; nothing here
changes the row's lifecycle state.
Exact current-code and test/evidence anchors inspected: VT_OP_PROVIDER_STATS
output for the real model run; tests/vt/test_backend_cross_device.cpp
(existing embedding-gather coverage confirmed the real failing row is
bit-exact); dense_attn_block.h:181 (ResidentWeight, the
is_cuda()-vs-is_cpu() host-pointer defect, already fixed generically,
confirmed correct for kROCM); src/vllm/v1/sample/sampler.cpp
(greedy_argmax_host + the async Sampler::forward path — CPU and ROCm
take different code paths by default); src/vllm/model_executor/models/qwen3.cpp
(ForwardLayers' per-layer loop).

What changed

Adds .agents/specs/rocm-gfx1200-m2-correctness.md, closing out the RX 9060 XT
(gfx1200) M2 attempt on Qwen3ForCausalLM (Qwen3-0.6B): the wrong greedy
completion is a genuine near-tie flip (Paris vs. a content-free space token,
~0.07% of the max logit) from ordinary bf16/reduction-order drift
accumulating uniformly across all 28 decoder layers — not a kernel defect.
Links the new issue #269 (Row: BACKEND-ROCM) from the roadmap's issue table.
No source code changes.

Evidence

scripts/agent-preflight.sh passes for the checks relevant to this
diff (check-doc-checkpoint, commit-trailers, now-current range,
check-public-doc-tables, check-env-doc, all green). The full run
also reports 8 pre-existing failures — verified unrelated by re-running
against the parent commit: audit-live-rows already flags the same
abandoned-ACTIVE row (ENG-LOAD-DIRECT-UPLOAD, untouched by this PR)
before this change exists; check-test-registration fails because
cmake is not on PATH in this environment at all (needs a local
build tree this sandbox doesn't have); role-undeclared because no
role was formally claimed this session.
tests that cover this change: none directly — it's a docs-only PR. The
underlying investigation's evidence (embedding-gather bit-exactness,
per-layer drift table) came from the existing
tests/vt/test_backend_cross_device.cpp harness plus ad-hoc
instrumentation that was reverted before this PR (not landed,
documented in the spec).
same-change doc obligations: N/A — BACKEND-ROCM's lifecycle state does
not change (stays ACTIVE); no feature/model/backend/quant surface
moved. This is investigation-record-only.

Speed claims

This PR makes NO speed claim.

Honest gaps

The project's own decisive teacher-forcing near-tie check
(scripts/qwen3-neartie-gap-transformers.py) was not run — it needs a real
vLLM oracle install, not set up on this board. Flagged explicitly in the
spec as PENDING, hardware/environment-blocked, not skipped silently. The
evidence gathered without it (identical top-5 logit candidate set on both
backends, sub-1% uniform per-layer drift present from layer 0, zero
NaN/Inf anywhere) is treated as strong enough to close this out, but it is
not the rigorous oracle-backed confirmation the project's own methodology
prefers.
This PR does not fix anything — there is nothing to fix. It documents a
closed investigation.
A real op was ported and gated during this same investigation
(GdnStateGather/Scatter for kROCM, bit-exact cross-device test) but is
not included in this PR — held back per project policy (feature code
needs the row/PR review process, this PR is docs-only) and to avoid
colliding with VikashLoomba's in-flight GDN kernel work for the same op
family on gfx1100.

localai-bot · 2026-08-10T22:37:17Z

Merged. This is exactly the kind of record I want in the tree, and the reason is the Qwen3-0.6B call.

You had a split between our CPU and ROCm backends on one prompt and you did not file it as a defect. Instead you ran two independent real vLLM-ROCm oracles and found they disagree with each other — 0.19.1 matching our ROCm, the exact pin matching our CPU — each internally deterministic at K=5. That's the strongest available evidence that the reference itself doesn't hold still on that input, so a token-exact bar cannot close it. Under the ratified distributional gate that's the correct disposition, and it's why this merges as filed rather than as a bug fix. Plenty of people would have shipped that as "ROCm is broken" or, worse, quietly tuned something until it matched.

The Gemma-3-1B-it result is the real headline: 48/48 tokens identical against both oracles, including one built from this project's own pinned commit 555967922 inside rocm/vllm-dev:base. That's a genuine M4, no caveat — and building the pin from source inside the same base image vLLM's own Dockerfile.rocm uses is a much better answer to the pip-vs-NixOS problem than fighting the host.

The docs/ROCM.md point about discrete boards is one I want to keep visible: with no reference tier, M2's bar is unchanged but its mechanism isn't — zero fallbacks under VT_OP_PROVIDER_STATS=1 is the evidence, precisely because a fallback can't exist to hide behind. That framing generalises to every future discrete target.

The rocm-shell devShell also lands as-is. The comment explaining why clr alone serves as ROCM_PATH — that CMake's enable_language(HIP) hard-requires a single root holding lib/cmake/hip-lang/hip-lang-config.cmake, while nixpkgs ships each component as its own store path — is going to save the next NixOS contributor a genuinely miserable afternoon, and the symlink overlay is idempotent and cached rather than a build step.

Issue #269 is now placed in the roadmap issue table, and backend-matrix.md / docs/FEATURES.md reflect gfx1200 M0–M4.

…dler#269 Adds .agents/specs/rocm-gfx1200-m2-correctness.md, closing out the RX 9060 XT (gfx1200) M2 attempt on Qwen3-0.6B: the wrong greedy completion is a genuine near-tie flip (`Paris` vs a content-free space token, ~0.07% of the max logit) from ordinary bf16/reduction-order drift accumulating uniformly across all 28 decoder layers, not a kernel defect. Embedding gather was checked bit-exact against the real safetensors weights at the real failing row; the per-layer L2-norm drift table shows no localized break, present from layer 0 onward. Links the new issue mudler#269 (Row: BACKEND-ROCM) from the roadmap's issue table, per the "no work without an open issue" rule. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: Claude:claude-sonnet-5 [ClaudeCode]

nixpkgs ships each ROCm component as its own store path instead of one /opt/rocm-shaped prefix, so a plain `nix-shell -p rocmPackages.*` needs manual ROCM_PATH/CPATH/CMAKE_PREFIX_PATH surgery every session. This folds that into one shellHook: ROCM_PATH is clr's own output (the only one with lib/cmake/hip-lang, which CMake's enable_language(HIP) requires at a single fixed root); hipBLAS/hipBLASLt/hipblas-common get merged into a small writable overlay at $XDG_CACHE_HOME/vllm-cpp-rocm-overlay once, idempotently, since hipblas.h includes hipblas-common's header but clang's --rocm-path probe does not reach it. Verified end to end on real gfx1200 (RX 9060 XT) hardware: configure, full build (1145/1145), and `ctest -R 'rocm|cross_device'` (3/3) all pass through this shell alone, no manual env exports. Local-only: no CI, no maintainer machine has ROCm, so nothing else changes. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: Claude:claude-sonnet-5 [ClaudeCode]

Adds a section to .agents/specs/rocm-gfx1200-m2-correctness.md broadening the RX 9060 XT (gfx1200) M2 attempt from Qwen3-0.6B to google/gemma-3-1b-it (via the ungated unsloth/gemma-3-1b-it mirror, SHA-256-verified against the upstream blob hash). Same board, same build-hip, same --device cpu vs --device auto (auto resolves to kROCM, confirmed, with VT_OP_PROVIDER_STATS=1 showing every op selected=vt-native — no reference-tier fallback, none exists on this discrete board). Greedy, the same 6-prompt battery the SACRED Gemma-3 gate uses: 48/48 tokens identical, CPU vs ROCm. This is the first exercise of the gemma (1+w) RmsNorm code path (sandwich norms + QK-norm), GeGLU, and dual per-layer RoPE-theta routing on this board, and unlike Qwen3-0.6B's single near-tie flip, it is a clean unanimous match across the whole battery. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: Claude:claude-sonnet-5 [ClaudeCode]

docs/ROCM.md's §4 hardware table and M2 milestone text were written before gfx1200 (RX 9060 XT, issue mudler#269) joined the effort, and only described M2's unified-memory/reference-tier mechanism (gfx1151/gfx1103). Adds a gfx1200 row to §4 and a discrete-board paragraph to §5's M2 description: on a dGPU there is no reference tier to fall back to, so VT_OP_PROVIDER_STATS=1 reporting selected=vt-native on every op, with zero fallbacks, is itself the evidence of full native kernel coverage. Cites both measured results on this board: Qwen3-0.6B (one near-tie flip, closed as ordinary bf16 drift) and Gemma-3-1B-it (clean 48/48, no near-tie) — .agents/specs/rocm-gfx1200-m2-correctness.md. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: Claude:claude-sonnet-5 [ClaudeCode]

…outlier A real vLLM-ROCm oracle (AMD's prebuilt gfx120X-targeted Docker image) is now available on this board, which the earlier Outcome explicitly could not get (recorded PENDING, hardware/environment-blocked). It settles the Qwen3-0.6B near-tie the earlier investigation could only eyeball: the oracle produces ' 1000000' deterministically (5/5 runs), matching our ROCm backend exactly and diverging from our CPU backend's 'Paris'. The prior conclusion had it backwards — it treated CPU's answer as the implicit ground truth and ROCm's divergence as an acceptable coin-flip loss. It is CPU, not ROCm, that diverges from the real reference on this input. Also upgrades the Gemma-3-1B-it finding from CPU-vs-ROCm agreement to a real M4 result: the oracle matches both of our backends byte-identically across the same 6-prompt battery, full text and token ids. Rewrites rocm-gfx1200-m2-correctness.md's Outcome in place (correction prepended, original mechanism analysis kept and recontextualized, new Real-oracle verification section with the reproducible Docker recipe) and updates docs/ROCM.md's §4 table and M2 section to match. Flags the CPU-side divergence as a small separate open question, not a ROCm-row defect. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: Claude:claude-sonnet-5 [ClaudeCode]

Rewrites rocm-gfx1200-m2-correctness.md fresh instead of layering a second correction on the first: Gemma-3-1B-it is unambiguous M4 (48/48 vs TWO independent real oracles, a prebuilt AMD image and a from-source build at this project's exact pinned commit 555967922, both on this board). Building that exact-pin oracle turned out to be the interesting part -- vLLM's own requirements/rocm.txt carries no torch pin, rocm/vllm-dev:base already ships ROCm 7.2.3 (matches this board's native build exactly) with gfx1200 in PYTORCH_ROCM_ARCH, and the whole build compiled clean in ~6.5 minutes. Qwen3-0.6B is the real finding: the exact-pin oracle disagrees with the 0.19.1 oracle on the one near-tie prompt (0.19.1 matches our ROCm, the exact pin matches our CPU), each internally deterministic (K=5, 5/5). That is direct, measured proof this specific input is a genuine cross-version near-tie in the reference implementation itself, not a bug in either of our backends -- a stronger and more honest conclusion than either of the two prior single-oracle reads (this spec's own history: original investigation had no oracle at all; a same-day first correction generalized from the 0.19.1-only result and was itself wrong). Reframes the row's stance around the project's own near-tie-distributional-gate methodology instead of a strict token-exact bar on this one prompt. Updates docs/ROCM.md's gfx1200 table row and M2/M4 section to match. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: Claude:claude-sonnet-5 [ClaudeCode]

Adds gfx1200 (RX 9060 XT, mudler#269) to backend-matrix.md's BACKEND-ROCM row, independently of the mudler#41 four-board set: M0/M1 MET, M2 MET via the native-kernel path, M4 MET for Gemma3ForCausalLM (gemma-3-1b-it) against two independent real vLLM-ROCm oracles including this project's exact pinned commit, plus the Qwen3-0.6B cross-version near-tie finding. Links the new gfx1200 M2/M4 spec alongside the existing W0/unified-memory records. docs/FEATURES.md's two ROCm rows previously said only 'model e2e pending' -- updates both to note gfx1200 cleared that bar for one model, keeping the still-pending state honest for the other boards. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: Claude:claude-sonnet-5 [ClaudeCode]

Files .agents/specs/rocm-decode-graph.md, spec-before-code for implementing decode-graph capture on the ROCm backend. No implementation here. The measurement that motivates it: on gfx1200 against a real vLLM-ROCm oracle at this project's own pin (555967922, production config, graphs on), we are 2.99x slower on BOTH Gemma-3-1B-it (146.43 vs 438.09 tok/s) and Qwen3-0.6B (184.69 vs 552.65 tok/s). The same ratio to three significant figures across two unrelated architectures points at a backend-wide launch overhead, not kernel quality, and vLLM captures 51 piecewise + 35 full hipGraphs on this same board while SupportsGraphCapture() is false on ours. Scoped deliberately small: six virtuals in rocm_backend.hip mirroring cuda_backend.cu:194-286, one platform flag, one RED-first test. No model edit is needed because every decode-graph class already gates generically on support_static_graph_mode() && SupportsGraphCapture() with no is_cuda() anywhere, so Qwen3 dense/MoE, DeepSeek-V2/V4, Voxtral and Laguna pick it up for free. States plainly that Gemma-3 gains NOTHING until a Gemma3DecodeGraph exists (it does not, on any backend including CUDA). Carries the risks that were found by reading the code rather than assumed. D1: rocm_matmul_hipblaslt.hip:243-255 hipMallocs the hipBLASLt workspace lazily in the GEMM path, which is illegal mid-capture and is the most likely way this fails. D2: the known version-sensitive Qwen3-0.6B near-tie may move and must be re-checked against the real oracle rather than re-baselined. D4: "graph capture is the 2.99x" is an untraced hypothesis and the spec refuses to launder it into a conclusion. Stacks on mudler#273, where the linked gfx1200 correctness spec lands; based on main instead, check-agent-record fails with a dangling link. Issue is PENDING and filing it is W0, with draft text in the spec. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: Claude:claude-sonnet-5 [ClaudeCode]

@joral

…#273) External contribution from Justin Card (@joral). Docs, records and a Nix devShell only; no product code. What it lands: - .agents/specs/rocm-gfx1200-m2-correctness.md: the full gfx1200 (RX 9060 XT) investigation. Gemma-3-1B-it is 48/48 token-identical against TWO independent real vLLM-ROCm oracles -- AMD's prebuilt rocm/vllm gfx120X image, and a from-source build of this project's own pinned commit 555967922 inside rocm/vllm-dev:base. - The Qwen3-0.6B split is recorded as a genuine near-tie, NOT a defect: the two oracles disagree with EACH OTHER (0.19.1 matches our ROCm, the exact pin matches our CPU), each internally deterministic at K=5. The reference does not hold still on that input, so a token-exact bar cannot close it. That is the correct disposition under the ratified distributional gate, and it is the reason this merges as filed rather than as a bug fix. - docs/ROCM.md: discrete boards have no reference tier, so M2's mechanism is zero-fallback all-native execution under VT_OP_PROVIDER_STATS=1 rather than parity-behind-a-fallback. Records gfx1200 MET on that bar. - flake.nix: a rocm-shell devShell. nixpkgs ships ROCm as per-component store paths, but CMake's enable_language(HIP) hard-requires one root holding lib/cmake/hip-lang; clr provides it, and hipBLAS/hipBLASLt/ hipblas-common are symlink-merged into one idempotent cache overlay. - Issue #269 placed in the roadmap issue table; backend-matrix and docs/FEATURES.md reflect gfx1200 M0-M4. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: Claude:claude-opus-5 [ClaudeCode]

…spike (#41) The Qwen3.5-0.8B M2 blocker on discrete ROCm (throws at op 77; the reference tier cannot install on a dGPU) needs ten GDN ops with no ROCm kernel. This commit is the spike spec and the claim; implementation PRs follow one family per PR in throw-order (state-IO, conv, norm-gate, core, fused). Pre-claim validation on 4x gfx1100 / ROCm 7.14 (scratch, recorded in the spec): all ten kernels hand-translated from the CUDA donors and 68/68 standalone checks green (state paths and dtype conversions bit-exact vs an independent host reference); the drop-in TU compiles with the exact production flags; five red-first cross-device cases run 16/16 against the real library with unregistered ops skipping correctly. The key design finding: registering the full indexed set flips IndexedGdnOpsNative(kROCM) true, so the model takes the CUDA lane's device-resident path with zero model or runner edits. M2 disposition follows the #269/#273-ratified distributional gate for the near-tie regime. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: pi:kimi-k3 [pi]

localai-bot · 2026-08-11T11:48:27Z

Landed on main as e9011c8b5b72 — the gfx1200 M2/M4 record and the rocm-shell devShell is on main and verified present.

Closing manually: it landed as a rebased copy of this branch, so the head SHA here never became reachable from main and GitHub could not auto-close it. Nothing is outstanding; the PR was just left showing OPEN.

Thanks for the contribution.

Files .agents/specs/rocm-decode-graph.md, spec-before-code for implementing decode-graph capture on the ROCm backend. No implementation here. The measurement that motivates it: on gfx1200 against a real vLLM-ROCm oracle at this project's own pin (555967922, production config, graphs on), we are 2.99x slower on BOTH Gemma-3-1B-it (146.43 vs 438.09 tok/s) and Qwen3-0.6B (184.69 vs 552.65 tok/s). The same ratio to three significant figures across two unrelated architectures points at a backend-wide launch overhead, not kernel quality, and vLLM captures 51 piecewise + 35 full hipGraphs on this same board while SupportsGraphCapture() is false on ours. Scoped deliberately small: six virtuals in rocm_backend.hip mirroring cuda_backend.cu:194-286, one platform flag, one RED-first test. No model edit is needed because every decode-graph class already gates generically on support_static_graph_mode() && SupportsGraphCapture() with no is_cuda() anywhere, so Qwen3 dense/MoE, DeepSeek-V2/V4, Voxtral and Laguna pick it up for free. States plainly that Gemma-3 gains NOTHING until a Gemma3DecodeGraph exists (it does not, on any backend including CUDA). Carries the risks that were found by reading the code rather than assumed. D1: rocm_matmul_hipblaslt.hip:243-255 hipMallocs the hipBLASLt workspace lazily in the GEMM path, which is illegal mid-capture and is the most likely way this fails. D2: the known version-sensitive Qwen3-0.6B near-tie may move and must be re-checked against the real oracle rather than re-baselined. D4: "graph capture is the 2.99x" is an untraced hypothesis and the spec refuses to launder it into a conclusion. Stacks on mudler#273, where the linked gfx1200 correctness spec lands; based on main instead, check-agent-record fails with a dangling link. Issue is PENDING and filing it is W0, with draft text in the spec. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: Claude:claude-sonnet-5 [ClaudeCode]

mudler#273's content reached main directly (closed, not merged, but its commits are on main), so the stacking premise is gone: this branch is rebuilt on current upstream/main and carries only the decode-graph spec plus its mudler#332 intake row. The old branch also had to be dropped rather than merged -- main's docs/FEATURES.md ROCm rows have since been improved by other work ("W0-W1 verified on 5 gfx archs; classic-dense AND GDN-hybrid e2e run all-native") and the roadmap intake table was re-sorted and grown, so replaying the stale versions would have REGRESSED both. roadmap_v1.md is taken from main wholesale with only the mudler#332 row reapplied, per the keyed-record rule. Main moved under the spec's anchors and they are corrected here, which is the anchor-drift hazard the spec itself warns about, caught by re-verification rather than by a reviewer: the cuda_backend.cu capture block shifted +4 lines (194-286 -> 198-290, with every port-map row moved to match), backend.h's virtuals +7 (181-195 -> 188-202), and rocm_ops.hip's kPagedAttention registration moved 92 -> 148. One factual claim also moved: ROCm now registers 44 ops, not 23, so the spec says 44. The claim that MATTERS is unchanged and re-verified -- none of the 44 is quantized, so Qwen3-8B remains out of reach on this board in either direction. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: Claude:claude-sonnet-5 [ClaudeCode]

joral added 7 commits August 10, 2026 21:04

joral force-pushed the docs/rocm-gfx1200-m2-spec branch from ce4ec36 to 3880903 Compare August 11, 2026 01:07

This was referenced Aug 11, 2026

spec(rocm): BACKEND-ROCM GDN kernel family — spike + claim (gfx1100, issue #41) #334

Merged

ROCm (AMD GPU) backend #41

Open

localai-bot mentioned this pull request Aug 11, 2026

feat(rocm): MoeRouterTopK bf16 logits + MoeSiluMul — the MoE-path ops the GDN slice surfaced (issue #41) #348

Merged

1 task

localai-bot closed this Aug 11, 2026

Navigation Menu

docs(rocm): file the gfx1200 M2 near-tie investigation, link issue #269 - #273

docs(rocm): file the gfx1200 M2 near-tie investigation, link issue #269#273
joral wants to merge 7 commits into
mudler:mainfrom
joral:docs/rocm-gfx1200-m2-spec

joral commented Aug 10, 2026

Uh oh!

localai-bot commented Aug 10, 2026

Uh oh!

localai-bot commented Aug 11, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

Conversation

joral commented Aug 10, 2026

Row

Before starting

What changed

Evidence

Speed claims

Honest gaps

Uh oh!

localai-bot commented Aug 10, 2026

Uh oh!

localai-bot commented Aug 11, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants