docs(rocm): file the gfx1200 M2 near-tie investigation, link issue #269 - #273
Closed
joral wants to merge 7 commits into
Closed
docs(rocm): file the gfx1200 M2 near-tie investigation, link issue #269#273joral wants to merge 7 commits into
joral wants to merge 7 commits into
Conversation
…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]
joral
force-pushed
the
docs/rocm-gfx1200-m2-spec
branch
from
August 11, 2026 01:07
ce4ec36 to
3880903
Compare
joral
added a commit
to joral/vllm.cpp
that referenced
this pull request
Aug 11, 2026
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]
This was referenced Aug 11, 2026
joral
added a commit
to joral/vllm.cpp
that referenced
this pull request
Aug 11, 2026
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
added a commit
to joral/vllm.cpp
that referenced
this pull request
Aug 11, 2026
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 a commit
to joral/vllm.cpp
that referenced
this pull request
Aug 11, 2026
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
added a commit
to joral/vllm.cpp
that referenced
this pull request
Aug 11, 2026
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 a commit
to joral/vllm.cpp
that referenced
this pull request
Aug 11, 2026
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
added a commit
to joral/vllm.cpp
that referenced
this pull request
Aug 11, 2026
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 a commit
to joral/vllm.cpp
that referenced
this pull request
Aug 11, 2026
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
added a commit
to joral/vllm.cpp
that referenced
this pull request
Aug 11, 2026
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 a commit
to joral/vllm.cpp
that referenced
this pull request
Aug 11, 2026
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
added a commit
to joral/vllm.cpp
that referenced
this pull request
Aug 11, 2026
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]
Frenchy2k1
pushed a commit
to Frenchy2k1/vllm.cpp_sm_70
that referenced
this pull request
Aug 29, 2026
…mudler#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 mudler#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: AGENT:deepseek-v4-flash [edit bash]
Frenchy2k1
pushed a commit
to Frenchy2k1/vllm.cpp_sm_70
that referenced
this pull request
Aug 29, 2026
…spike (mudler#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 mudler#269/mudler#273-ratified distributional gate for the near-tie regime. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:deepseek-v4-flash [edit bash]
Frenchy2k1
pushed a commit
to Frenchy2k1/vllm.cpp_sm_70
that referenced
this pull request
Aug 29, 2026
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: AGENT:deepseek-v4-flash [edit bash]
Frenchy2k1
pushed a commit
to Frenchy2k1/vllm.cpp_sm_70
that referenced
this pull request
Aug 29, 2026
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: AGENT:deepseek-v4-flash [edit bash]
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
1 participant
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Row
BACKEND-ROCMBefore starting
BACKEND-ROCMumbrella)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-ROCMbugs, rocm_matmul_hipblaslt.hip:384:13: error: no matching function for call to 'hipblasGemmEx' #201 (hipblasGemmExoverload) andROCm: -O0 RmsNorm triggers a CLR HostcallListener teardown deadlock #132 (
-O0teardown deadlock) — neither reproduced on this board/toolchain(ROCm 7.2.3,
-DCMAKE_BUILD_TYPE=Release), and both are since fixed onmain(ce134e1d)..agents/roadmap_v1.mdissue table (BACKEND-ROCMrow added here);
.agents/backend-matrix.mdBACKEND-ROCMrow (ACTIVE,unchanged by this PR).
scripts/ready-for-helper.pynot run — this isn't ahelper claim on the row, it's a closed investigation writeup; nothing here
changes the row's lifecycle state.
VT_OP_PROVIDER_STATSoutput 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, theis_cuda()-vs-is_cpu()host-pointer defect, already fixed generically,confirmed correct for
kROCM);src/vllm/v1/sample/sampler.cpp(
greedy_argmax_host+ the asyncSampler::forwardpath — CPU and ROCmtake 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 greedycompletion is a genuine near-tie flip (
Parisvs. 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.shpasses for the checks relevant to thisdiff (
check-doc-checkpoint,commit-trailers,now-currentrange,check-public-doc-tables,check-env-doc, all green). The full runalso reports 8 pre-existing failures — verified unrelated by re-running
against the parent commit:
audit-live-rowsalready flags the sameabandoned-ACTIVE row (
ENG-LOAD-DIRECT-UPLOAD, untouched by this PR)before this change exists;
check-test-registrationfails becausecmakeis not onPATHin this environment at all (needs a localbuild tree this sandbox doesn't have);
role-undeclaredbecause norole was formally claimed this session.
underlying investigation's evidence (embedding-gather bit-exactness,
per-layer drift table) came from the existing
tests/vt/test_backend_cross_device.cppharness plus ad-hocinstrumentation that was reverted before this PR (not landed,
documented in the spec).
BACKEND-ROCM's lifecycle state doesnot change (stays
ACTIVE); no feature/model/backend/quant surfacemoved. This is investigation-record-only.
Speed claims
Honest gaps
(
scripts/qwen3-neartie-gap-transformers.py) was not run — it needs a realvLLM oracle install, not set up on this board. Flagged explicitly in the
spec as
PENDING, hardware/environment-blocked, not skipped silently. Theevidence 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.
closed investigation.
(
GdnStateGather/ScatterforkROCM, bit-exact cross-device test) but isnot 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.