spec(BACKEND-ROCM): port the vt::Backend graph-capture seam to hipGraph (#332) by joral · Pull Request #390 · mudler/vllm.cpp

joral · 2026-08-11T15:38:38Z

Scopes #332. Rebased onto
b580452d and re-verified there: 7 commits, 2 files, +338, linear.

Row

BACKEND-ROCM

Before starting

Issue/PR search and existing claim: opened
#332 for this specific gap.
Checked the BACKEND-ROCM umbrella ROCm (AMD GPU) backend #41 and its thread, plus ROCm gfx1200: Qwen3-0.6B produces wrong greedy output despite all-native execution #269 (gfx1200
M2/M4, my own). Distinct from the two tracked 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 of which reproduced on
this board. Also disjoint from VikashLoomba's in-flight GDN kernel work for
gfx1100 — different ops, different board.
Roadmap or matrix row: .agents/backend-matrix.md BACKEND-ROCM
(ACTIVE, unchanged by this PR); the .agents/roadmap_v1.md intake row for
ROCm: no decode-graph capture — the hipGraph seam is unimplemented, costing ~3x decode throughput vs vLLM on gfx1200 #332 is added here. scripts/ready-for-helper.py not run — this is a
spec-before-code commit, not a helper claim on the row, and nothing here moves
the row's lifecycle state.
Exact current-code and test/evidence anchors inspected: vt::Backend's six
capture virtuals (SupportsGraphCapture through DestroyGraph,
include/vt/backend.h) and their throwing base impls in src/vt/backend.cpp;
the --- CUDA-graph capture/replay block in src/vt/cuda/cuda_backend.cu
(the implementation to mirror) with its "NO cudaMalloc/cudaFree inside the region" contract comment; support_static_graph_mode() in
src/vllm/platforms/interface.h, cuda.cpp and rocm.cpp; the
stays FALSE scope note in src/vt/rocm/rocm_backend.hip;
LtWorkspace() in src/vt/rocm/rocm_matmul_hipblaslt.hip (the D1 risk);
RegisterOp(OpId::kPagedAttention, DeviceType::kROCM, ...) in
src/vt/rocm/rocm_ops.hip; the ten consumers found by
grep -rn support_static_graph_mode src/vllm/model_executor/models/; and the
CUDA backend: graph capture/replay re-executes captured ops case in
tests/vt/test_cuda_backend.cpp (the test shape to mirror).

Anchors are symbols rather than file:line throughout, because an earlier
draft went stale inside a day on a main this fast; the spec header records
why. Each was re-verified with grep -F at b580452d.

What changed

Adds .agents/specs/rocm-decode-graph.md and its intake row. Spec only — no
implementation. It scopes porting the vt::Backend graph-capture virtuals to
hipGraph: six virtuals in rocm_backend.hip mirroring the CUDA-graph
capture/replay block in cuda_backend.cu, a support_static_graph_mode()
override in platforms/rocm.cpp, and one RED-first capture/replay test. No
model-level edit is needed because every
decode-graph class already gates on
support_static_graph_mode() && SupportsGraphCapture() with no is_cuda()
anywhere, so Qwen3 dense/MoE, DeepSeek-V2/V4, Voxtral and Laguna would pick it
up unchanged. This would also be the second implementation of that seam —
Metal, Vulkan and ROCm all carry the same stays FALSE note today, so the
abstraction currently has exactly one implementation and is unproven as an
abstraction.

Evidence

scripts/agent-preflight.sh — unticked deliberately: it does not fully
pass here, and ticking it would be false. All record and trailer gates
are green; 3 NixOS-local test failures remain, traced below.
tests that cover this change: records-only, so none are owed. The spec's
§6 defines the capture/replay test the implementation PR owes, red-first,
with the two mutations that must break it.
same-change doc obligations: no lifecycle change (BACKEND-ROCM stays
ACTIVE), so no STATUS/BENCHMARKS/NOW obligation and no ## Now
section under NOW.md is still a surface every PR must write: the doc-checkpoint gate marches them into it #374.

