fix(ENG-EXPERT-STREAM): the CUDA blocker is a missing vt::MatmulBTAlphaBeta, not a slower third path, and the arm now refuses by name (#1126, #1205, #1197)

[localai-bot]

CudaBackend still does not override vt::Backend::DeviceMemoryInfo, and this
change deliberately does not add it. What it fixes is the record that says how
to close #1126, because as filed that record would produce a defect — and, after
a fresh review, the first version of this record which got the reason wrong.

The one piece of product code here is unrelated to that decision and stands on
its own: vt::MatmulBTAlphaBeta refused an unimplemented arm with a message
that named nothing, and now names it, gated.

What was already true, and what was not

#1126 was filed against a seam comment reading "ROCm/CUDA override with
hipMemGetInfo/cudaMemGetInfo". That comment no longer exists. #1132 corrected it
in both places that carried it, include/vt/backend.h:78-93 and
gemma4_moe.cpp:440-448, and recorded the gap under ## Owed in
expert-streaming.md. Every anchor either of them cites was re-verified exact:
rocm_backend.hip:338-345, gemma4_moe.cpp:440-448, gemma4_moe.cpp:506,
backend.h:78-93, platforms/cuda.cpp:71 and platforms/rocm.cpp:74. Six
examined, six correct.

So the comment half of this issue is done, and the honest question left is
whether to build the capability. The answer is no, and the reason is neither the
one the issue gives nor the one this pull request first gave.

The blocker, which was recorded nowhere

vt::MatmulBTAlphaBeta has no CUDA implementation. src/vt/fused_ops.cpp:102
dispatches to rocm::MatmulBTAlphaBetaRocm (rocm_matmul_hipblaslt.hip:516,
declared include/vt/rocm/rocm_matmul_batch.h:28) under
#if defined(VLLM_CPP_HIP) and q.device.type == kROCM, and every other
device falls through to a throw. There is no CUDA, Vulkan, Metal or CPU arm.
This is a missing kernel, not a missing build flag.

An earlier draft of this record read ExpertGeGLUDeviceAccum
(gemma4_moe.cpp:76-93) as generic — vt::MatmulBT, GeluAndMul,
vt::MatmulBTAlphaBeta — and concluded it "would run" on CUDA, just on a third
path measured by nobody. It would not run. It would throw, mid-decode:

1. EnsureGemma4Fp8ExpertOnDevice (gemma4_moe.cpp:548-586) carries no HIP guard.
2. lru.MakeRoom at :565 starts succeeding as soon as FreeBytes can answer.
3. It uploads and returns true at :575.
4. The call site at :1486 takes ExpertGeGLUDeviceAccum at :1487.
5. That calls vt::MatmulBTAlphaBeta at :90.
6. Throw.

The try/catch (...) at :563-585 wraps only the upload. The compute at
:1487 sits outside it, so the exception propagates out of the decode step
rather than degrading to the host fallback.

It is latent today only because the other route into that function, same_dev
(:730-731), needs ex.gate_up_dev, which is assigned nowhere but
rocm_gemma4_experts.hip:207,226. Nothing in a non-HIP build sets it, so the
resident arm is unreachable rather than safe. Filed as #1205.

Three further corrections to the earlier draft

ExpertGeGLUDeviceBatched is not HIP-only. It was named beside three genuine
HIP-only stubs. It sits outside the file's only #ifndef VLLM_CPP_HIP block
(:1574-1628), in an anonymous namespace, with no HIP implementation and no
header declaration, and is unconditionally return false under its own lab note
at :237-239 (gather+strided produced wrong tokens at ~23 t/s; pointer-batch
~0.8 t/s; serial/fused-gelu kept at ~34 t/s). It is disabled everywhere,
ROCm included, so listing it inflated the ROCm/CUDA asymmetry. The other three —
RunGemma4FusedTopkExpertGeGLU (:1599), PeerCopyGemma4Fp8ExpertSlice
(:1607) and RunGemma4Fp8TopKOnExpertDevice (:1611) — are correct and stay.

"The two arms sum the top-k experts in a different order" is false. Both run
inside the same for (int i = 0; i < top_k; ++i) at :1431, so the summation
order is identical. The conclusion survives — the swap would not be token-neutral
— but the mechanism is rounding and where the routing weight is applied. The
resident arm folds ww into the GEMM alpha and accumulates in the epilogue
with beta (:1434, :1442, :90). The fallback writes the unweighted
product to a BF16 buffer (:67, :70) and applies ww afterwards with separate
BF16 kernels: vt::MulScalar at :1524 on the first expert, MulScalar plus
vt::Add at :1526-1527 on the rest. Two extra BF16 roundings per expert, with
the weight applied post-rounding.

"Never step 3" overstated the case, in the spec and in this pull request's
previous title. #1126 step 3 says the fit check may "read the budget from the
backend seam", and that seam returns both halves
(bool DeviceMemoryInfo(size_t* free_bytes, size_t* total_bytes),
include/vt/backend.h:94). Only one of them is illegitimate. Reading total is
contention-independent, it is the same quantity ResidencyPolicy already
carries, and a seam that answers it on every platform is a defensible source.
The defensible statement, and the one the row now asserts: step 3 may read only
total, never free, and must keep the value probed once at platform
registration (include/vllm/platforms/interface.h:70-72).

Two smaller ones. DevExpertLru::Note also calls slots.clear() at :522; it
is a device-index reset rather than an eviction, frees nothing, and is
unreachable in a single-device process, so the #1197 row now names it and the
next reader does not conclude that sweep missed it. And the host-to-device cost
the row stated in the present tense is conditional: the device LRU is
ex.is_fp8-gated (:969, :1484), and by the row's own point (1) no per-expert
FP8 Gemma-4 checkpoint is pinned, so on CUDA the fallback is never even asked for
today.

What survives unchanged

The decision, and four of the five other substantive claims.

1. There is nothing to run it on. The LRU needs ex.is_fp8, which only
   LoadMoeFp8PerExpert sets (gemma4_weights.cpp:210-215) from a per-expert
   F8_E4M3 export. No such Gemma-4 checkpoint is pinned anywhere in
   docs/USAGE.md: all 32 lines there matching gemma case-insensitively were
   swept, and every checkpoint among them is the LTX-2.5 text tower
   (gemma4-12b-with-proj*.safetensors), not a per-expert FP8 MoE decoder.
2. The headroom test does not mean the same thing on the CUDA device this
   project gates on. MakeRoom admits iff free_b >= need + 1.5 GiB (:514),
   a constant tuned on discrete dual R9700s where free VRAM is distinct from host
   RAM. On a GB10 it is not. Qwen3.8-2.4T on --device cuda: the model loads, then the first forward dies on cudaMalloc: out of memory #1123 measured that cudaMemGetInfo's total there
   is exactly /proc/meminfo MemTotal times 1024; free is therefore reported
   over the same unified pool, and the host BF16 expert cache the device upload
   exists to relieve is drawn from that pool too. The admission test would
   double-count, on a box whose unified-memory OOM takes the host down with it.
   The free half of that is an inference from the measured total rather than
   a second measurement, and it wants confirming on the device.
3. The gap is isolated, not a pattern. Comparing overrides one by one across
   src/vt/cuda/cuda_backend.cu and src/vt/rocm/rocm_backend.hip,
   DeviceMemoryInfo is the only optional vt::Backend seam ROCm answers and
   CUDA does not.
4. The authority split. vt::Backend::DeviceMemoryInfo is a live probe,
   authoritative for a runtime admission decision and nothing else.
   ResidencyPolicy::device_memory_total_bytes is a total probed once at
   registration, 0 == UNKNOWN, authoritative for a load-time budget verdict,
   which has to be reproducible.
5. Gemma4MoE device-expert LRU: the kMaxSlots test precedes the eviction loop, so VT_GEMMA4_EXPERT_EVICT=1 goes inert after 24 admissions #1197. Reading MakeRoom line by line turned it up: the slot cap is tested
   before the eviction loop (:498 against :499-500) and EvictOne (:457) is
   the only thing that shrinks slots, so after 24 admissions
   VT_GEMMA4_EXPERT_EVICT=1 is a no-op for the life of the process. Filed and
   not fixed: the one-line repair wakes more hipFree under load, which the
   surrounding comments call a permanent kfd_wait hang, so the ordering may be
   deliberate.

The code, and why it is here

fused_ops.cpp threw a bare std::runtime_error("vt::MatmulBTAlphaBeta: ROCm-only in this build") mid-decode for an unimplemented arm. AGENTS.md
requires the opposite: refuse an unimplemented arm with a message that names the
missing part, and never leave the missing path to be discovered later. That
message named neither the device that asked nor where an implementation would
go, so a caller who hit it on CUDA could not tell a missing kernel from a missing
build flag. It now names both, plus the issue.

tests/vt/test_gemma4_rocm_fp8_seams.cpp gates it on a posed CUDA queue and on
kCPU, kVULKAN and kMETAL. The refusal reads only q.device.type and
dereferences no buffer on that path, so a CPU build can pose as any device and no
ROCm hardware is needed. This is an in-flow gateable defect independent of #1126,
and when someone writes the CUDA implementation the same test is what it has to
satisfy — which is also the answer to whether the absence should be pinned: the
gate says what is missing, executably, without any checker that says do not fix
it.

Red before green: six assertions failed against the old message. Mutation-proven
twice, both applied (sha256 changed, non-empty git diff --stat), both
compiled (compile_err=0, enospc=0), three cases each, tree restored
byte-exact. Restoring the old message goes RED. Deleting the refusal outright
goes RED.

The implementation of a CUDA MatmulBTAlphaBeta stays owed and is what #1205
tracks. It is not written here: a beta-accumulating BT GEMM on cuBLASLt is a
kernel with its own correctness gate, point 1 above says there is no checkpoint
to exercise it on, and this host has neither a ROCm nor a CUDA device to measure
either arm.

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:claude-opus-5 [claude-code]