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)

include/vt/fused_ops.h

@@ -21,6 +21,19 @@ void DualRmsNormPlusRes(Queue& q, Tensor& out, const Tensor& x1, const Tensor& w
 void GeluMulSeparate(Queue& q, void* out, const void* gate, const void* up, int64_t n,
                      DType dtype);
 
+// Does `MatmulBTAlphaBeta` have an arm for this queue's device in THIS build?
+// It answers the question a caller has to ask BEFORE committing to a device
+// path, because the alternative is finding out from a throw: the only
+// implementation in the tree is `rocm::MatmulBTAlphaBetaRocm`, so on every other
+// device — and on ROCm in a build configured without `-DVLLM_CPP_HIP` — the call
+// below refuses instead of computing (issue #1205).
+//
+// It is not a device-name test that a reader has to keep in sync by hand.
+// `MatmulBTAlphaBeta` itself dispatches on this predicate, so the two cannot
+// disagree: false here means the very next line throws, and a future CUDA arm
+// makes both true in one edit.
+bool HasMatmulBTAlphaBeta(const Queue& q);
+
 void MatmulBTAlphaBeta(Queue& q, void* out, const void* a, const void* b, int M, int N, int K,
                        float alpha, float beta, DType dtype);
 

src/vllm/model_executor/models/gemma4_moe.cpp

@@ -547,6 +547,28 @@ DevExpertLru& ExpertLru() {
 
 bool EnsureGemma4Fp8ExpertOnDevice(Dev d, const Gemma4Fp8ExpertMats& ex, int64_t I,
                                    int64_t H) {
+  // Refuse BEFORE the upload on a device whose down-projection GEMM does not
+  // exist. Returning true here is a PROMISE that the caller may run the
+  // device-resident arm, and every caller that takes that promise ends in
+  // `vt::MatmulBTAlphaBeta` — `ExpertGeGLUDeviceAccum` (:76-93) and
+  // `ExpertGeGLUTopKFusedGelu` (:181-234) both do. That call has exactly one
+  // implementation in the tree, `rocm::MatmulBTAlphaBetaRocm`, so off ROCm it
+  // throws (issue #1205). The upload's own `try`/`catch (...)` below does NOT
+  // cover the compute, so without this line the exception leaves the decode step
+  // instead of degrading: the `else` arms at the call sites already fall back to
+  // `EnsureGemma4Fp8ExpertCached` + `ExpertGeGLUHost`, which is slower and
+  // rounds twice more, but answers.
+  //
+  // It is latent rather than live today only because `MakeRoom` needs
+  // `Backend::DeviceMemoryInfo`, which only ROCm overrides. #1126 step 1 is
+  // exactly the change that adds the CUDA override, which is why the refusal has
+  // to be here before it lands and not after.
+  //
+  // Keyed on whether the arm EXISTS, not on a device name or a build macro:
+  // `vt::HasMatmulBTAlphaBeta` is the same predicate the dispatch itself uses, so
+  // writing the CUDA kernel wakes this path with no edit here, and on ROCm the
+  // answer is true and nothing about this function changes.
+  if (!vt::HasMatmulBTAlphaBeta(d.q)) return false;
   // When device LRU disabled, do NOT host-cache-dequant here — that path was
   // unbounded (every expert forever) and OOM'd the 30G host (~27G RSS) under pollution.
   if (!ExpertLru().Enabled()) return false;

src/vt/fused_ops.cpp

@@ -1,6 +1,7 @@
 #include "vt/fused_ops.h"
 
 #include <stdexcept>
+#include <string>
 
 #include "vt/backend.h"
 #include "vt/dtype.h"
@@ -98,15 +99,25 @@ void GeluMulSeparate(Queue& q, void* out, const void* gate, const void* up, int6
   b.Free(tmp);
 }
 
+bool HasMatmulBTAlphaBeta(const Queue& q) {
+#if defined(VLLM_CPP_HIP)
+  return q.device.type == DeviceType::kROCM;
+#else
+  (void)q;
+  return false;
+#endif
+}
+
 void MatmulBTAlphaBeta(Queue& q, void* out, const void* a, const void* b, int M, int N, int K,
                        float alpha, float beta, DType dtype) {
 #if defined(VLLM_CPP_HIP)
-  if (q.device.type == DeviceType::kROCM) {
+  // Dispatch on the predicate rather than on a second copy of its condition, so
+  // `HasMatmulBTAlphaBeta` cannot drift from what this function actually does.
+  if (HasMatmulBTAlphaBeta(q)) {
     rocm::MatmulBTAlphaBetaRocm(q, out, a, b, M, N, K, alpha, beta, dtype);
     return;
   }
 #endif
-  (void)q;
   (void)out;
   (void)a;
   (void)b;
@@ -116,7 +127,33 @@ void MatmulBTAlphaBeta(Queue& q, void* out, const void* a, const void* b, int M,
   (void)alpha;
   (void)beta;
   (void)dtype;
-  throw std::runtime_error("vt::MatmulBTAlphaBeta: ROCm-only in this build");
+  // Two different absences reach this line, and telling a caller the wrong one
+  // sends them to fix the wrong thing.
+  //
+  // A kROCM queue arrives here only when the build was configured without
+  // `-DVLLM_CPP_HIP`, so the 'rocm' arm exists in the tree and is compiled out.
+  // That is a build-configuration problem, not a missing kernel, and it is not
+  // #1205 — the previous "ROCm-only in this build" got exactly this case right.
+  if (q.device.type == DeviceType::kROCM) {
+    throw std::runtime_error(
+        "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.");
+  }
+  // Every other device arrives here because no such kernel was ever written:
+  // the only implementation in the tree is rocm::MatmulBTAlphaBetaRocm. Name the
+  // device that asked, name the one arm that exists, and name the issue that
+  // owes the rest — "ROCm-only" alone left the caller unable to tell a missing
+  // kernel from a missing build flag. Reaching this on CUDA is issue #1205 and
+  // blocks #1126 step 1: waking Gemma4's device-expert LRU would route decode
+  // into ExpertGeGLUDeviceAccum, which lands here outside the upload's
+  // try/catch. `EnsureGemma4Fp8ExpertOnDevice` now refuses before that upload
+  // (gemma4_moe.cpp), so this throw is the backstop rather than the guard.
+  throw std::runtime_error(
+      std::string("vt::MatmulBTAlphaBeta: no implementation for device '") +
+      DeviceTypeName(q.device.type) + "'; no '" + DeviceTypeName(q.device.type) +
+      "' kernel has been written and the only arm in the tree is 'rocm' "
+      "(src/vt/rocm/rocm_matmul_hipblaslt.hip); see issue #1205.");
 }
 
 void MatmulBTFp8Channel(Queue& q, void* out, const void* a, const void* b_fp8,

tests/CMakeLists.txt

@@ -565,6 +565,13 @@ vllm_cpp_add_test(test_gemma_load vllm/models/test_gemma_load.cpp)
 vllm_cpp_add_test(test_gemma_forward vllm/models/test_gemma_forward.cpp)
 target_include_directories(test_gemma_forward PRIVATE ${CMAKE_SOURCE_DIR}/src)
 vllm_cpp_add_test(test_gemma4_honesty vllm/models/test_gemma4_honesty.cpp)
+# #1205: the Gemma-4 device-expert upload must refuse where `vt::MatmulBTAlphaBeta`
+# has no arm, entered through `vllm::RunGemma4Moe`. Its own binary: the layer's
+# env knobs are frozen into function-local statics on the first call, so the arm
+# it measures has to be decided before main() and cannot share a process with a
+# case that wants a different one.
+vllm_cpp_add_test(test_gemma4_moe_device_arm_guard
+                  vllm/models/test_gemma4_moe_device_arm_guard.cpp)
 vllm_cpp_add_test(test_qwen3_moe_load vllm/models/test_qwen3_moe_load.cpp)
 vllm_cpp_add_test(test_qwen3_moe_forward vllm/models/test_qwen3_moe_forward.cpp)
 target_include_directories(test_qwen3_moe_forward PRIVATE ${CMAKE_SOURCE_DIR}/src)