feat(rocm): implement the vt::Backend graph-capture seam on hipGraph (W1, #332)

[joral]

Implements W1 of #332, the
first work item under the spec that landed in #390. 8 commits, 5 files,
+375/-14.

Updated 2026-08-13 in response to review: the flake.nix rocwmma fix
(previously bundled here) is now split into
#638 against #444, and the W3
throughput refutation described below is now also recorded in the spec itself
(.agents/specs/rocm-decode-graph.md, D7) rather than living only in this
description — see localai-bot's review
comment
for the full ask.

Row

BACKEND-ROCM

Before starting

• Issue/PR search and existing claim:
  #332, whose spec
  (.agents/specs/rocm-decode-graph.md) landed on main via spec(BACKEND-ROCM): port the vt::Backend graph-capture seam to hipGraph (#332) #390. W1 is its
  §9 second item. Distinct from rocm_matmul_hipblaslt.hip:384:13: error: no matching function for call to 'hipblasGemmEx' #201 and ROCm: -O0 RmsNorm triggers a CLR HostcallListener teardown deadlock #132, and from the in-flight gfx1100
  GDN kernel work — different ops, different board.
  #444 was opened/used here
  originally but its fix has since moved to
  #638 (see Honest gaps).
• Roadmap or matrix row: BACKEND-ROCM stays ACTIVE; this moves no
  lifecycle state, 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.
• Anchors inspected: the --- CUDA-graph capture/replay block in
  src/vt/cuda/cuda_backend.cu (the mirror source) and its
  "NO cudaMalloc/cudaFree inside the region" contract comment; the six
  virtuals SupportsGraphCapture through DestroyGraph in
  include/vt/backend.h with their throwing base impls in src/vt/backend.cpp;
  LtWorkspace() in src/vt/rocm/rocm_matmul_hipblaslt.hip;
  support_static_graph_mode() in platforms/interface.h, cuda.cpp,
  rocm.cpp; the CUDA backend: graph capture/replay re-executes captured ops
  case in tests/vt/test_cuda_backend.cpp.

What changed

The six vt::Backend capture virtuals implemented against hipGraph in
rocm_backend.hip, mirroring cuda_backend.cu call for call, plus the
RED-first test that makes the mirror checkable. SupportsGraphCapture() becomes
true for kROCM. No model-level behaviour changes — RocmPlatform still
inherits support_static_graph_mode() == false, so no decode-graph class
engages; flipping that is W2. This is the seam's second implementation;
Metal and Vulkan still carry the stays FALSE note.

VT_BENCH_PROFILE_CONTROL is deliberately not ported (spec §2).

Evidence

• scripts/agent-preflight.sh — unticked deliberately. The declared
  W1 gates (spec §7 gates 1, 2, 7) are green; 3 NixOS-local host failures
  remain, all confirmed on an unmodified base: test_release_archive and
  test_release_metadata (our ELF binaries carry /nix/store/... RPATHs
  where the release-bundle validator wants bundle-relative), and
  test_agent_onboard (fixture inherits host init.defaultBranch=main
  against the test's hardcoded master).
• tests that cover this change: two new cases in
  tests/vt/test_rocm_backend.cpp — graph capture/replay re-executes captured ops and a pre-warmed GEMM captures and replays.
• same-change doc obligations: none; no feature/model/backend/quant surface
  moves and the row's lifecycle state is unchanged.

RED first, for the intended reason

Before implementation the case failed on
CHECK(rocm.SupportsGraphCapture()) == false, and BeginCapture threw
vt: graph capture unsupported on this backend from src/vt/backend.cpp — the
base impl, confirming kROCM inherited the throwing default rather than
overriding it.

Gates, rerun on the rebased head

$ cmake --build build-hip -j8                                     # gate 1
0 warnings, 0 errors (full HIP build, gfx1200)

$ ctest --test-dir build-hip -R 'rocm|cross_device'               # gate 2
100% tests passed, 0 tests failed out of 4

$ ./build-hip/tests/test_rocm_backend
test cases: 9 | 9 passed    assertions: 1058 | 1058 passed

$ python3 scripts/check-agent-record.py                           # gate 7
agent record OK: ENGINE=151 MODEL=362 QUANT=82 KERNEL=51 BACKEND=81
$ python3 scripts/check-device-leakage.py
OK (DSR 32 == baseline 32, ratchet holds)

Gate 1's other half: a non-HIP build still object-compiles the test file clean
via vllm_rocm_platform_syntax_check.

Mutation evidence

Every guarantee the tests pin was deleted or inverted in a scratch copy and the
focused test rerun, then the tree restored byte-for-byte. The load-bearing one:
making EndCaptureGraph execute once and freeze its result — the "replays a
snapshot" defect — leaves step 4 passing and step 5 failing
(CHECK(17 == 34)). That is why step 5 exists: a graph that replays baked
values instead of re-executing over persistent buffers would pass a naive test
and then serve stale inputs on every decode step after the first.

An independent reviewer ran 10 further mutations against the same head and
returned PASS. Three findings came back and all are addressed in-branch: a
[[nodiscard]] attribution error in a code comment, the DestroyGraph
asymmetry now recorded as D6, and a coverage gap where returning a stale
exec_ from EndCaptureGraph passed because both paths captured an identical
copy — the handle path now targets a distinct buffer, and re-running that
mutation turns it red. The full review write-up is available on request.

Speed claims

• This PR makes NO speed claim. It changes no decode path — the model-level
  gate stays false. And the row's performance rationale has since been
  refuted — see "What W3 measured" below. Read that before weighing this
  change on throughput grounds.
• the operator ran the numbers under ${GPU_LOCK} — N/A, contributor's own
  board (RX 9060 XT, gfx1200).

What W3 measured — the rationale for this row does not hold

This PR is W1. W2 and W3 have since run on a follow-up branch, and the
pre-registered performance gate failed: capture engaged correctly
(confirmed by 126 replays over 128 decode steps and a host/user/sys split
showing it removes no host CPU work) but moved throughput only 0-2% across
three model sizes, not the predicted ~1.36x convergence — indistinguishable
from noise. Not a ceiling claim: the gap is unexplained, not irreducible.

This is now recorded as D7 in .agents/specs/rocm-decode-graph.md, with
the full A/B table, the replay-count and CPU-split evidence, the correction
history (an initial +3.2% for Qwen3-0.6B was overstated — an independent
reviewer traced it to one low outlier and the honest figure is 0-2%), and the
ordered next hypotheses (same-tool rocprof trace; where the host CPU time
goes). Reading it there instead of duplicating it here so this description
and the spec cannot drift apart.

What this PR still delivers, unaffected: the seam has a second
implementation, mirrored call-for-call, mutation-tested, and behaviour-neutral
across four models and a 6.7x parameter range (Qwen3-0.6B/1.7B/4B dense and
Qwen3.5-0.8B GDN hybrid, capture ON vs OFF byte-identical). D1 and D6 are
documented. What died is the reason it was built, not the code. If the project
would rather not carry an unused capability, that is a fair call to make on
this PR — say so and I will close it.

Honest gaps

• One board, one arch. gfx1200 only, and no CI job builds ROCm at all
  (grep -i 'hip\|rocm' .github/workflows/ci.yml matches only comment prose),
  so nothing here is machine-verified. The four ROCm (AMD GPU) backend #41 boards have not run it.
• flake.nix no longer carries the rocwmma fix. It was here because
  adding rocwmma to the shell is what made W1's gates runnable in the first
  place, but it was unrelated scope — now split into
  #638 against ROCm: main does not build for gfx1200/gfx1201 — rocm_paged_attn.hip includes rocwmma unconditionally on arch, not availability #444. That fix
  is still only the narrow half of ROCm: main does not build for gfx1200/gfx1201 — rocm_paged_attn.hip includes rocwmma unconditionally on arch, not availability #444: it fixes one shell, while every
  other rocWMMA-free gfx12 environment still fails the same way.
• D1 was verified and is worse than the spec predicted. Capturing a cold
  GEMM fails on hipMalloc, hipFree and hipblasCreate — two lazy
  initialisations, not one. All fail loudly; none corrupts. A pre-warm of the
  identical GEMM clears both. Recorded in the spec, with the consequence for W2:
  the pre-warm must reach handle creation, not just workspace growth.
• The new GEMM test pins the mitigation, not the hazard, deliberately.
  Asserting that the cold path throws would forbid a future capture-safe
  allocator.
• D6 (new, from review): DestroyGraph/EndCapture Check() the destroy
  where CUDA's leg silently ignores it, so a teardown that destroys an exec
  still in flight would throw on HIP and succeed on CUDA. Not reachable in W1 —
  the model path is not engaged and the tests synchronise first — and recorded
  in the spec as a W2 constraint.
• Nothing consumes this yet. Until W2 flips
  support_static_graph_mode(), the capture path is reachable only from tests.

[localai-org-maint-bot]

@mudler This implementation is technically well evidenced, but it needs a maintainer call before merge: the follow-on measurements refute the throughput rationale and no production path consumes the seam yet. I also recommend splitting the unrelated flake.nix rocWMMA dependency fix into #444 so this PR remains the graph-capture seam only. If we still want the unused backend capability for parity/future work, the capture/replay tests and mutation evidence look solid.

[localai-bot]

Reviewed this as part of a sweep over the open external-contributor PRs. Short version: the code is in good shape and I could not fault the isolation — what is holding it is one records item and one decision that is not mine to make.

What I verified. Isolation checks out three independent ways: every decode-graph call site ANDs SupportsGraphCapture() with support_static_graph_mode(), which RocmPlatform never overrides, so flipping the backend bool alone engages nothing; src/vt/rocm/rocm_backend.hip appears zero times in a CPU build's compile_commands.json; and there is zero drift on every file this touches despite the 194-commit gap, including the seam it implements against. The mirror of cuda_backend.cu reads faithfully call for call, and HIP's 5-arg hipGraphInstantiate is the right mapping for CUDA's 3-arg form.

The new bit-rot test has real teeth. I changed the EndCaptureGraph seam signature in a scratch copy and it failed to compile with a clear error, then restored and verified by checksum — so this genuinely guards seam drift on machines with no AMD hardware, which is most of them. CI is fully green: 16 pass, 5 release-lane skips, zero failures, and agent-record passes here.

Worth saying plainly: reporting the W3 result the way you did — measuring instead of assuming, then retracting the earlier prediction against your own interest — is exactly the behaviour this protocol is trying to produce. That is not a mark against the PR.

What needs to change before it can merge.

The W3 refutation lives only in the PR description. .agents/specs/rocm-decode-graph.md on this branch still opens with the 2.99x / 1.90x / 1.46x table and the intact "launch overhead dominates" argument, and §7's gate-5 prediction still stands. Under "History is git", a PR body is not a record — so as things stand, the next agent to claim W2 or W3 opens that spec, reads a live rationale, and re-derives a closed negative result.

You have already written the analysis; it just needs to live in the file. A D7 under §8 (or a dated note under §1) carrying the A/B table, the 126-replays-over-128-steps proof that capture actually engaged, the user/sys/wall split, and the next hypotheses would close it.

Two smaller things while you are in there:

1. src/vllm/platforms/rocm.cpp:67-69 still carries the W0 placeholder comment saying hipGraph "is not implemented" and that the address-baking concern must be resolved before the bool can flip. After this PR both statements are out of date — you updated the twin comment in rocm_backend.hip:22-28 but not this one. It is the exact file a W2 implementer opens first.
2. Issue ROCm: main does not build for gfx1200/gfx1201 — rocm_paged_attn.hip includes rocwmma unconditionally on arch, not availability #444, referenced by the flake.nix rocWMMA commit, is open but absent from the roadmap issue table and every .agents/ file. One row under BACKEND-ROCM satisfies the three-place rule.

The open decision. Because the row's own performance rationale is refuted, whether to carry the capability at all is a product call for @mudler rather than something a reviewer should settle. I have put the case to him with a recommendation to take it — the argument being that the runtime cost today is zero (nothing engages it) and that the vt::Backend graph-capture seam gaining a second real implementation has value independent of the decode number, since a seam with one implementation is a guess. But that is his call, and I would rather say so than have you guess at it.

One thing to be aware of either way: #523's per-call hipMalloc/hipFree and hipStreamSynchronize are illegal under hipGraph stream capture, so if the capability is kept, that PR has to be reconciled with it before this can ever be switched on.

Not asking for a rebase — the merge is clean and nothing you touch has moved.

[joral]

Thanks for the thorough review — addressed all items:

1. D7 added to .agents/specs/rocm-decode-graph.md with the full W3 A/B table, the 126-replays-over-128-steps proof capture engaged, the user/sys/wall split, and the ordered next hypotheses. §1's fit paragraph and §7's gate-5 prediction table are marked SUPERSEDED rather than rewritten. Also corrected the Qwen3-0.6B figure from the original +3.2% to the review-verified honest 0-2% while I was in there.
2. rocm.cpp:67-69 updated — no longer says hipGraph capture 'is not implemented'.
3. ROCm: main does not build for gfx1200/gfx1201 — rocm_paged_attn.hip includes rocwmma unconditionally on arch, not availability #444 added to the roadmap issue table under BACKEND-ROCM.
4. flake.nix split out into build(nix): add rocwmma to the ROCm dev shell so gfx12 builds again (#444) #638 against ROCm: main does not build for gfx1200/gfx1201 — rocm_paged_attn.hip includes rocwmma unconditionally on arch, not availability #444, per your recommendation — this PR is graph-capture-seam-only again.