gate(ROCM): M4 for Qwen3.5-0.8B GDN -- oracle-backed green near-tie gate by VikashLoomba · Pull Request #559 · mudler/vllm.cpp

VikashLoomba · 2026-08-13T00:24:05Z

M4 for Qwen3.5-0.8B (GDN) — the oracle-backed gate, landed green-shaped (issue #41)

Stacked: #549 (0.6B lane + oracle) ← the dispatch fix (#570) ← this gate. The fix lands below the gate, so every commit in the stack is individually green on the gate host and the committed goldens are ones the tree can actually reproduce.

Gate

test_qwen35_paged_engine.cpp — the GDN sibling of the Qwen3-dense SACRED gate: standard 16-prompt battery through the full paged engine, hard anchor + ≤0.5-nat near-tie band, backend proof over the 15-op GDN set. Goldens captured from the pinned vLLM-ROCm oracle on this box (K=10 per-prompt, deterministic in every cell) with manifest.json recording the oracle identity; the checkpoint resolves via parity::HfSnapshot at the pinned HF revision 2fc06364 (check-snapshot-pins clean).

Verdict: 16/16 PASS (15/16 strict token-exact vs the oracle, 1/16 near-tie band, max gap 0.125 nats, 0 forward-divergent).

Review-sweep rework (all three findings accepted)

CI-red manifest fixed: qwen35_0_8b_greedy is in PendingRunnerOps() (the qwen36_gguf_greedy precedent) — the manifest itself stays; its oracle identity is the point.
Fail-safe by device: the gate exits 77 (CTest Skipped) on any non-ROCm device and when the pinned checkpoint is absent — never a false green, never a foreign engine compared against ROCm-derived goldens.
Green-shaped landing: the committed goldens are the FIXED engine's oracle-re-derived pair. The pre-fix capture (13/16 forward-divergent; first-token divergences on 6 prompts at 0.375–1.062 nats, 5 of 6 over band) is kept as evidence in the spec + parity ledger, not as goldens no code can pass. The anchor-drift REQUIRE now reads "REGRESSION SUSPECTED — bisect first"; golden re-derivation is the last step of a justified re-capture, not the response to a failure.

Prose corrections from review

First-divergence maximum is 1.062 nats (not 1.188), 5 of the 6 token-0 divergences are over the band (not all), and the headline gap is at the first over-band position rather than a post-divergence maximum.

In-flow gate repair

Issue #1222 found three missing-artifact paths that returned normally and let doctest report false success. The gate now routes absent greedy, anchor, and gap artifacts through exit 77. A no-GPU subprocess regression drives the real executable and asserts 77/77/77; its complete-artifact control requires hard sentinel exit 86 so probe mode can never bypass snapshot/model/device checks as zero-assertion success. Bootstrap writes its token IDs and then exits skipped until correctness artifacts exist. Malformed present arrays remain hard failures.

Fixes #1222.

Current-main integration cutoff (2026-08-18)

Rebased from 5b31bc30349d527ef0b4c1c8fcd7c86749b4cfa3 to 4673a35dd47087cae3f6f9808e4fcf64667bcbcf on recorded current-main cutoff 65d6cdaed3e20e9bc70b4f9374fccafefefa7bd0. Exact six-commit replay of #549 → #570 → the green 0.8B gate → the three #1222 spec/repair commits. Fresh cutoff-head builds, focused gates, full repository preflight, and independent mutation review all pass.

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: pi:kimi-k3 [pi]
Assisted-by: codex:gpt-5.6-sol [codex]
Assisted-by: pi:gpt-5.6-sol [pi]

localai-bot · 2026-08-13T17:14:50Z

Reviewed as part of a sweep over the open external PRs. The measurement behind this is real — I recomputed the verdict independently from the committed goldens and got exactly your numbers: 13/16 fail, 3 pass, max gap 14125 mnats at (prompt 8, tok 8), 6 prompts differing at token 0. Golden provenance also checks out: greedy_ids.npy and our_ids.npy differ in 196 of 256 positions so they are not the same source, and the prompt token ids are distinctly Qwen3.5's rather than copied from the 0.6B lane. The manifest recording the oracle identity is a genuine improvement over the existing goldens, which record none.

Three things before this can land.

1. It would turn main's CPU CI red for every subsequent PR.

tests/parity/goldens/qwen35_greedy_0_8b/manifest.json declares "op": "qwen35_0_8b_greedy", and no runner is registered for it, so test_op_parity.cpp:1933 hits FATAL ERROR: no runner for op 'qwen35_0_8b_greedy' — add one before committing goldens. That takes down build-test-cpu and both sanitize-cpu jobs — all three are the same single ctest failure, not sanitizer diagnostics (no ASan report, no TSan race anywhere in the logs), and all three are green on main.

The guard is behaving correctly: the existing qwen3_greedy_0_6b/ goldens only avoid it by having no manifest at all. So please fix it by adding the runner arm next to the existing qwen3_5_mtp_head / gdn_packed_decode_bf16 precedents, or a PendingRunnerOps() entry — not by deleting the manifest, which is the best thing in this PR.

2. The gate does not fail safe by device.

const bool other_partial = metal || tenstorrent; — kCPU and kCUDA set neither flag, so nothing returns and the gate proceeds to teacher-force a ROCm-captured anchor against whatever device is running, with the entire backend proof (OpRegistered, selections>0, declines==0) sitting inside if (rocm) and therefore skipped. On a CPU or CUDA box with that checkpoint cached you get either a hard REQUIRE "anchor drift" failure attributed to a CUDA regression that does not exist, or a green gate with no proof of which device ran. The only thing preventing this today is that the pinned revision is not cached on anyone else's box, which is luck rather than a fail-safe. Two lines fixes it: if (!rocm) { MESSAGE(...); std::exit(77); }.

Related, and cheap: the checkpoint-absent path at :118-123 returns, so doctest prints assertions: 0 | 0 passed | 0 failed and Status: SUCCESS! — indistinguishable in a log from a gate that loaded a model and matched an oracle. tests/CMakeLists.txt:34 already wires SKIP_RETURN_CODE 77 for every test and test_voxtral_e2e.cpp:55 has the canonical SkipGate() helper; this is issue #463's exact pattern.

3. The RED anchor does not have teeth against the defect it names.

The pass/fail decision is if (mn > kNearTieMnats) prompt_ok = false; where mn is read from the committed neartie_gap_mnats.npy. The engine's output enters only via the anchor REQUIRE(got == our_ids). So the committed RED is a property of a data file rather than of the code, and no forward fix can turn it green — the anchor REQUIRE fires first. #570 confirms this by replacing the goldens rather than passing the gate, which means the gate never had the chance to fail for the original defect.

The shape that works is to land the gate with GREEN goldens alongside #570, keeping the RED capture as recorded evidence in the spec and the parity ledger, where it already lives in full.

One more worth fixing while you are here: the anchor-drift message ends "— re-run qwen3-neartie-gap.py to refresh the gap golden". After #570, that REQUIRE is the only thing catching a reintroduction of the dtype-dispatch bug, and its failure text tells the next engineer to regenerate the golden. It should say a drift is a regression until proven otherwise.

Numbers to correct in the prose: the six token-0 divergences are 0.375/0.750/0.750/0.938/1.000/1.062 nats, so the max is 1.062 not 1.188, and only 5 of the 6 are over the 0.5-nat band. And worst_gap scans positions after the first over-band token, where the oracle is teacher-forced on an already-diverged prefix — the defensible headline is the gap at the first over-band position.

On CI: agent-record and pr-size are the unrebased-fork artifact (now fixed on our side in #619) and Windows is the known-broken arm — but build-test-cpu and both sanitizers are genuinely yours, per finding 1.

No AMD hardware here, so nothing was executed on ROCm; findings 1–3 are static, and the verdict arithmetic above came from your committed files.

VikashLoomba · 2026-08-14T09:08:46Z

All three findings accepted and reworked (commit a77a3b3, rebased onto current main):

CI-red fixed: qwen35_0_8b_greedy is in PendingRunnerOps() next to the qwen36_gguf_greedy precedent — the manifest stays (its oracle identity is the improvement you credited).
Fail-safe by device: the gate exits 77 (CTest Skipped) on any non-ROCm device AND on absent checkpoint (the test_voxtral_e2e reports SUCCESS with zero assertions when its weights env var is unset (and ~40 more gates share the shape) #463 pattern) — no more foreign-engine-vs-ROCm-goldens comparison, no more 0-assertion "SUCCESS".
Green-shaped landing: the stack is reordered — the kernel fix (fix(ROCM): AttnQkNormRopeGate output-dtype dispatch -- the 0.8B GDN forward divergence root-caused and fixed (RED anchor turns GREEN) #570) now lands BELOW this gate, so the committed goldens are the fixed engine's oracle-re-derived pair and every commit in the stack is individually green on the gate host. The RED capture (13/16, max first-divergence 1.062 nats — corrected per your arithmetic, 5 of 6 over band) lives as evidence in the spec + parity ledger. The anchor-drift REQUIRE now reads "REGRESSION SUSPECTED — bisect the engine change first"; golden re-derivation is the last step of a justified re-capture, never the response to a failure.

Gate on gfx1100: 16/16 PASS (15/16 strict, 1/16 band, max gap 0.125 nats); 0.6B gate re-run green.

… validates an outside contributor's trailers (#773) (#782) Closes #773. `check-pr-size.py` and `check-commit-trailers.py` both required the base revision to be an ANCESTOR of head. CI passes `github.event.pull_request.base.sha`, the TIP of the base branch, which stops being an ancestor the moment main advances after the branch was cut -- continuously, on this repo. Measured on three open PRs before changing anything. Base is not an ancestor in any of them, and a merge base exists in all three: #506 ancestor=NO merge-base=e1087a881 #523 ancestor=NO merge-base=fdd452637 #559 ancestor=NO merge-base=fafa16f0f The consequence was not a noisy red check. Both checkers aborted BEFORE examining anything, so CI had never validated commit trailers on an external contribution: the gate enforcing FOLLOWING_AGENTS_PROTOCOL and Assisted-by: exited before reading a single commit. Across the external PRs reviewed this week, hand-checking by a reviewer was the only verification those trailers received. pr-size aborted identically, so path classification and the checker-evidence contract went unenforced on forks too. THE FIX. Diff from the merge base, which is what a pull request IS: `git diff A...B` is defined as `git diff $(git merge-base A B) B` and is what GitHub shows. Two-dot diffing against a moved main is not merely stricter, it is WRONG -- main's own commits render as reversions inside the contributor's diff, so paths they never touched get classified and charged to them. The new pr-size test asserts both halves: the PR's file present, main's absent. `executable_evidence` gets the same treatment, since the BASE version of a checker for the red-before half is the one at the merge base. WHAT DELIBERATELY DOES NOT MOVE. The old rule conflated two situations: ordinary divergence (merge base exists) now examines merge_base..head; unrelated histories (no merge base) STILL RAISES. Absence of information must never look like absence of work -- the script's own require_origin_main() docstring already states that principle for the other input. test_missing_and_nonancestor_objects_fail_closed uses an ORPHAN branch, so it still raises; only its regex changed, because the message now names what is actually wrong. Its assertRaises(ValueError) is untouched. The trailers case was SPLIT, not deleted: its divergent-branch half built two branches off a common root -- which share a merge base and are the ordinary shape of every PR -- so that half now asserts it validates, with the genuinely-unrelated case asserted separately. Nothing that used to fail closed stopped failing closed. Range changed, contract unchanged: test_a_bad_trailer_in_the_merge_base_range_is_still_reported puts a trailerless commit inside the new range and requires it still be reported. RED before on the unmodified checkers, GREEN after (74 passed, 148 subtests). Stop conditions checked individually rather than inferred. Full tests/scripts: 9 failed / 1359 passed, all nine pre-existing and reproduced on main. CI: agent-record and pr-size both SUCCESS on this PR -- the two checkers it repairs passing on a live PR. Remaining red is baseline only: windows-msvc-* are the PR-only arm (#584), and sanitize-cpu is red on main itself for #775 (test_nemotron_h_scaffold, nemotron_h_registry.cpp:112 downcasting a doctest StubModel to NemotronHLoadedModel). This PR touches no C++.

…d of aborting the pass (#776) (#853) Closes #776. `test_op_parity` THREW instead of failing an assertion: tests/parity/test_op_parity.cpp:1989: ERROR: test case THREW exception: [json.exception.type_error.302] type must be string, but is null The malformed input was tests/parity/goldens/minimax_music3_oracle/manifest.json -- an oracle capture with no "op" key at all, the only non-op manifest among the 66 committed. Attributed decisively rather than by inference: rebuilding the pre-fix source and moving that one directory aside turned the same binary green, 70 assertions -> 123. The throw was hiding a third of the gate. #776 IS A DUPLICATE, AND THAT IS NOT THE END OF IT. #755 is the same defect, closed by 043e568 (#761) eight minutes after #776 was filed, and the manifest is already repaired on main. Reporting only that would have been accurate and useless, because #761 closed the walker's INPUT SET and not its EXCEPTION SURFACE. Reproduced on the FIXED tree by nulling one tensor dtype in rmsnorm_f32_8x128: :2177: ERROR: test case THREW exception: [...type_error.302...] [doctest] assertions: 37 | 37 passed 37 where 153 should run -- 45 committed goldens never reached, on a tree where this was supposedly closed. The gate was one bad field away from silently not gating. WHY THAT MATTERS: test_op_parity is the guard that refuses a golden whose manifest names an op with no registered runner. It is what caught PR #559's missing runner arm. A throw at manifest-read time aborts the case, so every golden after the malformed one goes unchecked -- not merely a red job, a gate that had stopped gating. THE FIX, both halves. The two remaining throw sites (json::parse on the manifest, and any runner field read) funnel through GuardGoldenStage, which converts a std::exception into a FAIL_CHECK NAMING goldens/<case>/manifest.json, quotes the original exception, and continues to the next golden. The dispatch chain moved into RunGoldenCase so its `continue`s became `return false`, which is what permits a per-case catch. std::exception is the deliberate width: doctest::detail::TestFailureException is NOT derived from it (third_party/doctest/doctest.h:2563). That is load-bearing -- the "no runner for op" refusal is a FAIL, and a wider catch would have cleared the red while REMOVING the gate. MUTATION EVIDENCE, tree restored byte-for-byte each time (823 files, md5sum -c): dtype: null in a real golden -> named by path; 45 op cases STILL RAN, 152 assertions still checked (was: THREW, 37) invalid JSON manifest -> named by path; 46 cases ran, no THREW {"op": null} / missing "op" -> refused by name, pass continues {"op": "no_such_runner_op_559"} -> FATAL ERROR: no runner for op ... -- the #559 guard STILL BITES The last row is the one that mattered; without it the others would only prove the red was silenced. GREEN: focused 13 cases / 153 assertions, with 46 op cases + 1 non-op skip -- counts IDENTICAL to before, so the extraction changed nothing about what is gated. sanitize-cpu (address,undefined) rebuilt with the job's own flags and run under its env: SUCCESS. pytest tests/scripts: 1400 passed, 9 failed, all three distinct failures reproduced on a clean origin/main worktree. Survey: ~167 unguarded m.at(...) reads live inside runners in this file and the guard covers them generically. Other json::parse sites in tests/ each read one named fixture owned by their own case, so a throw there fails one already-named test rather than a walker over an open input set. No follow-up owed. CI: merged with checks queued -- the runner pool has been saturated for hours and no job started on this head. The sanitize lane, which is the one this change exists for, was reproduced locally with that job's own build flags and environment. Note main moved mid-gate: #840 relocated the issue table out of roadmap_v1.md into the append-only .agents/issue-index.md. roadmap_v1.md was taken from origin/main wholesale and the #776 row reapplied as an append to the new index.

VikashLoomba · 2026-08-18T06:43:02Z

Current-main rebase and fresh review

Rebased from acea2bef to 14d2850b on current-main pin f22c6cc8 and force-updated with an explicit lease. The PR title/body now describe the green-shaped stack rather than the historical RED anchor.

Qwen3.5-0.8B gate ran on ROCm, not skipped: 16/16 prompts, 137/137 assertions; 15 strict, 1 ratified near-tie, maximum 0.125 nats, zero forward-divergent.
HF revision remains pinned to 2fc06364715b967f1860aea9cf38778875588b17; all 21 golden/manifest/prompt artifacts remain committed.
Removing the complete CMake gate block regenerated cleanly but the explicit configured-CTest assertion went RED at exact count 0. Restoration produced exact count 1 and the real model gate passed.
M4 topology is exact: gate(ROCM): M4 -- pinned vLLM-ROCm oracle on gfx1100 + ROCm device-golden lane in the Qwen3-0.6B near-tie gate #549 → fix(ROCM): AttnQkNormRopeGate output-dtype dispatch -- the 0.8B GDN forward divergence root-caused and fixed (RED anchor turns GREEN) #570 → this PR. Derived NOW/coordination records are untouched.
Repository preflight: All gates green.
Fresh independent integration review: PASS; final tree clean.

…lden lane in the Qwen3-0.6B near-tie gate Issue mudler#41 M4 milestone. The oracle is upstream vLLM AT THE PIN (555967922, runtime 0.23.1rc1.dev1511+g555967922 -- verified == the parity-pin block) built from source as a ROCm wheel for gfx1100 in rocm/vllm-dev:base and committed as the image vllm-rocm-oracle:555967922-gfx1100 (recipe: .agents/specs/rocm-m4-oracle.md). Gate result on this box (4x RX 7900 XTX, ROCm 7.14, enforce_eager, batch=1, flock ~/gpu.lock): oracle K=10 per-prompt greedy deterministic in every cell; our ROCm engine 16/16 prompts PASS -- 11/16 strict token-exact vs base greedy, 5/16 near-tie band, max teacher-forced gap 0.125 nats, 0 forward-divergent, 125/125 assertions, backend proof all 8 Qwen3-dense ops selections>0 and declines==0 on device 5. The p0 France/Italy flip is a literal 0.0000-nat tie in the oracle's own logits (the Metal M3b signature, reproduced on ROCm). Change shape: the parity gate's Metal/Tenstorrent device-golden lane gains the kROCM arm (identical anchor+band logic, ROCm's own oracle-backed golden pair); dgx base goldens byte-untouched; the *_rocm goldens are additive. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: pi:kimi-k3 [pi]

…he 0.8B GDN forward divergence, root-caused and fixed Issue mudler#41. The ROCm dispatcher for the fused full-attention preamble keyed its template on the SOURCE dtype (qgate.dtype) instead of the OUTPUT dtype like the CUDA lane (LaunchAttnPreambleOut switches on q_out.dtype). The 0.8B bf16 model runs bf16 QKV-projection output + f32 q/k/gate outs (the f32-attention path; no FA-2 on ROCm), which the src-keyed dispatch mis-launched as all-bf16 -- writing bf16 bits through the f32 out pointers. Evidence chain (gfx1100, flock): CPU==pinned-oracle token-exact / ROCm diverges at tok0; VT_DUMP_ACT per-layer residual dumps: layers 0-2 clean (<=9.6e-3), layer 3 (first full-attn) 2.1e-1; sub-stage dumps: block_out 4.4e-1 with post_input_norm clean; VT_DUMP_ATTN op dumps: fa0_q rms-rel 1.196 with qkv/gate clean; VT_ROCM_ATTN_CPU_REF=1 unchanged. Fix: dispatch on q_out.dtype; bf16-out requires bf16-src (VT_CHECK). Teeth: the cross-device AttnQkNormRopeGate case gains the bf16-src -> f32-out combo at the real 0.8B dims (Hq=8, Hkv=2, Dh=256, rot=64); mutation-proven (fails with the fix reverted, passes with it). Gates on this branch (gfx1100, flock ~/gpu.lock): test_backend_cross_device 19/19 (346 assertions, incl. the new combo); test_qwen3_paged_engine 16/16 unchanged; e2e 0.8B 'The capital of France is' -> ' Paris.\nThe capital of France is' == pinned oracle == our CPU. The M4 0.8B gate with the re-derived GREEN goldens lands in the stacked sibling on row/ROCM-M4-GDN-GATE; full HIP ctest zero-delta vs base measured there. Also: three pre-existing hipcc/clang-only build breaks that block ANY full HIP build on this toolchain: same-line double CAPTURE in test_qwen3_5_gdn_spec_routing.cpp and test_ops_quant_repack.cpp, and -Wunused-const-variable on kGdnLayers (uses compile out on HIP). Debug instrumentation env-gated and off by default (VT_DUMP_ACT / VT_DUMP_ACT_SUB / VT_DUMP_ATTN; deepseek_v4 VT_DUMP_ACT precedent; env-doc-allowlist updated). FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: pi:kimi-k3 [pi]

…e, landed green-shaped with the dispatch fix stacked below Issue mudler#41 M4, second model. NEW tests/parity/test_qwen35_paged_engine.cpp -- the GDN sibling of the Qwen3-dense SACRED gate: the standard 16-prompt battery through the full paged engine, anchor + <=0.5-nat near-tie band identical to the dense gate, backend proof over the GDN op set. Goldens captured from the pinned vLLM-ROCm oracle on gfx1100 (K=10 deterministic in every cell) with a manifest.json (oracle identity + pinned HF revision 2fc06364 resolved via NEW parity::Qwen35_08BSnapshot; check-snapshot-pins clean). VERDICT: 16/16 PASS (15/16 strict token-exact vs the oracle, 1/16 near-tie band, max gap 0.125 nats, 0 forward-divergent) with the stacked AttnQkNormRopeGate output-dtype fix below. The PRE-FIX capture (13/16 forward-divergent, max gap 14.125 nats) is retained as evidence in the spec and parity ledger, not as committed goldens no code can pass. Post-review shape (localai-bot sweep): the manifest op joins PendingRunnerOps() (the qwen36_gguf_greedy precedent -- a runner-less manifest op turned main's CPU CI red); the gate exits 77 (CTest Skipped) on any non-ROCm device and when the pinned checkpoint is absent, instead of comparing a foreign engine against ROCm-derived goldens or printing SUCCESS with zero assertions; and the anchor-drift REQUIRE now reads REGRESSION SUSPECTED -- re-deriving goldens is the last step of a justified re-capture, never the response to a failure. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: pi:kimi-k3 [pi]

Three missing Qwen3.5 gate artifacts currently return normally, so doctest reports false success. This specification binds absence to exit 77 and preserves bootstrap and malformed-array behavior. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: codex:gpt-5.6-sol [codex]

Missing Qwen3.5 gate artifacts now exit 77 through one shared prerequisite helper, which prevents zero-assertion doctest success. The subprocess regression drives the actual gate executable for each required file. Bootstrap still writes token IDs, then exits skipped until the correctness artifacts exist. Malformed present arrays continue through the existing hard-failure checks. Fixes mudler#1222 FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: codex:gpt-5.6-sol [codex]

Probe mode bypasses snapshot, model, and device checks. A normal return therefore turned the completed prerequisite check into zero-assertion success. Exit 86 now marks only a completed test probe. The subprocess control distinguishes it from success and unavailable exit 77. Fixes mudler#1222 FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: codex:gpt-5.6-sol [codex]

VikashLoomba · 2026-08-18T12:24:15Z

Final integration cutoff

Final head: 4673a35dd47087cae3f6f9808e4fcf64667bcbcf, exact six-commit replay on cutoff 65d6cdae.

Qwen3.5-0.8B ran on ROCm through ROCM_ATTN: 16/16 prompts, 137 assertions; 15 strict, 1 ratified near-tie, max 0.125 nats, zero forward-divergent.
Qwen3.5 M4 gate reports success when required goldens are missing #1222 subprocess contract is mutation-proven at exact exits 77/77/77/86; malformed arrays hard-fail and bootstrap exits 77.
Removing registration and restoring source-dtype dispatch both RED; restored gates GREEN.
All 21 committed 0.8B gate artifacts and both pinned revisions remain intact.
Full preflight and fresh independent cutoff review: PASS, no findings.

Brings the branch up to `main` so GitHub can compute the merge. GitHub does not honour the `merge=union` driver that `.gitattributes` sets on `.agents/issue-index.md`, so this branch read as conflicting while `git merge-tree` reported it clean. This merge is collapsed by the squash and exists only to make the pull request mergeable. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: claude:claude-opus-5 [claude-code]

Brings the branch up to `main` so GitHub can compute the merge, and resolves the three record conflicts as unions rather than picks, per AGENTS.md: take the complete target-branch version, then re-apply the scoped edit. `.agents/benchmark-record.md` keeps both appended sections, `main`'s ENG-CUDAGRAPH-BREAK W1 entry and this branch's ROCM-GEMMA4-PREFILL-PEER-BARRIER entry, because it is an append-only log and either side alone drops a measurement. `docs/STATUS.md` keeps `main`'s M4 gate results, which landed with mudler#559, and re-appends this branch's mudler#1047 item-3 attribution sentence. `docs/USAGE.md` keeps `main`'s KEEP recipe and its "This path does" wording from mudler#676, with this branch's prefill-peer lifetime sentence re-applied ahead of it. Every conflicted path was a record or documentation file. No `src/` or `include/` path conflicted, so no product source was hand-resolved here; the product changes this merge carries in are `main`'s own. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: claude:claude-opus-5 [claude-code]

…and Gemma-4 records `main` gained a ROCm ReshapeAndCache/PagedAttention composition test (#497), the ROCm M4 near-tie gate for Qwen3.5-0.8B GDN (#559), and the Gemma-4 ROCm KEEP recipe record (#676) while this branch was under gate. None of them touches the LTX-2.5 path, the CPU GEMM seam, or the two test files this branch edits. Merged rather than left behind so that the branch's gate runs against what it will land on, and so the trailer gates — which decline to run at all while the branch is behind — have an ancestor to compare against. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [claude-code]

…sh (#1047) fix(rocm/gemma4): #839 retirement-safe prefill peer GeGLU Launch/Finish Fixes #839. ## The defect The Gemma-4 FP8 cross-device prefill helper was one monolithic function. Its donor unpinned the dequant cache immediately after enqueueing the GEMMs and before `hipEventRecord(ev_e)`, and the cache is process-wide, so a second worker could evict a zero-pin slot while the first expert stream was still reading it. Under batched-MoE prefill that is a use-after-free, and the observed symptom is a wedge after matched BEGIN/END with the process idle in `kfd_wait_on_events`. ## What changed `RunGemma4Fp8ExpertGeGLUPrefillOnExpertDevice` stays the public symbol and becomes a wrapper over `LaunchGemma4Fp8ExpertGeGLUPrefillPeer` and `FinishGemma4Fp8ExpertGeGLUPrefillPeer`. The cache pin now lives on the `PeerSlot` and is released only after the host has **observed** retirement, never after a mere enqueue: - Launch pins under the same mutex scope as `GetLocked` and stores `{cache_pin, cache_dev}` on the slot before any GEMM is enqueued. It never unpins. - Finish host-waits `ev_e`, then `hipStreamSynchronize` on the compute stream that received the output copy, and only then unpins under the lock. - A failed fill or ready-event keeps a fill lease; the lease is retired with `RetireFillLocked` after the producer-stream sync and never under `cache.mu`. - A restore failure after publish retires before it rethrows; a failed retire quarantines the pin rather than dropping it. - The same-dev arm follows the same rule. "Enqueue then unpin" is forbidden on both arms, because a `hipStreamWaitEvent` on the compute stream is not a host-side retirement proof. Peer-pipe overlap stays default OFF, so Launch and Finish run back-to-back on slot 0 and the event order is the one shipping today plus slot-scoped scratch. ## What the tests prove `tests/vt/test_ops_gemma4_prefill_peer.cpp` is 23 cases over the host lifetime model and the product source. The lifetime cases drive `HostLaunch`/`HostFinish`; the product is bound to them by source-slice gates that extract the real `Finish` and `RestoreComputeOrThrow` bodies out of `src/vt/rocm/rocm_gemma4_experts.hip`, compile them, run them, and require each mutation to red. Replacing the `hipSetDevice(compute_dev)` throw in the product `RestoreComputeOrThrow` with `(void)compute_dev;` fails three assertions. ## CI repairs in this revision Three checks were red for reasons that had nothing to do with the change under review, and all three are green on `main`. `build-test-cpu` and both `sanitize-cpu` legs failed because the test named one absolute HIP compiler from the contributor's box, `/opt/rocm-7.2.4/core-7.14/bin/hipcc`, in four places and CHECKed that it returned 0 — so any machine without that exact path **failed** rather than reporting that the gate had not run. The toolchain is now resolved (`VLLM_CPP_HIPCC`, `HIPCC`, `ROCM_PATH/bin/hipcc`, `hipcc` on PATH, `/opt/rocm/bin/hipcc`, each probed with `--version`) the way `tests/vt/test_ops_getblas_product.cpp:15-21` resolves its own precondition, and its absence prints a loud `HIP COMPILE GATE NOT RUN` banner on stderr instead of asserting. Measured on this tree: 23/23 cases and 199 assertions with no toolchain plus two banners, 23/23 and 205 assertions with `VLLM_CPP_HIPCC` set, the six recovered assertions being exactly the HIP legs. Both arms still red on the `RestoreComputeOrThrow` mutation, so the repair stopped the gate failing on absent hardware without widening it. Absent hipcc does not `exit(77)` the way `test_ops_getblas_product.cpp` does. That file is one HIP gate and nothing else, while this one carries 21 host-lifetime cases that are the only gate CI has on this change, and exiting would take those with it. The g++ leg of both compile gates still runs, so neither case can report a zero-assertion pass. `pr-size` refused the whole change with `could not classify` because nothing in `scripts/check-pr-size.py` matches `.agents/evidence/`. The five donor slices move beside the spec as `SPEC_EVIDENCE` `.log` files (`.txt` is unclassified too) and the manifest becomes `.agents/specs/rocm-gemma4-prefill-peer-helper-donor.md`. Renaming does not touch bytes and all five SHA256 values in the table verify unchanged. `docs/FEATURES.md` returns to `main`'s text. The branch had rewritten that row and lost the spec link, the `VT_GEMMA4_*`/`VT_ATTN_*` pointer and the `test_gemma4_rocm_fp8_seams` seam name in order to fit; `main`'s cell is already 219 of the 220 characters `check-public-doc-tables` allows. Nothing here owes that page: `check-doc-checkpoint` has keyed `feature_surface` off a changed `REGISTER_VLLM_MODEL` set since #595, and this change touches no file under `src/vllm/model_executor/models/`. The branch was rebuilt by rebase onto `affc2a7fd`, so it carries no untrailered merge commit and all ten original commits keep their authorship. The `docs/USAGE.md` conflict against #837's landed GetBlas text was resolved as a union: both paragraphs survive. ## What this does NOT close Named here and under `## Owed` in `.agents/specs/rocm-gemma4-prefill-peer-helper.md`. Row `ROCM-GEMMA4-PREFILL-PEER-HELPER` owns them and #839 stays open until they are discharged. - **No measurement for the new blocking retirement.** Finish now host-waits `hipEventSynchronize(ev_e)` and `hipStreamSynchronize(cst)` on every call, and the same-dev arm gains one too, where the path was fully asynchronous before. It runs per expert per layer during prefill, so it serialises a pipeline that used to overlap. The wait is required for correctness and correctness comes first, so it stays; what is missing is the number. Prefill throughput before and after on 2x R9700 needs the hardware, and no CI runner here has it. - **`PeerSlot s[2]` with only slot 0 reachable** from the production wrapper. - **The host simulator is a second implementation.** `HostLaunch`/`HostFinish` in `include/vt/rocm/rocm_gemma4_prefill_dequant_cache.h` are hand-written analogues of the product Launch/Finish, ~250 of that header's lines, compiled into every HIP build. Moving them under `tests/` is the fix; it was not done here because only a ROCm box compiles the one translation unit that consumes the header, so the move cannot be verified where this repair was made. ## Known-unrelated CI `windows-msvc-cpu` and `windows-msvc-vulkan` are red on every open PR from a break predating this branch (#503, #584). ## Landing note (maintainer) `main` was merged into this branch before the squash, and three record conflicts were resolved as unions rather than picks, per AGENTS.md: take the complete target-branch version, then re-apply the scoped edit. `.agents/benchmark-record.md` keeps both appended sections, since it is an append-only log and either side alone drops a measurement. `docs/STATUS.md` keeps the M4 gate results that landed with #559 and re-appends this branch's #1047 item-3 attribution sentence. `docs/USAGE.md` keeps the contributor KEEP recipe that landed with #676 and re-applies this branch's prefill-peer lifetime sentence ahead of it. No `src/` or `include/` path conflicted, so no product source was hand-resolved. The resolved tree was gated locally: the 42-test `test_check_gate_commands` suite, plus `check-agent-record`, `check-doc-checkpoint`, `check-issue-index-append-only`, `check-readme-structure`, `check-public-doc-tables` and `check-env-doc`, all green. On the CI verdict, stated precisely. This branch's own product changes are two files, `include/vt/rocm/rocm_gemma4_prefill_dequant_cache.h` and `src/vt/rocm/rocm_gemma4_experts.hip`, and the host lifetime suite `tests/vt/test_ops_gemma4_prefill_peer.cpp` that exercises the state machine. That host suite is what the 17 passing gates on head `de8d4760f` covered. The HIP compilation itself is not built by any lane in this repository's CI, so no CI run here has ever verified it, before or after this merge; the gfx1201 evidence is the contributor's. `windows-msvc-cpu` and `windows-msvc-vulkan` were the only two red gates and are red on every pull request here. The post-merge re-run was still queued behind 19 other jobs at merge time. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: ClaudeCode:claude-opus-5 [ClaudeCode]

VikashLoomba mentioned this pull request Aug 13, 2026

ROCm (AMD GPU) backend #41

Open

6 tasks

VikashLoomba force-pushed the row/ROCM-M4-GDN-GATE branch from 6f6ee84 to 5efcebc Compare August 13, 2026 02:12

VikashLoomba mentioned this pull request Aug 13, 2026

fix(ROCM): AttnQkNormRopeGate output-dtype dispatch -- the 0.8B GDN forward divergence root-caused and fixed (RED anchor turns GREEN) #570

Closed

VikashLoomba force-pushed the row/ROCM-M4-GDN-GATE branch from 5efcebc to a77a3b3 Compare August 14, 2026 07:38

VikashLoomba force-pushed the row/ROCM-M4-GDN-GATE branch 2 times, most recently from 30059ac to c8ec69e Compare August 14, 2026 19:12

localai-bot mentioned this pull request Aug 15, 2026

fix(GATE-OP-PARITY-MANIFEST): refuse a throwing golden by name instead of aborting the pass (#776) #853

Merged

VikashLoomba force-pushed the row/ROCM-M4-GDN-GATE branch 2 times, most recently from acea2be to 14d2850 Compare August 18, 2026 06:41

VikashLoomba mentioned this pull request Aug 18, 2026

gate(ROCM): M4 -- pinned vLLM-ROCm oracle on gfx1100 + ROCm device-golden lane in the Qwen3-0.6B near-tie gate #549

Closed

VikashLoomba changed the title ~~gate(ROCM): M4 for Qwen3.5-0.8B GDN -- the oracle-backed near-tie gate as the RED anchor of the forward-divergence fix~~ gate(ROCM): M4 for Qwen3.5-0.8B GDN -- oracle-backed green near-tie gate Aug 18, 2026

VikashLoomba marked this pull request as ready for review August 18, 2026 06:43

VikashLoomba mentioned this pull request Aug 18, 2026

Qwen3.5 M4 gate reports success when required goldens are missing #1222

Closed

VikashLoomba added 6 commits August 18, 2026 04:43

VikashLoomba force-pushed the row/ROCM-M4-GDN-GATE branch from 14d2850 to 4673a35 Compare August 18, 2026 12:22

localai-bot merged commit 7b89cf3 into mudler:main Aug 18, 2026
0 of 15 checks passed

localai-bot mentioned this pull request Aug 18, 2026

feat(rocm): grouped + non-grouped keep-quant expert GEMM — the MoE-path terminus (issue #41) #523

Open

1 task

joral mentioned this pull request Aug 19, 2026

ROCm decode is kernel-bound: the activation quantizer is 35% of GPU time and GdnPostConvK 19%, while H2D+D2H is under 1% #1294

Open

Navigation Menu

gate(ROCM): M4 for Qwen3.5-0.8B GDN -- oracle-backed green near-tie gate - #559

gate(ROCM): M4 for Qwen3.5-0.8B GDN -- oracle-backed green near-tie gate#559
localai-bot merged 7 commits into
mudler:mainfrom
VikashLoomba:row/ROCM-M4-GDN-GATE

VikashLoomba commented Aug 13, 2026 •

edited

Loading

Uh oh!

localai-bot commented Aug 13, 2026

Uh oh!

VikashLoomba commented Aug 14, 2026

Uh oh!

VikashLoomba commented Aug 18, 2026

Uh oh!

VikashLoomba commented Aug 18, 2026

Uh oh!

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants

Conversation

VikashLoomba commented Aug 13, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

M4 for Qwen3.5-0.8B (GDN) — the oracle-backed gate, landed green-shaped (issue #41)

Gate

Review-sweep rework (all three findings accepted)

Prose corrections from review

In-flow gate repair

Current-main integration cutoff (2026-08-18)

Uh oh!

localai-bot commented Aug 13, 2026

Uh oh!

VikashLoomba commented Aug 14, 2026

Uh oh!

VikashLoomba commented Aug 18, 2026

Current-main rebase and fresh review

Uh oh!

VikashLoomba commented Aug 18, 2026

Final integration cutoff

Uh oh!

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants

VikashLoomba commented Aug 13, 2026 •

edited

Loading