Run inside the dev shell (nix develop .#rocm-shell -c scripts/agent-preflight.sh) at b580452d. Every record, doc and trailer gate
is green over origin/main..HEAD, including check-agent-record
(ENGINE=149 MODEL=362 QUANT=82 KERNEL=51 BACKEND=80), which is load-bearing
here: it fails on a dangling link, and that is what forced this branch to follow
the gfx1200 correctness record rather than assume it.

The 3 failures are host-environment, pre-existing on an unmodified base:
test_release_archive and test_release_metadata (our NixOS ELF binaries carry
absolute /nix/store/... RPATHs where the release-bundle validator wants
bundle-relative ones — release packaging, untouched here), and
test_agent_onboard (its fixture repo inherits host init.defaultBranch=main
against the test's hardcoded master; run with GIT_CONFIG_COUNT=1 GIT_CONFIG_KEY_0=init.defaultBranch GIT_CONFIG_VALUE_0=master and it passes).

Speed claims

This PR makes NO speed claim. It does contain numbers, but they record a
deficit against vLLM rather than a win, are labelled indicative in the
spec (single run per cell, board also driving a display, no rocprof
trace either side), and are not entered in docs/BENCHMARKS.md.
Gate 5 is what owes real numbers, and it is written as a pre-registered
prediction so the result cannot be retrofitted.
the operator ran the numbers under ${GPU_LOCK} — N/A, this is a
contributor's own board (RX 9060 XT, gfx1200).

Honest gaps

The attribution is evidenced, not profiled. The premise was tested before
committing to it: across a 7.9x span of per-step compute the gap falls
monotonically 2.99x → 1.90x → 1.46x (Qwen3-0.6B / 1.7B / 4B), and the
0.6B/1.7B pair isolates it at constant layer count (28), which a
kernel-quality-dominated gap would not do. But no same-tool trace separates
launch overhead from everything else capture changes, so D4 requires the
write-up to say "consistent with", not "because of", until one is run. The
~1.36x asymptote is provisional — a two-point fit gave ~1.54x and Qwen3-4B
landed below it, forcing the refit (spec §1).
One board, one arch. gfx1200 only. hipGraph is not arch-specific and the
four ROCm (AMD GPU) backend #41 boards are likelier-supported parts, but none has run any of this.
The other boards stay PENDING-community.
D1 is a real unverified risk, not a formality. LtWorkspace() in
rocm_matmul_hipblaslt.hip hipMallocs the hipBLASLt workspace lazily
inside the GEMM path, which is illegal mid-capture. The decode-graph pre-warm
should grow it to high-water mark first; that is asserted, not verified, and
W1 verifies it explicitly. It fails loudly at hipStreamEndCapture rather than
corrupting, which is the acceptable direction.
Gate 3 could look like a regression. Qwen3-0.6B's
'The capital of france is' is a known version-sensitive near-tie — our CPU
and ROCm backends disagree, and two real vLLM builds disagree with each other.
A captured graph may move it again. The required outcome is an honest report
against the real oracle, not a specific token.
Two adjacent defects found in passing and deliberately NOT fixed here:
get_attn_backend_priority in platforms/rocm.cpp returns {} (the M3
attention NAME registration), and the comment just above it claiming
kPagedAttention is "not registered for kROCM" — it is registered
(RegisterOp(OpId::kPagedAttention, DeviceType::kROCM, ...) in
rocm_ops.hip) and ran selected=vt-native on this board. Both want their
own issue.
Qwen3-8B was considered and is out of reach, in both directions: ~16.4 GB
of bf16 weights against 15.92 GiB of VRAM, and no quantized path exists on
ROCm (44 registered ops to CUDA's 84, none quantized, and a discrete board has
no reference tier, so an unregistered op throws rather than degrading).

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]

Adds a controlled scaling experiment run BEFORE any implementation, and rewrites the spec's expectations around what it found. Qwen3-0.6B vs Qwen3-1.7B is a clean control: identical layer count (28), so an identical kernel-launch count per decode step, with ~4x the compute per step (2x hidden, 2x intermediate). The gap against the real oracle fell 2.99x -> 1.90x. A kernel-quality-dominated gap would not move under that change; a fixed-overhead-dominated one must. The premise survives. More usefully it now BOUNDS the win. Fitting ratio = alpha*(1 + L/C) to the two points gives alpha ~= 1.54x size-independent (kernel quality, inductor fusion, Triton attention) and a launch term of ~1.45x at 0.6B falling to ~0.36x at 1.7B. So the realistic outcome is 0.6B 2.99x -> ~1.54x and 1.7B 1.90x -> ~1.54x: real, worth doing, and NOT parity. Gate 5 is re-calibrated against that pre-registered prediction rather than against closing the gap, with convergence of the two sizes as the load-bearing signal, and section 10 gains a numeric stop threshold (0.6B not below ~2.2x) fixed in advance so it cannot become a judgement call made after seeing the number. Also retracts an inference of my own. The spec previously argued the matching 2.99x on Gemma-3-1B and Qwen3-0.6B proved a backend-wide, architecture-independent overhead. The 1.7B point refutes that reasoning: Gemma-3-1B has ~2.5x the MLP compute of Qwen3-0.6B, so it should sit near 2.1x on the fitted curve and it measures 2.99x. There is an architecture-dependent component after all, the matching ratio was substantially coincidence, and the residual is recorded as new D6 rather than dropped for being inconvenient. D4 is rewritten to separate what is now evidenced from what is still unprofiled. Measurement provenance added: stock upstream checkpoints, SHA-256-verified against HF blob hashes; single run per cell on a display-driving board, which is explicitly NOT the reproduced-idle standard gate 5 demands of the real measurement. Qwen3-8B recorded as considered and not fitting (~16.4 GB bf16 against 15.92 GiB VRAM). FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: Claude:claude-sonnet-5 [ClaudeCode]

Third scaling point. Qwen3-4B measures 1.46x against the oracle (98.49 vs 143.36 tok/s, batch 8), so the gap now falls monotonically 2.99x -> 1.90x -> 1.46x across a 7.9x span of per-step compute. The 36-layer count looked like a confound against the 28-layer pair; it is not, because launches and compute both scale with layers, so the layer count cancels and L/C depends only on 1/(hidden x intermediate). All three sit on one curve. The third point falsified the earlier two-point fit: it measured 1.46x, below that fit's own 1.54x asymptote, which a curve cannot do. Refitting across three points gives alpha ~= 1.36x and beta ~= 1.65, so the predicted post-capture ratio improves from ~1.54x to ~1.36x and gate 5's prediction table gains a third row. The stop threshold stays ~2.2x. Drops Gemma-3 throughout, per direction to focus on Qwen3: the model row, the off-curve D6 discussion, the gate 3 regression bullet (Gemma has no decode-graph sibling, so capture cannot affect it), and its provenance and issue-draft mentions. Records that Qwen3-8B is out of reach in either direction: ~16.4 GB bf16 against 15.92 GiB, and no quantized path exists on ROCm, which registers 23 ops to CUDA's 84, none quantized, with no reference tier on a discrete board to fall back to. Prose tightened throughout, 401 -> 335 lines, with no technical content dropped beyond the Gemma material. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: Claude:claude-sonnet-5 [ClaudeCode]

Issue mudler#332 is open, so the three places that must agree now do: the spec header carries the link instead of PENDING, the roadmap intake table gains a BACKEND-ROCM / perf row, and the PR body owes it when the PR is opened. W0's inline draft-issue text is replaced by a link to the filed issue -- it was scaffolding for writing mudler#332 and is redundant now that mudler#332 exists. 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]

main runs 75-207 commits/day (measured 2026-08-07..11), which makes a file:line anchor wrong within days rather than months. This spec already proved it: inside one day an earlier draft drifted +4 lines in cuda_backend.cu, +7 in backend.h, and rocm_ops.hip's kPagedAttention registration moved 92 -> 148, and a factual op count went stale 23 -> 44. Converts all 35 line anchors to greppable symbols -- function signatures, the RegisterOp call, the test-case title, distinctive comments, and in one place the grep a reader should run instead (`grep -rn support_static_graph_mode src/vllm/model_executor/models/`, which finds all 10 decode-graph consumers). The port-map table already named every CUDA function, so its line numbers were pure noise and are simply dropped. Verified: every replacement anchor resolves with grep -F against the current tree. The header now says anchors are symbols and asks editors to keep it that way, so this does not quietly rot back. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: Claude:claude-sonnet-5 [ClaudeCode]

…ment The symbol-anchor conversion introduced a defect the anchor check could not see. Replacing a `file:line` with a phrase that itself contained the search text nested one anchor inside another: `vt::Backend`'s capture virtuals (`vt::Backend`'s capture virtuals (...)) `grep -F` passed on it, because a duplicated anchor still resolves in the tree. Existence was the wrong property to verify; the prose was never read. The same edit pushed ~20 lines past 80 columns in a file wrapped at 80, invisible to that check for the same reason. Both are repaired here, and the line count rises 330 -> 337 because rewrapping adds lines while removing words. D4 restated §1 rather than adding to it — the monotonic fall, the 0.6B/1.7B constant-layer control and the 1.54x -> 1.36x refit are all established two sections earlier. It now carries only what is unique to it: no same-tool trace exists, W4 must say "consistent with" rather than "because of", and ~1.36x is the next target rather than a floor. Net 2786 -> 2749 words. Deliberately small: an implementer works from this file, so the capture contract, D1's allocator trap and the pre-registered gate-5 thresholds all stay whole. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [claude-code]

