fix(ENG-EXPERT-STREAM): the CUDA blocker is a missing vt::MatmulBTAlphaBeta, and the Gemma-4 device-expert upload now refuses before it happens (#1126, #1205, #1197, #1218)

[localai-bot]

FOLLOWING_AGENTS_PROTOCOL

#1126 says adding CudaBackend::DeviceMemoryInfo would let Gemma4MoE's
device-expert LRU stop refusing on CUDA and wake a slower third path. Reading the
chain to establish that claim found something else: it would wake a throw, in
the middle of a decode step. vt::MatmulBTAlphaBeta has no CUDA implementation at
all, and that is what step 1 is actually blocked on. That is #1205, and it was
recorded nowhere.

What lands

The hazard is guarded, before the upload rather than after it.
EnsureGemma4Fp8ExpertOnDevice returns false at
src/vllm/model_executor/models/gemma4_moe.cpp:571 when
vt::HasMatmulBTAlphaBeta(d.q) is false, so the caller takes the host fallback
already sitting in the else at :1515-1521 — EnsureGemma4Fp8ExpertCached plus
ExpertGeGLUHost, slower and two extra BF16 roundings per expert — and the step
answers instead of throwing. The upload's own try/catch (...) at :585-607
never covered the compute, so without this the exception left the decode step
rather than degrading.

The predicate is not a device-name test and not a build-macro test.
vt::MatmulBTAlphaBeta dispatches on HasMatmulBTAlphaBeta itself
(src/vt/fused_ops.cpp:102-109, dispatch at :117), so the two cannot disagree:
false there means the very next line throws, and writing a CUDA arm wakes the
device path in one edit with nothing to change at the call site.

The refusal names the right absence for each caller. "ROCm-only in this build"
named neither the device that asked nor where an implementation would go, so a
caller on CUDA could not tell a missing kernel from a missing build flag. The two
messages now are:

vt::MatmulBTAlphaBeta: no implementation for device 'cuda'; no 'cuda' kernel has
been written and the only arm in the tree is 'rocm'
(src/vt/rocm/rocm_matmul_hipblaslt.hip); see issue #1205.

vt::MatmulBTAlphaBeta: the 'rocm' arm (src/vt/rocm/rocm_matmul_hipblaslt.hip) is
compiled out of this build; reconfigure with -DVLLM_CPP_HIP to enable it.

The second is for a kROCM queue, which reaches the same line in a build configured
without -DVLLM_CPP_HIP. For that caller the kernel exists and the flag is
missing; telling them to write a kernel the tree already has would send them to fix
the wrong thing. The old message got that one case right and the first draft of
this change broke it.

Say plainly what the message gate pins: a contract in a unit test, not observable
behaviour. With the guard in place the throw is unreachable off ROCm in any
shipped configuration, so no production run prints either string. What
tests/vt/test_gemma4_rocm_fp8_seams.cpp fixes is what a CUDA implementation has
to satisfy when someone writes one. That limitation is recorded under ## Owed
rather than left for a reader to discover.

Four records are corrected, and one of them was the reason for this branch. The
comment on vt::Backend::DeviceMemoryInfo claimed "(ROCm/CUDA)" in both places
that carried it. #1126's own closing plan, step 3, would source a load-time
budget verdict from a live contention-dependent probe; it is narrowed here to
total only, probed once at platform registration, rather than rejected outright.
#1197 and #1205 are filed with their chains traced.

How it is proven

tests/vllm/models/test_gemma4_moe_device_arm_guard.cpp enters through
vllm::RunGemma4Moe — the layer entry
src/vllm/model_executor/models/gemma4.cpp:634 calls — with a real
Gemma4MoeLayerWeights, the real router, the real device-expert LRU and the real
fallback. Nothing in it constructs the LRU, the Dev or vt::MatmulBTAlphaBeta by
hand.

It decorates the registered CPU backend so vt::Backend::DeviceMemoryInfo
answers, and changes nothing else. That is the whole point: the hazard is latent on
every machine that can run this suite, because MakeRoom refuses unless that probe
answers and only ROCm overrides it. On a stock build the memory probe refuses the
upload and the guard never binds, so deleting the guard would be invisible — an
unmoved-mutation trap, not a proof. #1126 step 1 IS the change that makes the
probe answer on a device with no MatmulBTAlphaBeta arm, so the test puts the tree
in exactly that state. Both directions of the probe are asserted before anything is
concluded, so an instrument that failed to arm cannot report the guard proven by a
run that never reached it.

Mutations on the merged tree. Each was applied, asserted applied by sha256 and
git diff --stat, rebuilt (a mutation that does not build is INVALID, not a pass),
run with its exit code captured directly rather than through a pipe, then restored
and re-asserted by sha256. Every one compiled.

#	Mutation	seams	guard
M1	Restore the pre-PR "ROCm-only in this build" message	RED rc 1, 4 cases / 1 failed	green
M2	Delete the refusal outright (return;)	RED rc 1, 4 cases / 2 failed	green
M3	Delete the production call site gemma4_moe.cpp:90	green	green
M4	Force HasMatmulBTAlphaBeta to true	green	RED rc 1
M5	Delete the guard	green	RED rc 1
M6	Delete the kROCM build-flag branch	RED rc 1, 4 cases / 1 failed	green
M7	Delete the guard and gemma4_moe.cpp:90	green	RED rc 1

M5 is the reachability proof, and its failure text is the evidence: the guard test
goes red with vt::MatmulBTAlphaBeta: no implementation for device 'cpu' thrown
out of vllm::RunGemma4Moe. M4 shows the guard consults the predicate on the
executed path rather than passing for some incidental reason. M7 shows :90 is not
the only production route into the refusal — with :90 gone the chain still
reaches it through ExpertGeGLUTopKFusedGelu.

M3 is reported as it came out, because it is the finding that started this
repair. Deleting gemma4_moe.cpp:90 leaves both suites green. For the guard test
that is correct and expected: with the guard in place the layer never reaches :90
at all, which is the point of the guard. For the seams suite it is the honest
consequence of what that suite is — a unit contract on a message that no correct
program reaches. A refusal a correct program never reaches has no production path
by construction, so no mutation can make one appear; the alternative would be
leaving the hazard unguarded so the string could be observed. That is stated under
## Owed rather than presented as reach.

Green after restore: test_gemma4_rocm_fp8_seams 4 cases / 37 assertions rc 0,
test_gemma4_moe_device_arm_guard 1 case / 11 assertions rc 0. Both source files
back to their exact pre-mutation sha256.

Full gate on the merged tree: cmake --build build -j 8 rc 0, 0 ENOSPC,
0 errors; ctest --test-dir build -j 4 523/523 passed, 0 failed, rc 0, with the two
pre-existing checkpoint-gated skips (test_modelopt_mixed_precision_checkpoint,
test_voxtral_e2e). An earlier merge in this sequence changed
cmake/CompilerWarnings.cmake, so that round was a clean rebuild from an empty
build/ rather than an incremental one, and it was also 523/523 rc 0.

scripts/check-symbol-anchors.py, which arrived on main during this work and
gates exactly the failure mode this row keeps producing, is green over these
records: 618 citations, in-repo checked 93, fresh 93, stale 0.

test_cpu_x86_llamacpp_floor failed once during this session under box load and
printed its own precondition: NO_QUIET_WINDOW after 30s (busy=110% builders=0 load=42.82 50.24 77.96) on a 20-core host — the load-dependent red of #618. It
is isolated rather than dismissed: the three files it reads have identical blob
SHAs on this branch and on origin/main
(scripts/cpu-x86-llamacpp-floor.sh bd7a5192…,
docs/bench-evidence/cpu-x86-llamacpp-20260811.md 94f8509c…,
tests/scripts/test_cpu_x86_llamacpp_floor.py 358927d1…), so this branch cannot
move it, and it passed on the quieter reruns.

What is owed

#1218 — EnsureGemma4Fp8NativeOnDevice (gemma4_moe.cpp:611) has the same shape
and no guard, and it is the default arm: VT_GEMMA4_FP8_NATIVE returns true
when unset (:969-974), so #1126 step 1 wakes it before it wakes the guarded one.
Filed and not fixed here, and not for effort: the BF16 guard keys on ONE predicate
that is the same condition its own dispatch uses, which is what makes it honest,
whereas the twin depends on three different ops and needs a predicate per op.
Reusing HasMatmulBTAlphaBeta there would be a guard naming the wrong arm — the
exact defect this branch's review corrected in a refusal message.

#1205 — the CUDA MatmulBTAlphaBeta itself. A beta-accumulating BT GEMM on
cuBLASLt is a kernel with its own correctness gate, no per-expert FP8 Gemma-4
checkpoint is pinned anywhere in docs/USAGE.md to exercise it on, and the host
that found this has neither a ROCm nor a CUDA device.

#1197 — the device LRU's slot cap sits ahead of its eviction loop, so
VT_GEMMA4_EXPERT_EVICT=1 goes inert after 24 admissions. Filed and not fixed: the
one-line repair wakes more hipFree under load, which the surrounding comments
record as a permanent kfd_wait hang, and deciding that needs the dual-RDNA4 box.

Anchors

This row has now produced three stale-anchor defects, and the third was not caused
by its own diff. The audit therefore states its scope and uses two instruments.

The first repair covered the four ## Owed rows in this cluster: 53 examined, 53
exact, 0 stale. It never included a GitHub issue body, which is exactly how
#1205's body came to be eleven lines stale at 7beada17c. This repair re-derived
every anchor in those four rows, in the two new ## Owed rows, in the #1197,
#1205 and #1218 index rows and in the #1205 and #1218 issue bodies:
168 endpoints across nine rows resolved for range, 0 out of range, plus 69
load-bearing anchors asserted against the substring the prose claims lives at that
line, 0 stale.

The second instrument is what earned its place. Merging origin/main moved
RocmBackend::DeviceMemoryInfo from :338-345 to :358-365, and the range-only
check reported OK because the old lines still existed — a checker that only asks
"does this line exist" passes exactly the failure it was written to find.

Several bare :NNN anchors also inherited the wrong file after a mid-sentence
mention of a different one; gemma4.cpp:634 in particular resolved to the 123-line
tool parser, where it is out of range. Those are spelled out in full.

Three .agents/issue-index.md rows are appended: #1197, #1205 and #1218. All
three were added by this branch and have not merged, so editing them while the pull
request is open cannot duplicate anything. Two index defects were caught by the
gates rather than by review and are repaired: the anchor sweep edited the MERGED
#1126 row, which is restored byte-for-byte from origin/main and deliberately
keeps its pre-move anchor because a duplicated row is worse than one stale number;
and a union merge left two byte-identical #1193 rows, one of which is removed.
The file now removes and edits nothing that main has, adds exactly three rows,
and contains no duplicate.

Closes nothing. #1126, #1197, #1205 and #1218 all stay open with their
remaining work stated above.

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:claude-opus-5-1m [Claude Code]