From 2fd64af99b24ed0c242e0481856b525e82f3929f Mon Sep 17 00:00:00 2001 From: Justin Card Date: Tue, 11 Aug 2026 12:30:49 -0400 Subject: [PATCH 1/8] feat(rocm): implement the vt::Backend graph-capture seam on hipGraph (W1, #332) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit The six capture virtuals mirrored from cuda_backend.cu call for call, plus the RED-first test that makes the mirror checkable. SupportsGraphCapture() goes true for kROCM; the model-level path stays gated on Platform's support_static_graph_mode(), which is W2. RED first, for the intended reason: before this change 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. Green after: 10/10 assertions across the stored-graph and handle paths. Both mutations from the spec's §6 were run in a scratch copy and both turn the test red: * ReplayGraph made a no-op -> 3 failures, dst never updates. * EndCaptureGraph executing once and freezing the result (the "replays a snapshot" defect) -> step 4 PASSES, step 5 FAILS (CHECK(17 == 34): dst holds pattern A where it must hold B). The second is why step 5 exists. A graph that replays baked values instead of re-executing over persistent buffers passes a naive test and then silently serves stale inputs on every decode step after the first. rocm_backend.hip was restored byte-for-byte afterwards (sha256 4a6abb18... before and after). D1 VERIFIED rather than trusted, and it is worse than the spec assumed. A cold captured GEMM fails with hipMalloc AND hipFree "operation not permitted when stream is capturing" AND hipblasCreate INTERNAL_ERROR -- there are TWO lazy initialisations in that path, not just LtWorkspace's growth. All fail loudly; none corrupts. Pre-warming the identical GEMM clears both and the captured graph replays numerically correct. The new test pins the MITIGATION, not the hazard: asserting that the cold path throws would forbid a future capture-safe allocator. The spec records the consequence for W2 -- the pre-warm must reach handle creation, not only workspace growth. DEVIATION from the CUDA source: hipGraphExecDestroy and hipGraphDestroy are [[nodiscard]] where their cudaGraph* counterparts are not, so the mirrored bare calls do not compile. They are Check()ed, matching Free/FreePinned in the same file. VT_BENCH_PROFILE_CONTROL is deliberately not ported (spec §2). Gates: (1) HIP gfx1200 -Werror 0 warnings; non-HIP build object-compiles the test file clean via vllm_rocm_platform_syntax_check; CPU ctest 380/382, the two failures pre-existing and structurally unrelated (rocm_backend.hip is compiled 0 times in a CPU build, and test_safetensors links nothing this touches) -- test_safetensors is an mmap-residency measurement, test_serve_low_tools wants shellcheck. (2) ctest -R 'rocm|cross_device' 3/3. (7) check-agent-record, check-public-doc-tables, check-test-registration, check-device-leakage all OK; DSR holds at 32, no new device predicate. Issue: https://github.com/mudler/vllm.cpp/issues/332 Spec: .agents/specs/rocm-decode-graph.md FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [claude-code] --- .agents/specs/rocm-decode-graph.md | 24 +++++ src/vt/rocm/rocm_backend.hip | 88 ++++++++++++++++- tests/vt/test_rocm_backend.cpp | 152 +++++++++++++++++++++++++++++ 3 files changed, 260 insertions(+), 4 deletions(-) diff --git a/.agents/specs/rocm-decode-graph.md b/.agents/specs/rocm-decode-graph.md index 60fec59e1..a0e4fe12f 100644 --- a/.agents/specs/rocm-decode-graph.md +++ b/.agents/specs/rocm-decode-graph.md @@ -280,6 +280,30 @@ capture, which should grow `cap` to its high-water mark. *If it does not:* direction. **W1 verifies this rather than trusting it**; a shape appearing only at capture time would be a latent, board-specific trap. +**VERIFIED in W1 on gfx1200, and it is worse than written above — there are TWO +lazy initialisations, not one.** Capturing a cold `MatmulBT` ([1,2048] x +[2048,2048]^T) fails with all three of: + +```text +vt rocm: hipMalloc: operation not permitted when stream is capturing +vt rocm: hipFree: operation not permitted when stream is capturing +vt rocm: matmul: hipblasCreate: hipblas 6 (INTERNAL_ERROR) +``` + +`hipblasCreate` was not anticipated: the handle initialises on first use in the +same path, so a pre-warm must cover **handle creation as well as workspace +growth**. Both fail loudly; neither corrupts. Running the identical GEMM once +beforehand clears both, and the captured graph then replays numerically correct +(0.409606 vs 0.409600 expected, f32). Pinned by `ROCm backend: a pre-warmed GEMM +captures and replays` in `test_rocm_backend.cpp`, which asserts the MITIGATION +rather than the hazard — a future capture-safe allocator would be an +improvement, and a test forbidding it would ratchet the wrong way. + +*Consequence for W2:* the decode-graph pre-warm must reach every GEMM shape the +captured region will execute, including the first call that creates the handle. +A shape reached only under capture still aborts, so W2's gate 4 replay count is +what proves the pre-warm was complete. + **D2 — the Qwen3-0.6B near-tie may move.** Covered by gate 3. Separate because it is the one outcome that could look like a regression while being nothing of the kind, and quietly re-baselining a golden is what "never weaken a checker" diff --git a/src/vt/rocm/rocm_backend.hip b/src/vt/rocm/rocm_backend.hip index 9239f313e..a4e44636b 100644 --- a/src/vt/rocm/rocm_backend.hip +++ b/src/vt/rocm/rocm_backend.hip @@ -19,10 +19,14 @@ // backend, written out by hand so a reader can diff the two. // // SCOPE / what is NOT here, stated plainly: -// * `SupportsGraphCapture()` stays FALSE. hipGraph exists and is the eventual -// mapping, but a graph that captures a wrong stream is a silent correctness -// bug (see .agents/ on CUDA-graph capture baking stack addresses), so it is -// not something to write blind. +// * `SupportsGraphCapture()` is TRUE as of BACKEND-ROCM W1 — see the hipGraph +// capture/replay block below and .agents/specs/rocm-decode-graph.md. The +// concern that kept it false (a capture that bakes addresses instead of +// re-executing over persistent buffers is a SILENT correctness bug) is now +// an assertion rather than a worry: the mutate-src-then-replay step in +// tests/vt/test_rocm_backend.cpp fails if replay ever returns a snapshot. +// What is NOT claimed here is the model-level path — Platform's +// support_static_graph_mode() still gates that, and flipping it is W2. // * The async-output primitives (AllocPinned / events) inherit the vt::Backend // defaults. src/vt/backend.cpp documents those as already correct for // unified-memory backends; on a discrete AMD card they are correct but @@ -218,6 +222,79 @@ class RocmBackend final : public Backend { Check(hipStreamSynchronize(AsStream(q)), "hipStreamSynchronize"); } + // --- hipGraph capture/replay (BACKEND-ROCM W1) ------------------------------ + // The hipGraph mirror of cuda_backend.cu's capture block, call for call. Every + // name below is a long-stable HIP runtime API with the same signature and + // semantics as its CUDA counterpart — the property that lets upstream compile + // csrc/ for both through hipify. + // + // Capture contract, identical to the CUDA one (the caller must honour it, else + // capture aborts): + // * every op in the region runs ASYNC on THIS stream (no Synchronize, no + // null-stream work, no host<->device blocking copies); + // * NO hipMalloc/hipFree inside the region — the scratch pool must be + // pre-warmed so every allocation is a pool hit. hipBLASLt's workspace is + // the known hazard here: LtWorkspace() in rocm_matmul_hipblaslt.hip grows + // it lazily inside the GEMM path, so a shape first seen during capture + // allocates and invalidates the capture. It fails loudly at + // hipStreamEndCapture rather than corrupting, which is the acceptable + // direction (spec .agents/specs/rocm-decode-graph.md, D1); + // * captured pointers stay valid + FIXED across replays (persistent + // buffers); only their CONTENTS change between replays. + // + // hipStreamCaptureModeThreadLocal matches CUDA's choice deliberately: it makes + // an illegal op during capture a loud failure on THIS thread rather than a + // process-wide mode change. + bool SupportsGraphCapture() const override { return true; } + void BeginCapture(Queue& q) override { + Check(hipStreamBeginCapture(AsStream(q), hipStreamCaptureModeThreadLocal), + "hipStreamBeginCapture"); + } + void EndCapture(Queue& q) override { + hipGraph_t graph = nullptr; + Check(hipStreamEndCapture(AsStream(q), &graph), "hipStreamEndCapture"); + if (exec_ != nullptr) { + Check(hipGraphExecDestroy(exec_), "hipGraphExecDestroy"); + exec_ = nullptr; + } + Check(hipGraphInstantiate(&exec_, graph, nullptr, nullptr, 0), + "hipGraphInstantiate"); + Check(hipGraphDestroy(graph), "hipGraphDestroy"); + } + void Replay(Queue& q) override { + Check(hipGraphLaunch(exec_, AsStream(q)), "hipGraphLaunch"); + } + + // Multi-graph handle API (batched decode graph): instantiate the just-captured + // stream graph and hand the exec back as an opaque handle the caller owns and + // selects per padded batch size. Unlike EndCapture, nothing is stored here. + void* EndCaptureGraph(Queue& q) override { + hipGraph_t graph = nullptr; + Check(hipStreamEndCapture(AsStream(q), &graph), "hipStreamEndCapture"); + hipGraphExec_t exec = nullptr; + Check(hipGraphInstantiate(&exec, graph, nullptr, nullptr, 0), + "hipGraphInstantiate"); + Check(hipGraphDestroy(graph), "hipGraphDestroy"); + return reinterpret_cast(exec); + } + // NOT ported: the CUDA leg's VT_BENCH_PROFILE_CONTROL block, which drives + // cudaProfilerStart/Stop around a chosen replay. A rocprofiler equivalent is + // later work and deliberately out of W1's scope (spec §2). + void ReplayGraph(Queue& q, void* graph) override { + Check(hipGraphLaunch(reinterpret_cast(graph), AsStream(q)), + "hipGraphLaunch"); + } + // hipGraphExecDestroy is [[nodiscard]] where cudaGraphExecDestroy is not, so + // the CUDA leg's bare call would not compile here. Checked rather than + // discarded, matching Free/FreePinned above: a failing destroy is a leak this + // backend would rather report than swallow. + void DestroyGraph(void* graph) override { + if (graph != nullptr) { + Check(hipGraphExecDestroy(reinterpret_cast(graph)), + "hipGraphExecDestroy"); + } + } + // THE LOAD-BEARING BOOL. This is what decides whether the portable CPU // reference tier installs for kROCM (include/vt/op_provider.h:197-201): a CPU // kernel dereferences HOST pointers, so it is correct only where host and @@ -263,6 +340,9 @@ class RocmBackend final : public Backend { bool managed_alloc_ = false; int major_ = 0; int minor_ = 0; + // Single-graph slot for the EndCapture/Replay pair. The handle API + // (EndCaptureGraph) stores nothing here — its caller owns the exec. + hipGraphExec_t exec_ = nullptr; }; // Registers every visible AMD GPU at its own Device{kROCM, i} slot, mirroring diff --git a/tests/vt/test_rocm_backend.cpp b/tests/vt/test_rocm_backend.cpp index a59ff0a52..97d715217 100644 --- a/tests/vt/test_rocm_backend.cpp +++ b/tests/vt/test_rocm_backend.cpp @@ -284,3 +284,155 @@ TEST_CASE("the ROCm platform self-registers and is selected over CPU") { // is the reminder to update it deliberately. CHECK(rocm.get_attn_backend_priority({}).empty()); } + +// Mirrors "CUDA backend: graph capture/replay re-executes captured ops" in +// tests/vt/test_cuda_backend.cpp assertion for assertion. Same shape, same +// persistent-buffer contract, hipGraph underneath. +// +// Step 5 is the load-bearing one and the reason this test exists. Replaying +// must RE-EXECUTE the captured copy over the persistent buffers, not replay a +// snapshot of their contents — that is precisely how a decode graph picks up +// each new token's inputs. A capture that bakes values instead of addresses +// passes step 4 and fails step 5, which is the silent-correctness-bug shape +// rocm_backend.hip's scope note warns about. +TEST_CASE("ROCm backend: graph capture/replay re-executes captured ops") { + if (NoDevice()) return; + Backend& rocm = vt::GetBackend(DeviceType::kROCM); + CHECK(rocm.SupportsGraphCapture()); + + Queue q = rocm.CreateQueue(); + constexpr size_t kBytes = 64 * 1024; + + // Allocated ONCE; the pointers stay fixed across every replay below. Only + // their CONTENTS change — the capture contract. + void* src = rocm.Alloc(kBytes); + void* dst = rocm.Alloc(kBytes); + + std::vector pattern_a(kBytes, 0x11); + std::vector pattern_b(kBytes, 0x22); + std::vector back(kBytes, 0); + + rocm.Copy(q, src, pattern_a.data(), kBytes); + rocm.Memset(q, dst, 0, kBytes); + rocm.Synchronize(q); + + // Recorded, NOT executed: dst must still be zero after EndCapture. + rocm.BeginCapture(q); + rocm.Copy(q, dst, src, kBytes); + rocm.EndCapture(q); + rocm.Copy(q, back.data(), dst, kBytes); + rocm.Synchronize(q); + CHECK(back.front() == 0x00); + + // Replay #1 -> pattern A. Proves the graph ran at all. + rocm.Replay(q); + rocm.Synchronize(q); + rocm.Copy(q, back.data(), dst, kBytes); + rocm.Synchronize(q); + CHECK(back.front() == 0x11); + CHECK(back.back() == 0x11); + + // Mutate src in place (SAME address) -> replay must observe the new contents. + rocm.Copy(q, src, pattern_b.data(), kBytes); + rocm.Replay(q); + rocm.Synchronize(q); + rocm.Copy(q, back.data(), dst, kBytes); + rocm.Synchronize(q); + CHECK(back.front() == 0x22); + CHECK(back.back() == 0x22); + + // Handle variant — the path decode graphs actually take, since they keep one + // exec per padded batch size rather than a single stored graph. + rocm.Copy(q, src, pattern_a.data(), kBytes); + rocm.Memset(q, dst, 0, kBytes); + rocm.Synchronize(q); + + rocm.BeginCapture(q); + rocm.Copy(q, dst, src, kBytes); + void* graph = rocm.EndCaptureGraph(q); + REQUIRE(graph != nullptr); + + rocm.ReplayGraph(q, graph); + rocm.Synchronize(q); + rocm.Copy(q, back.data(), dst, kBytes); + rocm.Synchronize(q); + CHECK(back.front() == 0x11); + + rocm.Copy(q, src, pattern_b.data(), kBytes); + rocm.ReplayGraph(q, graph); + rocm.Synchronize(q); + rocm.Copy(q, back.data(), dst, kBytes); + rocm.Synchronize(q); + CHECK(back.front() == 0x22); + CHECK(back.back() == 0x22); + + rocm.DestroyGraph(graph); + rocm.Free(src); + rocm.Free(dst); + rocm.DestroyQueue(q); +} + +// The capture contract's allocation clause, asserted rather than assumed +// (.agents/specs/rocm-decode-graph.md D1). hipBLASLt sizes its workspace lazily +// inside the GEMM path — LtWorkspace() in rocm_matmul_hipblaslt.hip does +// hipFree+hipMalloc when a shape needs more than the current high-water mark — +// and hipblasCreate() likewise initialises on first use. Both are illegal +// mid-capture. +// +// MEASURED on gfx1200 during W1: capturing a cold GEMM fails loudly, with +// `hipMalloc: operation not permitted when stream is capturing` (and hipFree, +// and hipblasCreate INTERNAL_ERROR) — never silent corruption. Running the +// identical GEMM once beforehand grows the workspace and creates the handle, so +// the in-capture call is a pure pool hit. +// +// This case pins the MITIGATION, not the hazard: it asserts that a pre-warmed +// GEMM captures and replays correctly. Deliberately not asserting that the cold +// path throws — a future capture-safe allocator would be an improvement, and a +// test that forbade it would be a ratchet in the wrong direction. +TEST_CASE("ROCm backend: a pre-warmed GEMM captures and replays") { + if (NoDevice()) return; + Backend& rocm = vt::GetBackend(DeviceType::kROCM); + if (!rocm.SupportsGraphCapture()) return; + + Queue q = rocm.CreateQueue(); + const Device dev{DeviceType::kROCM, 0}; + constexpr int kM = 1, kN = 2048, kK = 2048; // a decode-shaped GEMM + + const std::vector ha(kM * kK, 0.01f); + const std::vector hb(kN * kK, 0.02f); + const float expect = 0.01f * 0.02f * static_cast(kK); + + void* da = rocm.Alloc(ha.size() * sizeof(float)); + void* db = rocm.Alloc(hb.size() * sizeof(float)); + void* dc = rocm.Alloc(kM * kN * sizeof(float)); + rocm.Copy(q, da, ha.data(), ha.size() * sizeof(float)); + rocm.Copy(q, db, hb.data(), hb.size() * sizeof(float)); + rocm.Synchronize(q); + + Tensor ta = Tensor::Contiguous(da, DType::kF32, dev, {kM, kK}); + Tensor tb = Tensor::Contiguous(db, DType::kF32, dev, {kN, kK}); + Tensor tc = Tensor::Contiguous(dc, DType::kF32, dev, {kM, kN}); + + // Pre-warm: grows LtWorkspace's cap and creates the hipBLAS handle. + vt::MatmulBT(q, tc, ta, tb); + rocm.Synchronize(q); + rocm.Memset(q, dc, 0, kM * kN * sizeof(float)); + rocm.Synchronize(q); + + rocm.BeginCapture(q); + vt::MatmulBT(q, tc, ta, tb); + rocm.EndCapture(q); + + rocm.Replay(q); + rocm.Synchronize(q); + std::vector back(kM * kN, 0.0f); + rocm.Copy(q, back.data(), dc, back.size() * sizeof(float)); + rocm.Synchronize(q); + CHECK(back.front() == doctest::Approx(expect).epsilon(0.01)); + CHECK(back.back() == doctest::Approx(expect).epsilon(0.01)); + + rocm.Free(da); + rocm.Free(db); + rocm.Free(dc); + rocm.DestroyQueue(q); +} From 654c2642273825f18f0ab27c566ea8ed0d1fc73f Mon Sep 17 00:00:00 2001 From: Justin Card Date: Tue, 11 Aug 2026 14:15:39 -0400 Subject: [PATCH 2/8] fix(rocm): correct nodiscard attribution in DestroyGraph comment (LOW-1) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Finding: LOW-1 in review-rocm-decode-graph-w1.md claimed the comment above DestroyGraph in rocm_backend.hip is "factually wrong" because a bare hipGraphExecDestroy call compiled with 0 warnings, and asked for the compile-failure rationale to be deleted. That remediation is itself wrong and was NOT applied — checked and confirmed: $ hipcc -std=c++20 -Werror -Wall -Wextra --offload-arch=gfx1200 -c probe.hip error: ignoring return value of type 'hipError_t' declared with 'nodiscard' attribute [-Werror,-Wunused-value] Under this project's actual build flags (-Werror; the review's probe most likely omitted it, or compiled below C++17 where __HIP_NODISCARD expands to nothing) a bare call to any HIP function returning hipError_t is a hard error. What WAS wrong is the attribution, not the claim of a compile failure. Verified directly against $ROCM_PATH/include/hip/hip_runtime_api.h (ROCm 7.2.3, /nix/store/8nqihd79gkvmjpc3i9icz6y7pc6rf4ma-clr-7.2.3): line 293: #define __HIP_NODISCARD [[nodiscard]] (guarded __cplusplus >= 201703L) line 305: typedef enum __HIP_NODISCARD hipError_t { ... } hipError_t; line 8321: hipError_t hipGraphExecDestroy(hipGraphExec_t graphExec); -- no attribute of its own [[nodiscard]] sits on the hipError_t return TYPE, not on hipGraphExecDestroy's declaration — so it applies to every HIP API returning hipError_t under C++17+, not to hipGraphExecDestroy specifically. CUDA's cudaError_t carries no such attribute, which is why the mirrored bare calls compile on the CUDA leg and not here. The same imprecision (source-attributed to hipGraphExecDestroy/hipGraphDestroy specifically, rather than to hipError_t generally) is present in 150b8f40's commit message. 150b8f40 is the reviewed, immutable head and is not amended; this commit is the correction of record. Remediation applied: rewrote the comment to attribute [[nodiscard]] correctly to hipError_t via __HIP_NODISCARD, name the -Werror consequence, and contrast with cudaError_t — while keeping the Free/FreePinned consistency rationale for Check() unweakened, since that design-choice justification was never disputed. Issue: https://github.com/mudler/vllm.cpp/issues/332 Spec: .agents/specs/rocm-decode-graph.md FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-sonnet-5 [claude-code] --- src/vt/rocm/rocm_backend.hip | 13 +++++++++---- 1 file changed, 9 insertions(+), 4 deletions(-) diff --git a/src/vt/rocm/rocm_backend.hip b/src/vt/rocm/rocm_backend.hip index a4e44636b..614317f33 100644 --- a/src/vt/rocm/rocm_backend.hip +++ b/src/vt/rocm/rocm_backend.hip @@ -284,10 +284,15 @@ class RocmBackend final : public Backend { Check(hipGraphLaunch(reinterpret_cast(graph), AsStream(q)), "hipGraphLaunch"); } - // hipGraphExecDestroy is [[nodiscard]] where cudaGraphExecDestroy is not, so - // the CUDA leg's bare call would not compile here. Checked rather than - // discarded, matching Free/FreePinned above: a failing destroy is a leak this - // backend would rather report than swallow. + // hipError_t itself — not hipGraphExecDestroy specifically — is + // [[nodiscard]] at C++17+ ($ROCM_PATH/include/hip/hip_runtime_api.h:293-305, + // __HIP_NODISCARD on the `typedef enum ... hipError_t` line; the function + // declaration at line ~8321 carries no attribute of its own). CUDA's + // cudaError_t has no such attribute, so the mirrored bare call compiles + // there and not here: under this project's -Werror it is a hard error + // (-Werror,-Wunused-value). Checked rather than discarded, matching + // Free/FreePinned above: a failing destroy is a leak this backend would + // rather report than swallow. void DestroyGraph(void* graph) override { if (graph != nullptr) { Check(hipGraphExecDestroy(reinterpret_cast(graph)), From 485a23c10797e18476c022df5a96b72df0127d4d Mon Sep 17 00:00:00 2001 From: Justin Card Date: Tue, 11 Aug 2026 14:15:48 -0400 Subject: [PATCH 3/8] spec(BACKEND-ROCM): record D6, the DestroyGraph/EndCapture Check() asymmetry (INFO-1) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Finding: INFO-1 in review-rocm-decode-graph-w1.md — DestroyGraph and the prior-exec_ destroy inside EndCapture (src/vt/rocm/rocm_backend.hip) both Check() hipGraphExecDestroy's return and throw on failure, where the CUDA leg silently ignores it. A decode-graph teardown or column-change recapture (W2+) that destroys/recaptures an exec still in flight would throw on HIP where CUDA succeeds silently. Not a W1 defect: the model-level path is not engaged (support_static_graph_mode() stays false until W2), and W1's tests always Synchronize before destroying or recapturing, so the asymmetry is inert here. The reviewer's own remaining_concern section made the same point about the in-EndCapture destroy. Evidence: reviewer's static read of rocm_backend.hip's Check() calls on hipGraphExecDestroy in both DestroyGraph and EndCapture, contrasted against cuda_backend.cu's bare (unchecked) calls at the same call sites. Remediation: recorded as D6 in .agents/specs/rocm-decode-graph.md §8, matching D1's "Consequence for W2" shape — the decode-graph class must synchronize before destroying or recapturing an exec on HIP. Issue: https://github.com/mudler/vllm.cpp/issues/332 Spec: .agents/specs/rocm-decode-graph.md FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-sonnet-5 [claude-code] --- .agents/specs/rocm-decode-graph.md | 14 ++++++++++++++ 1 file changed, 14 insertions(+) diff --git a/.agents/specs/rocm-decode-graph.md b/.agents/specs/rocm-decode-graph.md index a0e4fe12f..83d4c903e 100644 --- a/.agents/specs/rocm-decode-graph.md +++ b/.agents/specs/rocm-decode-graph.md @@ -330,6 +330,20 @@ the four #41 boards are likelier-supported RDNA3/CDNA parts, but none has run this. Records say gfx1200; the other boards stay `PENDING-community` exactly as the W1 approach-(b) delta does today. +**D6 — `DestroyGraph`/`EndCapture` `Check()` the destroy where CUDA silently +ignores it.** `rocm_backend.hip`'s `DestroyGraph` and the prior-`exec_` destroy +inside `EndCapture` both `Check()` `hipGraphExecDestroy`'s return and throw on +failure; the CUDA leg ignores it. Destroying (or recapturing over) an exec still +in flight would throw on HIP where CUDA would succeed silently. Not a W1 defect +— the model-level path is not engaged (`support_static_graph_mode()` is false +until W2), and W1's tests always `Synchronize` before destroying or recapturing. +Found in the W1 review (`review-rocm-decode-graph-w1.md`, INFO-1). + +*Consequence for W2:* the decode-graph class must synchronize before destroying +or recapturing an exec on HIP — a teardown or column-change recapture that races +an in-flight replay is exactly the case this asymmetry would surface as a thrown +`Check()` instead of a silent no-op. + ## 9. Work breakdown - **W0 — DONE.** [#332](https://github.com/mudler/vllm.cpp/issues/332) filed and From 2fb23e5801a39d5f08f316cb015d9617a9939687 Mon Sep 17 00:00:00 2001 From: Justin Card Date: Tue, 11 Aug 2026 14:16:02 -0400 Subject: [PATCH 4/8] test(rocm): distinguish the handle-path capture from the stored-graph path (INFO-2) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Finding: INFO-2 / mutation 8 in review-rocm-decode-graph-w1.md — EndCaptureGraph returning the stale member exec_ instead of the freshly captured local exec went GREEN, 10/10, in "ROCm backend: graph capture/replay re-executes captured ops" (tests/vt/test_rocm_backend.cpp). Both the stored-graph path (EndCapture, Copy(dst, src)) and the handle-path (EndCaptureGraph) captured the identical operation, so the test could not tell a correct EndCaptureGraph from one that silently handed back the previous graph. Remediation: the handle-path section now allocates its own destination buffer (dst2) and captures Copy(dst2, src) instead of Copy(dst, src), then reads back from dst2. A stale exec_ (still bound to the stored-graph path's Copy(dst, src)) now writes the wrong buffer, leaving dst2 at its pre-capture zero and failing the post-replay CHECKs. IMP-TEST-FIRST / IMP-MUTATE evidence — mutation 8 reapplied to a scratch edit of EndCaptureGraph's return (`return reinterpret_cast(exec_);` in place of the local `exec`), rebuilt, and rerun in isolation: before fix: GREEN, 10/10 (reviewer's original finding) after fix: RED, 7/10, 3 failed — CHECK(back.front() == 0x11) at line 366: values 0 == 17; CHECK(back.front() == 0x22) at line 373: 0 == 34; CHECK(back.back() == 0x22) at line 374: 0 == 34 (dst2 never written by the stale exec_, which still targets dst) Restored byte-for-byte, sha256 confirmed equal before and after the mutation for both src/vt/rocm/rocm_backend.hip and tests/vt/test_rocm_backend.cpp: rocm_backend.hip: 7421f4d78c8f42d16c4d66c954ad5e3aac90bddb3daf1aa493f481736a9f1b5e test_rocm_backend.cpp: 27dc8fc115cccd8da85849e2ac2d799a1e0b91739c7a946bee03b8e3628beec4 Post-fix, unmutated: capture case 10/10 (unchanged assertion count), gate 2 (ctest -R 'rocm|cross_device') 3/3. Reviewer mutations 9 (leaked hipGraph_t) and 10 (skipped prior-exec_ destroy) are explicitly out of scope per the finding — leak-only, not visible to a unit test, not built here. Issue: https://github.com/mudler/vllm.cpp/issues/332 Spec: .agents/specs/rocm-decode-graph.md FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-sonnet-5 [claude-code] --- tests/vt/test_rocm_backend.cpp | 18 +++++++++++++----- 1 file changed, 13 insertions(+), 5 deletions(-) diff --git a/tests/vt/test_rocm_backend.cpp b/tests/vt/test_rocm_backend.cpp index 97d715217..28aa93886 100644 --- a/tests/vt/test_rocm_backend.cpp +++ b/tests/vt/test_rocm_backend.cpp @@ -342,26 +342,33 @@ TEST_CASE("ROCm backend: graph capture/replay re-executes captured ops") { CHECK(back.back() == 0x22); // Handle variant — the path decode graphs actually take, since they keep one - // exec per padded batch size rather than a single stored graph. + // exec per padded batch size rather than a single stored graph. Captures + // into a DIFFERENT destination (dst2) than the stored-graph path above + // (dst), not the same op replayed twice: a mutation that returns the stale + // member exec_ (still the stored-graph path's Copy(dst, src)) instead of the + // freshly captured local exec must be distinguishable from the correct + // behaviour, and only fails here because the two graphs write different + // buffers. + void* dst2 = rocm.Alloc(kBytes); rocm.Copy(q, src, pattern_a.data(), kBytes); - rocm.Memset(q, dst, 0, kBytes); + rocm.Memset(q, dst2, 0, kBytes); rocm.Synchronize(q); rocm.BeginCapture(q); - rocm.Copy(q, dst, src, kBytes); + rocm.Copy(q, dst2, src, kBytes); void* graph = rocm.EndCaptureGraph(q); REQUIRE(graph != nullptr); rocm.ReplayGraph(q, graph); rocm.Synchronize(q); - rocm.Copy(q, back.data(), dst, kBytes); + rocm.Copy(q, back.data(), dst2, kBytes); rocm.Synchronize(q); CHECK(back.front() == 0x11); rocm.Copy(q, src, pattern_b.data(), kBytes); rocm.ReplayGraph(q, graph); rocm.Synchronize(q); - rocm.Copy(q, back.data(), dst, kBytes); + rocm.Copy(q, back.data(), dst2, kBytes); rocm.Synchronize(q); CHECK(back.front() == 0x22); CHECK(back.back() == 0x22); @@ -369,6 +376,7 @@ TEST_CASE("ROCm backend: graph capture/replay re-executes captured ops") { rocm.DestroyGraph(graph); rocm.Free(src); rocm.Free(dst); + rocm.Free(dst2); rocm.DestroyQueue(q); } From 0c93c3dc7f3a87da9c61c1ff1898a4b0a6eb0015 Mon Sep 17 00:00:00 2001 From: Justin Card Date: Tue, 11 Aug 2026 17:36:53 -0400 Subject: [PATCH 5/8] fix(rocm): correct my own overstated -Werror claim in the DestroyGraph comment (LOW-1, take 2) e2160d87 fixed the [[nodiscard]] ATTRIBUTION (hipError_t via __HIP_NODISCARD, not hipGraphExecDestroy specifically) but kept an overstated consequence: "under this project's -Werror it is a hard error." That claim came from a standalone `hipcc -Werror -Wall -Wextra ...` probe, which does NOT reflect this project's actual HIP build flags. Caught by the reviewer's own corrected revision of the review (see previous commit) and independently reconfirmed here two ways: 1. Static: cmake/CompilerWarnings.cmake's vllm_cpp_set_warnings gates -Wall -Wextra -Werror on $ only -- no HIP branch. build-hip/compile_commands.json's actual entry for rocm_backend.hip.o carries no -Wall/-Wextra/-Werror at all: clang++ -DVLLM_CPP_HIP ... -O3 -DNDEBUG -std=c++20 --offload-arch=gfx1200 -fPIC -ffp-contract=off -x hip -c rocm_backend.hip 2. Dynamic: applied a scratch mutation (bare hipGraphExecDestroy call, no Check()) to DestroyGraph and rebuilt through this project's own `vllm` cmake target (not a standalone hipcc invocation): rocm_backend.hip:298:7: warning: ignoring return value of type 'hipError_t' declared with 'nodiscard' attribute [-Wunused-value] [100%] Built target vllm <- exit 0, 0 errors, build succeeds So in this project's actual build, a bare call WARNS and does not fail the build. It is a hard error only under an explicitly-added -Werror, such as the standalone hipcc probe e2160d87 relied on -- which is not what this project's CMake applies to .hip translation units. Remediation: rewrote the comment to state the warning as fact, name the exact gate that would need to change for it to become a build failure (vllm_cpp_set_warnings has no HIP branch), and keep the Check() rationale on its own merits (Free/FreePinned consistency, leak-vs-swallow) rather than on an overstated compile-failure claim. Restored byte-for-byte after the scratch mutation: sha256 7421f4d78c8f42d16c4d66c954ad5e3aac90bddb3daf1aa493f481736a9f1b5e before and after (matches the hash recorded in 93ad9ccc, confirming no drift from the mutation probe). Issue: https://github.com/mudler/vllm.cpp/issues/332 Spec: .agents/specs/rocm-decode-graph.md FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-sonnet-5 [claude-code] --- src/vt/rocm/rocm_backend.hip | 15 ++++++++++----- 1 file changed, 10 insertions(+), 5 deletions(-) diff --git a/src/vt/rocm/rocm_backend.hip b/src/vt/rocm/rocm_backend.hip index 614317f33..0ee462aff 100644 --- a/src/vt/rocm/rocm_backend.hip +++ b/src/vt/rocm/rocm_backend.hip @@ -288,11 +288,16 @@ class RocmBackend final : public Backend { // [[nodiscard]] at C++17+ ($ROCM_PATH/include/hip/hip_runtime_api.h:293-305, // __HIP_NODISCARD on the `typedef enum ... hipError_t` line; the function // declaration at line ~8321 carries no attribute of its own). CUDA's - // cudaError_t has no such attribute, so the mirrored bare call compiles - // there and not here: under this project's -Werror it is a hard error - // (-Werror,-Wunused-value). Checked rather than discarded, matching - // Free/FreePinned above: a failing destroy is a leak this backend would - // rather report than swallow. + // cudaError_t has no such attribute, so the mirrored bare call diagnoses + // here (-Wunused-value) and not there. This project's -Werror does NOT + // reach it, though: vllm_cpp_set_warnings (cmake/CompilerWarnings.cmake) + // gates -Werror on COMPILE_LANGUAGE CXX/OBJCXX/CUDA only, no HIP branch, so + // a bare call here would warn, not fail the build (verified against the + // actual compile_commands.json entry for this file, and by building a + // scratch bare-call mutation through this project's own cmake target). + // Checked anyway, matching Free/FreePinned above: a failing destroy is a + // leak this backend would rather report than swallow, independent of what + // the build flags happen to enforce. void DestroyGraph(void* graph) override { if (graph != nullptr) { Check(hipGraphExecDestroy(reinterpret_cast(graph)), From 46d2f4c6eeab2a2f7f301a0d135bc54ab3cf3fa7 Mon Sep 17 00:00:00 2001 From: Justin Card Date: Wed, 12 Aug 2026 08:44:48 -0400 Subject: [PATCH 6/8] build(nix): add rocwmma to the ROCm dev shell so gfx12 builds again (#444) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Since 9302732f, rocm_paged_attn.hip includes whenever the target is gfx1200/gfx1201. The guard is on ARCH, not on availability, so targeting a gfx12 board fires the include whether or not rocWMMA exists. clr — the store path this shell uses as ROCM_PATH — does not ship it, so `nix develop .#rocm-shell` could not build main at all on a gfx12 board: src/vt/rocm/rocm_paged_attn.hip:8:10: fatal error: 'rocwmma/rocwmma.hpp' file not found Reproduced on pristine upstream/main at 0f2b12ed in a clean worktree with no other commits applied, so it is not a local-branch artifact. rocwmma is header-only, so it rides the existing overlay: rocmOverlayInputs symlinks each input's include/ into $ROCM_OVERLAY/include, which is already on CPATH. Adding it to the shell's packages list too keeps the two lists reading the same. Note the overlay is cached behind a .complete sentinel, so an existing shell needs $ROCM_OVERLAY removed once to pick this up. This is the NARROW half of #444 and does not close it. It fixes this shell only; every other rocWMMA-free environment targeting gfx12 still fails the same way. The issue asks for CMake-level detection that fails configure with a message naming the package, which is the real fix and stays open. Deliberately NOT done here: gating the include on __has_include. It works (verified under hipcc, which correctly reports the header absent), but VT_ROCWMMA_OK also selects the kernel BODY — PagedAttnPrefillWmmaWave at rocm_paged_attn.hip:900 casts every parameter to (void) and returns when the macro is undefined. Deciding that macro by availability rather than arch would compile a paged-attention kernel that silently writes nothing on a gfx12 board with no rocWMMA. A build failure is the correct outcome there; this commit supplies the missing dependency instead of hiding the symptom. Verified: with this change `nix develop .#rocm-shell` builds the HIP target on gfx1200 with 0 warnings, and `ctest -R 'rocm|cross_device'` is 4/4. Issue: https://github.com/mudler/vllm.cpp/issues/444 FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [claude-code] --- flake.nix | 7 ++++++- 1 file changed, 6 insertions(+), 1 deletion(-) diff --git a/flake.nix b/flake.nix index c023a18e6..5d1967ac4 100644 --- a/flake.nix +++ b/flake.nix @@ -67,8 +67,12 @@ # entry (idempotent — skipped if already populated). This is the # one Nix-specific step; a standard /opt/rocm install needs none of # it. + # rocwmma is header-only and rides the same overlay: rocm_paged_attn.hip + # includes whenever the target is gfx1200/gfx1201 + # (arch-gated, not availability-gated), and clr alone does not ship it, + # so a gfx12 build fails at that include without this. See issue #444. rocmOverlayInputs = - [ rocm.hipblas rocm.hipblaslt rocm.hipblas-common ]; + [ rocm.hipblas rocm.hipblaslt rocm.hipblas-common rocm.rocwmma ]; in { default = pkgs.mkShell { packages = commonPackages ++ [ pkgs.gcc ]; @@ -122,6 +126,7 @@ rocm.hipblas rocm.hipblaslt rocm.hipblas-common + rocm.rocwmma rocm.rocminfo ]; From b41e38b2331ce4e6cc2ea9f28cf6592f4253f4c5 Mon Sep 17 00:00:00 2001 From: Justin Card Date: Thu, 13 Aug 2026 11:10:41 -0400 Subject: [PATCH 7/8] revert(rocm): drop the flake.nix rocwmma fix from this PR, split to #444 localai-org-maint-bot review on PR #473 recommended this PR stay scoped to the graph-capture seam; the rocwmma dev-shell fix (46d2f4c6) is unrelated scope for #444 and belongs in its own PR. Reverting it here; it will land via a fresh branch targeting #444 instead. This reverts commit 46d2f4c6eeab2a2f7f301a0d135bc54ab3cf3fa7. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-sonnet-5 [claude-code] --- flake.nix | 7 +------ 1 file changed, 1 insertion(+), 6 deletions(-) diff --git a/flake.nix b/flake.nix index 5d1967ac4..c023a18e6 100644 --- a/flake.nix +++ b/flake.nix @@ -67,12 +67,8 @@ # entry (idempotent — skipped if already populated). This is the # one Nix-specific step; a standard /opt/rocm install needs none of # it. - # rocwmma is header-only and rides the same overlay: rocm_paged_attn.hip - # includes whenever the target is gfx1200/gfx1201 - # (arch-gated, not availability-gated), and clr alone does not ship it, - # so a gfx12 build fails at that include without this. See issue #444. rocmOverlayInputs = - [ rocm.hipblas rocm.hipblaslt rocm.hipblas-common rocm.rocwmma ]; + [ rocm.hipblas rocm.hipblaslt rocm.hipblas-common ]; in { default = pkgs.mkShell { packages = commonPackages ++ [ pkgs.gcc ]; @@ -126,7 +122,6 @@ rocm.hipblas rocm.hipblaslt rocm.hipblas-common - rocm.rocwmma rocm.rocminfo ]; From 6e500d2c533e2f76896133d1f0bd96c194e188d9 Mon Sep 17 00:00:00 2001 From: Justin Card Date: Thu, 13 Aug 2026 11:13:41 -0400 Subject: [PATCH 8/8] =?UTF-8?q?spec(BACKEND-ROCM):=20record=20D7=20--=20th?= =?UTF-8?q?e=20=C2=A71/gate-5=20throughput=20rationale=20is=20refuted=20(#?= =?UTF-8?q?332)?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit localai-bot's review on PR #473 flagged that the W3 refutation (capture moved throughput 0-2%, not the predicted ~1.36x convergence) lived only in the PR description while this spec still opened with the live 2.99x/1.90x/1.46x table and an unretracted 'the premise survived' argument. A future agent opening this file to pick up W2/W3 would re-derive a result already closed. Adds D7 (§8) with the A/B table, the 126-replays-over-128-steps proof capture engaged, the user/sys/wall split showing capture removes no host CPU work, and the ordered next hypotheses (same-tool rocprof trace; where the host CPU time goes). Marks §1's fit paragraph and §7 gate 5's prediction table SUPERSEDED rather than rewriting them -- a falsified pre-registered prediction that gets quietly reworded is worthless. The 0.6B figure uses the review-corrected 0-2% (not the original run's overstated +3.2%, later traced to one low outlier in the OFF arm and confirmed independently) -- see the row's W2/W3 review findings. Also two records fixes from the same review: - src/vllm/platforms/rocm.cpp:67-69 still said hipGraph capture 'is not implemented' after this PR implemented it (the twin comment in rocm_backend.hip:22-28 was updated, this one was not). - Issue #444 (main does not build for gfx1200/gfx1201) is carried by this row's flake.nix commit but was absent from the roadmap issue table. Issue: https://github.com/mudler/vllm.cpp/issues/332 Spec: .agents/specs/rocm-decode-graph.md FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-sonnet-5 [claude-code] --- .agents/roadmap_v1.md | 1 + .agents/specs/rocm-decode-graph.md | 82 +++++++++++++++++++++++++++--- src/vllm/platforms/rocm.cpp | 10 ++-- 3 files changed, 83 insertions(+), 10 deletions(-) diff --git a/.agents/roadmap_v1.md b/.agents/roadmap_v1.md index d13b8e948..81d30f346 100644 --- a/.agents/roadmap_v1.md +++ b/.agents/roadmap_v1.md @@ -43,6 +43,7 @@ issue is not yet placed. Keyed record: update in place, never append. | [#201](https://github.com/mudler/vllm.cpp/issues/201) | `BACKEND-ROCM` | `hipblasGemmEx` overload mismatch in `rocm_matmul_hipblaslt.hip` | bug | | [#269](https://github.com/mudler/vllm.cpp/issues/269) | `BACKEND-ROCM` | ROCm gfx1200: Gemma-3 is 48/48 exact vs two vLLM-ROCm oracles; Qwen3-0.6B exposes a deterministic cross-version near-tie, not a backend defect | verification | | [#332](https://github.com/mudler/vllm.cpp/issues/332) | `BACKEND-ROCM` | ROCm: no decode-graph capture — the hipGraph seam is unimplemented, costing ~3x decode throughput vs vLLM on gfx1200 | perf | +| [#444](https://github.com/mudler/vllm.cpp/issues/444) | `BACKEND-ROCM` | ROCm: main does not build for gfx1200/gfx1201 — `rocm_paged_attn.hip` includes rocwmma unconditionally on arch, not availability | bug | | [#125](https://github.com/mudler/vllm.cpp/issues/125) | `BACKEND-VULKAN` | Vulkan on AMD Strix Halo (gfx1151) does not load | bug | | [#203](https://github.com/mudler/vllm.cpp/issues/203) | `BACKEND-VULKAN` | Vulkan on unified memory holds TWO copies of the weights: 27B peaks at 100.8 GiB RSS and OOM-reboots a Spark | bug | | [#310](https://github.com/mudler/vllm.cpp/issues/310) | `BACKEND-VULKAN` | docs/FEATURES.md understates Vulkan: says decode 4.24 vs 4.35 where the binding figure is 4.36 vs 4.35 | bug | diff --git a/.agents/specs/rocm-decode-graph.md b/.agents/specs/rocm-decode-graph.md index 83d4c903e..743d37773 100644 --- a/.agents/specs/rocm-decode-graph.md +++ b/.agents/specs/rocm-decode-graph.md @@ -59,13 +59,16 @@ like a confound but are not — launches and compute both scale with layers, so the layer count cancels and `L/C` depends only on per-layer width (`1 / (hidden x intermediate)`). All three sit on one curve. -**It also bounds the win.** Fitting `ratio = alpha + beta / (hidden x inter)` -across the three points gives **alpha ~= 1.36x** (size-independent: kernel -quality, inductor fusion, the Triton attention path) and **beta ~= 1.65** in -units where 0.6B's `hidden x inter` = 1 — an overhead contribution of ~1.65x at -0.6B, ~0.41x at 1.7B, ~0.21x at 4B. So roughly half the 0.6B gap is fixed -overhead, and the expected outcome is all three sizes converging on **~1.36x**: -a real win, and **not parity**. A ~1.4x residual would remain. +**It also bounded the win — WRONGLY; this paragraph is the falsified +prediction, preserved verbatim rather than quietly reworded.** Fitting +`ratio = alpha + beta / (hidden x inter)` across the three points gave +**alpha ~= 1.36x** (size-independent: kernel quality, inductor fusion, the +Triton attention path) and **beta ~= 1.65** in units where 0.6B's +`hidden x inter` = 1 — an overhead contribution of ~1.65x at 0.6B, ~0.41x at +1.7B, ~0.21x at 4B. The reasoning WAS that roughly half the 0.6B gap is fixed +overhead, so the expected outcome WAS all three sizes converging on **~1.36x**. +**That did not happen.** D7 (§8) has the measurement: capture moved throughput +0-2%, indistinguishable from zero, not the predicted convergence. Treat the fit as provisional. An earlier two-point version gave `alpha ~= 1.54x`; Qwen3-4B then measured 1.46x, below that asymptote, which a curve cannot do, so @@ -258,6 +261,13 @@ mutations that must turn it red, in a scratch copy, restored byte-for-byte: fixed-overhead term is smaller than the fit implies, and redirect to D4. **Do not describe any outcome as parity**; ~1.36x is the predicted floor for this change alone. + + **SUPERSEDED.** This is the pre-registered prediction, kept verbatim per + AGENTS.md ("never trade correctness for throughput" applies equally to + quietly rewriting a call before the result). W2/W3 ran it: capture engaged + correctly on all three sizes and moved throughput 0-2%, not toward ~1.36x. + See D7 (§8) for the measurement and why. Any future W2/W3 claim on this row + starts from D7, not from this table. 6. **`GetReferenceTierHits()` == 0** in any perf measurement — structurally impossible on a discrete board, assert anyway. 7. **Records green:** `agent-preflight.sh --staged`, `check-agent-record.py`, @@ -344,6 +354,64 @@ or recapturing an exec on HIP — a teardown or column-change recapture that rac an in-flight replay is exactly the case this asymmetry would surface as a thrown `Check()` instead of a silent no-op. +**D7 — the §1/gate-5 rationale is REFUTED. W2 and W3 ran (on a follow-on +branch, not shipped in this PR) and the pre-registered convergence prediction +did not hold.** Recorded here so the next agent to open this spec does not +re-derive a closed negative result from a still-live-looking prediction. + +Same-binary A/B on gfx1200, 128in/128out batch 8, capture ON vs +`VLLM_CPP_CUDAGRAPH=0`, gate 4's `VT_DECODE_GRAPH_STATS=1` confirming capture +engaged (**126 replays over 128 decode steps**, one capture at padded size +`S=8`, not eager fallback): + +| Model | capture ON | capture OFF | delta | +|---|---|---|---| +| Qwen3-0.6B | 190-194 tok/s | 188-192 tok/s | **0-2%**, indistinguishable from zero | +| Qwen3-1.7B | 148.72 tok/s | 147.80 tok/s | +0.6% | +| Qwen3-4B | 100.22 tok/s | 101.27 tok/s | -1.0% | + +Qwen3-0.6B needed a second pass: the first same-binary run measured +3.2% +(193.67 vs 187.62), but an independent reviewer's own 3-rep A/B measured ++0.7%. The discrepancy traced to one low outlier in the original OFF arm +(reps 191.88 / 179.29 / 191.70); dropping it moves the original delta to ++1.0%, and pooling both reviewers' reps gives +2.0%. **0-2% is the honest +figure; +3.2% must not be quoted as a measured win.** + +Not batch-size dilution: single-stream is capture's best case and shows no +gain either (TPOT 13.99 ms OFF vs 13.82-14.98 ms ON). And the host/device +split says why capture isn't paying off here — it removes no host CPU work: + +```console +capture ON wall 11.31s user 13.81s sys 14.09s 247% CPU +capture OFF wall 11.36s user 13.78s sys 13.54s 240% CPU +``` + +Collapsing hundreds of per-step launches into one call should cut `sys` time +visibly; it is marginally *higher*. §1's scaling-curve argument — that roughly +half the 0.6B gap is fixed launch overhead recoverable by capture — is +directly refuted by this intervention. `cuda_backend.cu`'s +"88%-of-wall host-API overhead" figure is a CUDA measurement and does not +transfer to this board. + +Against §1's oracle figures the ratios are 2.85x / 1.93x / 1.43x versus +2.99x / 1.90x / 1.46x before capture — **unmoved**. §10's stop condition +(Qwen3-0.6B below ~2.2x) was not met; W3 stopped there per the spec's own +rule, rather than iterating blind. + +**Not a ceiling claim.** The gap is unexplained, not irreducible. Next +hypotheses, in order: (1) where the ~14 ms decode step actually goes on the +device — same-tool `rocprof` both sides, the trace D4 already flagged as +missing; (2) what burns ~2.5 cores of host CPU to produce ~50 tok/s with `sys` +exceeding wall. + +**What this leaves standing, unaffected by the refutation:** the seam itself +— mirrored call-for-call, mutation-tested, behaviour-neutral across four +models (Qwen3-0.6B/1.7B/4B dense, Qwen3.5-0.8B GDN hybrid) and a 6.7x +parameter range, capture ON vs OFF byte-identical. What died is the reason it +was built, not the code. Whether to carry an unused capability given the +refuted rationale is a product call, not a technical one — see the row's PR +discussion for that decision; it does not belong in this spec. + ## 9. Work breakdown - **W0 — DONE.** [#332](https://github.com/mudler/vllm.cpp/issues/332) filed and diff --git a/src/vllm/platforms/rocm.cpp b/src/vllm/platforms/rocm.cpp index 43ce027df..295f32bbd 100644 --- a/src/vllm/platforms/rocm.cpp +++ b/src/vllm/platforms/rocm.cpp @@ -64,9 +64,13 @@ class RocmPlatform final : public Platform { // supports_fp8() stays false: gfx942/gfx950 have hardware fp8 and rocm.py lists // "fp8" in supported_quantization (rocm.py:457-467), but we have no ROCm fp8 // kernel, and this predicate gates a fused path that would then not exist. - // support_static_graph_mode() stays false: hipGraph is the mapping and is not - // implemented. Capture that bakes a wrong address is a silent correctness - // bug, so this flips only alongside a real capture implementation. + // support_static_graph_mode() stays false: the vt::Backend hipGraph capture + // seam is implemented as of BACKEND-ROCM W1 (rocm_backend.hip; see + // .agents/specs/rocm-decode-graph.md) and the address-baking concern that + // used to justify leaving this false is now an assertion, not a worry — + // the mutate-src-then-replay test step fails if replay ever returns a + // snapshot. This flag still stays false because flipping it to engage a + // real model's decode-graph path is W2, not W1. // needs_weight_staging() stays false: this is the memory-model POLICY that // selects the device-resident forward over the host-resident reference path. // HIP's programming model does stage (hipMalloc hands back a distinct