…ec (#332) Lands Justin Card's work from PR #390. Merged --no-ff rather than squashed so `git log` keeps `author=Justin Card` across his seven commits; a squash would re-author the result to whoever pressed the button, and a squash of a multi-commit PR also concatenates every commit's trailer block and reds the gate. Spec only: ports the vt::Backend graph-capture seam onto hipGraph for BACKEND-ROCM, scoping issue #332. Two files, +338, no product source. Nothing was wrong with this branch. It was two commits behind, its checks had not finished under the queue backlog (#274), and every gate passes on the merged tree: check-agent-record, check-doc-checkpoint, check-public-doc-tables, check-now-current, check-commit-trailers and check-pr-size, with scripts/agent-preflight.sh rc=0. It is landed on that evidence rather than on a CI green that the runner queue is not currently able to produce. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [Claude Code]

localai-bot · 2026-08-11T22:05:49Z

Landed on main as 31a2b493 via row/BACKEND-ROCM-DECODE-GRAPH, with your authorship preserved — merged --no-ff so git log keeps author=Justin Card across all seven commits rather than re-authoring them to whoever pressed the button.

Nothing was wrong with this PR. It was two commits behind and its checks simply had not finished under the runner-queue backlog (#274). Every gate passes on the merged tree: check-agent-record, check-doc-checkpoint, check-public-doc-tables, check-now-current, check-commit-trailers and check-pr-size, with scripts/agent-preflight.sh rc=0.

It is landed on that evidence rather than on a CI green, because the queue is not currently able to produce one — per AGENTS.md the operator's own gate run is the authority and CI is convenience. Worth stating plainly rather than implying it went green.

joral mentioned this pull request Aug 11, 2026

ROCm: no decode-graph capture — the hipGraph seam is unimplemented, costing ~3x decode throughput vs vLLM on gfx1200 #332

Open

joral force-pushed the row/ROCM-DECODE-GRAPH branch from a706350 to 94d0a31 Compare August 11, 2026 17:39

joral added 7 commits August 11, 2026 17:30

joral force-pushed the row/ROCM-DECODE-GRAPH branch from 94d0a31 to 94c7fe3 Compare August 11, 2026 21:31

localai-bot merged commit c628559 into mudler:main Aug 11, 2026
12 of 14 checks passed

Navigation Menu

spec(BACKEND-ROCM): port the vt::Backend graph-capture seam to hipGraph (#332) - #390

spec(BACKEND-ROCM): port the vt::Backend graph-capture seam to hipGraph (#332)#390
localai-bot merged 7 commits into
mudler:mainfrom
joral:row/ROCM-DECODE-GRAPH

joral commented Aug 11, 2026

Uh oh!

Uh oh!

localai-bot commented Aug 11, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

Conversation

joral commented Aug 11, 2026

Row

Before starting

What changed

Evidence

Speed claims

Honest gaps

Uh oh!

Uh oh!

localai-bot commented Aug 11, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants