vt::MatmulBTAlphaBeta is ROCm-only with no CUDA implementation, so #1126's step 1 would wake a throw in the Gemma4 decode step

[localai-bot]

Found while repairing the record for #1126, which claimed the opposite.

Every line anchor below is pinned to f1134c7db696e7e8eeb45a1ed3df26e961dc73af. The first version of this body was written against 0acf2ca1d and went eleven lines stale inside its own pull request, because the same change that filed it also moved the lines it cites. An unpinned anchor in an issue body is a moving target, and no gate reads an issue body.

vt::MatmulBTAlphaBeta has no CUDA implementation at all

src/vt/fused_ops.cpp:111-157 is the function; the dispatch is at :117, guarded on #if defined(VLLM_CPP_HIP) and q.device.type == DeviceType::kROCM. Every other device falls through to the refusal at :152.

The only implementation in the tree is rocm::MatmulBTAlphaBetaRocm (src/vt/rocm/rocm_matmul_hipblaslt.hip:516, declared include/vt/rocm/rocm_matmul_batch.h:28). There is no CUDA, Vulkan, Metal or CPU arm. This is not a build-configuration gap — it is a missing implementation.

A kROCM queue reaches the same line in a build configured without -DVLLM_CPP_HIP, and for that caller the opposite is true: the kernel exists and the flag is missing. Those two absences now get two different messages (src/vt/fused_ops.cpp:138 for the flag, :152 for the kernel), because telling a ROCm caller to write a kernel the tree already has sends them to fix the wrong thing.

Why it matters: it is the real blocker under #1126

#1126 proposes adding CudaBackend::DeviceMemoryInfo via cudaMemGetInfo so Gemma4MoE's device-expert LRU stops refusing on every CUDA device. Landing that override alone would not have woken a slower path. It would have woken a throw, mid-decode:

1. EnsureGemma4Fp8ExpertOnDevice (src/vllm/model_executor/models/gemma4_moe.cpp:548-608).
2. Its lru.MakeRoom(d, total) at :587 starts succeeding as soon as FreeBytes can answer.
3. It uploads and returns true at :597.
4. The call site at :1508 then takes ExpertGeGLUDeviceAccum at :1509.
5. ExpertGeGLUDeviceAccum calls vt::MatmulBTAlphaBeta at :90.
6. Throw.

The try/catch (...) at :585-607 wraps only the upload. The compute at :1509 sits outside it, so the exception propagated out of the decode step.

It is latent today only because the other route into ExpertGeGLUDeviceAccum, same_dev (gemma4_moe.cpp:752-753), requires ex.gate_up_dev, which is assigned nowhere but src/vt/rocm/rocm_gemma4_experts.hip:207,226. Nothing in a non-HIP build sets it, so the resident arm is unreachable rather than safe.

The chain is now cut at its first link

EnsureGemma4Fp8ExpertOnDevice refuses at src/vllm/model_executor/models/gemma4_moe.cpp:571 when vt::HasMatmulBTAlphaBeta(d.q) is false — before the upload, not after it — so the caller takes the host fallback already sitting in the else at :1515-1521 and the decode step answers. Slower, and two extra BF16 roundings per expert, but correct.

The predicate is not a device-name test. vt::MatmulBTAlphaBeta dispatches on HasMatmulBTAlphaBeta itself (src/vt/fused_ops.cpp:102-109), so the two cannot disagree, and writing the CUDA arm wakes the device path with nothing to change at the call site.

It is gated by tests/vllm/models/test_gemma4_moe_device_arm_guard.cpp, which enters through vllm::RunGemma4Moe — the production layer entry src/vllm/model_executor/models/gemma4.cpp:634 calls — and decorates the registered CPU backend so vt::Backend::DeviceMemoryInfo answers. That is the post-#1126 state and the only state in which the guard binds at all; on a stock build MakeRoom refuses on the absent probe and deleting the guard would be invisible. Deleting the guard turns that test RED with the exact no implementation for device 'cpu' throw.

Scope

• The missing CUDA MatmulBTAlphaBeta (a beta-accumulating BT GEMM; cuBLASLt already carries the epilogue). That is what this issue still tracks.
• CudaBackend never overrides DeviceMemoryInfo, and the seam's comment says it does — so Gemma4MoE's device-expert LRU is dead on CUDA #1126 step 1 can now land without deadlocking on it, because the guard degrades instead of throwing — but the device-resident arm stays unreachable on CUDA until the kernel exists.
• Gemma-4 MoE: EnsureGemma4Fp8NativeOnDevice has no arm-existence guard, so #1126 would still throw on the DEFAULT expert arm #1218 is the same shape in EnsureGemma4Fp8NativeOnDevice, which is the default arm and is still unguarded.
• The same file refuses six more arms the same way: MatmulBTFp8Channel at src/vt/fused_ops.cpp:177, DequantFp8ChannelBf16 at :194, MoeGatherRows at :286, MoeWeightedScatterAdd at :304, MoeZeroBf16 at :317, and DualRmsNormPlusRes at :46. Only MatmulBTAlphaBeta is on the CudaBackend never overrides DeviceMemoryInfo, and the seam's comment says it does — so Gemma4MoE's device-expert LRU is dead on CUDA #1126 path; the others are named here so the next reader does not have to re-derive the list.

Fixed in flow, partly

The refusal is gated by tests/vt/test_gemma4_rocm_fp8_seams.cpp on a posed CUDA queue, on kCPU, kVULKAN and kMETAL, and on kROCM. Be exact about what that gate is: a contract in a unit test, not observable behaviour. With the guard above in place the throw is unreachable off ROCm in any shipped configuration, so no production run prints either string. What the test fixes is what a CUDA implementation has to satisfy when someone writes one.

Owning row: ENG-EXPERT-STREAM (## Owed in .agents/specs/expert-streaming.md).