diff --git a/.agents/issue-index.md b/.agents/issue-index.md index d7a25b34c..d76a13b36 100644 --- a/.agents/issue-index.md +++ b/.agents/issue-index.md @@ -366,6 +366,9 @@ rather than merged. `scripts/check-agent-record.py` gates both. | [#1179](https://github.com/mudler/vllm.cpp/issues/1179) | `ENG-CUDAGRAPH-BREAK` | The hand-rolled decode-graph driver count recorded in `9bc4d7f44` is **eight** and is actually **nine**, and the row it feeds was framed as coverage-only when it is also correctness. The ninth is the DFlash draft graph, file-local with no header declaration, at `src/vllm/model_executor/models/qwen3_dflash.cpp:771,870,1038,1091,1095,1106` — its own `int g_state = 0` three-state machine (`:771`), its own `VT_DFLASH_GRAPH` kill switch (`:870`) instead of the `VLLM_CPP_CUDAGRAPH` the six batched drivers read, its own invalidate-on-block-width-change (`:1038-1047`) and its own `try { EndCaptureGraph(); } catch (...) {}` drain (`:1106`). The eight-count is stated in four places, all corrected here: [`sglang-breakable-cuda-graph.md`](specs/sglang-breakable-cuda-graph.md) §4 and `## Owed`, [`.agents/engine-matrix.md`](engine-matrix.md) rows `ENG-CUDAGRAPH-BREAK` and `ENG-CUDAGRAPH-DEDUP` ("times eight drivers", which sizes #1162's signature table), and [`.agents/roadmap_v1.md`](roadmap_v1.md) track `C12`. The reframing is the substantive half: `ENG-CUDAGRAPH-BREAK` was recorded as a COVERAGE row, and the duplication has already cost a SHIPPED model its decode graph. `src/vllm/model_executor/models/qwen3.cpp:961-986` declines the decode graph outright whenever the asynchronous device-token mirror is live, on its own measured battery — `depth-1, graph ON PASS 78/78`; `depth-2, graph OFF PASS 82/82`; `depth-2, graph ON FAIL, slots 1-3 degenerate` — because `Step()` replays against the HOST `input.token_ids` and the combine has patched the DEVICE ids. The comment names the real fix as reading the identifiers at replay time from a stable device buffer, and that fix exists, in exactly one sibling driver, as `StepDevInputs` (`src/vllm/model_executor/models/qwen3_5.cpp:3894`): `grep -c StepDevInputs` returns 41 lines there and 0 in each of `qwen3_moe.cpp`, `qwen3.cpp`, `deepseek_v2.cpp` and `voxtral.cpp`. One capability, written once, unavailable to four models, with a live mitigation standing in its place. This does NOT weaken the framing rule that `ENG-CUDAGRAPH` established: the row still makes no throughput claim, and the prefill refutation (GB10 3.8% host-idle between launches, GPU-busy >96%, 27B prefill gap 92.5% non-GEMM glue) stands unchanged. Coverage AND correctness, never speed. Fixed in flow with the [`eng-cudagraph-break.md`](specs/eng-cudagraph-break.md) review repair ([#1163](https://github.com/mudler/vllm.cpp/issues/1163)) | record | | [#1181](https://github.com/mudler/vllm.cpp/issues/1181) | `FIX-READ-F32-SCALAR-GUARD` | `ReadF32Scalar` (`src/vllm/model_executor/models/qwen3_5_weights.cpp:312-318` @ `ab6e65216`) bounds its input with `t.data != nullptr && t.nbytes >= sizeof(float)`, a LOWER bound, and then `memcpy`s four bytes into a `float`. Two silent wrong-value paths follow and neither fails: an ARRAY is reduced to element 0, so a block-wise FP8 scale grid of shape `[ceil(N/128), ceil(K/128)]` passes and stands in for the whole weight (measured under [#1166](https://github.com/mudler/vllm.cpp/issues/1166) on `Qwen/Qwen3.8-27B-FP8` @ `017b9c7af6b5689d5dd426a76e0bc077eb5ca20a`, `q_proj.weight_scale_inv` is `[96, 40]`), and ANY dtype is reinterpreted, since that same tensor is `BF16` and its four bytes are two bf16 values read as one float. Both return a finite plausible float, so the output is fluent, plausible and wrong, which is what a token gate cannot see. Upstream makes both facts structural rather than optional: a per-tensor scale is a distinct parameter TYPE that asserts `loaded_weight.shape[0] == 1` (`vllm/model_executor/parameter.py:260-272,304-309` @ `555967922`, plus the `_assert_and_load` shape assert at `:93-96`), the slot is allocated `torch.float32` so a narrow on-disk dtype is VALUE-converted rather than reinterpreted (`utils/fp8_utils.py:1276`), and the declared strategy TENSOR/CHANNEL/BLOCK picks the parameter type before a byte is read (`compressed_tensors/schemes/compressed_tensors_w8a8_fp8.py:63,128`). The AUDIT corrects the issue's own framing twice. The 27 grep hits across five files are 5 definitions, 20 call sites and 2 comment references, and both counts are short: `ReadCtF32Scalar` (`include/vllm/model_executor/models/dense_weight_loaders.h:376`) is a SIXTH copy of the same defect under another name, reached from a SIXTH model file (`src/vllm/model_executor/models/qwen3_weights.cpp:100,126-128` through `LoadCtNvfp4W4A16`). Of the six, three check nothing, `LnReadF32Scalar`/`ShReadF32Scalar` check dtype but not count, and only `nemotron_h_weights.cpp:557-573` is correct, which makes it the model the shared guard generalizes. No call site legitimately passes a multi-element or non-F32 tensor, and every existing fixture emits rank-0 or `{1}` `F32`, so nothing in the tree needed the leniency. It is NOT merely latent: `dense_weight_loaders.h:73-74` and `docs/BENCHMARKS.md:52` both record `unsloth/Qwen3.6-27B-NVFP4` @ `ccdaab7e` as FP8 W8A8 throughout with BF16 PER-OUTPUT-CHANNEL scales, and `LoadAttnDense` branches on the weight dtype alone (`qwen3_5_dense_weights.cpp:478-480`), so those projections enter the per-tensor arm and hit both defects at once under the tensor name the loader actually asked for, with no misspelling to stop them. Fixed in flow by one `dense_loaders::ReadF32Scalar(get, name)` that refuses `numel != 1` naming the shape, refuses a non-`F32` dtype naming the dtype, and requires exactly four readable bytes, with the other five copies deleted onto it and `nemotron_h`'s `Loader`-based twin kept as the one tracked exception. A narrow dtype is refused rather than converted, because a one-element BF16 scale has never been read correctly here and the BF16 layout that IS shipped is per-channel, which the count check refuses first. Per-channel FP8, block-wise FP8 and any explicit narrow-dtype conversion stay owed. Spec [`read-f32-scalar-guard.md`](specs/read-f32-scalar-guard.md) | bug | | [#1185](https://github.com/mudler/vllm.cpp/issues/1185) | `ENV-ORACLE-WHEEL-IN-LEASE` | The pinned vLLM oracle BUILDS, installs, imports and sees the GPU inside an `rc` lease on `dgx:gpu0`, measured 2026-08-18, which falsifies the `nvcc` clause four records carried. [`lease-runtime-staging.md`](specs/lease-runtime-staging.md) said the oracle "needs `nvcc`, which the worker still lacks", and `.agents/environment.md`, [`mtp-k-gt-1.md`](specs/mtp-k-gt-1.md) and [`gpu-lease-methodology.md`](specs/gpu-lease-methodology.md) each derived a blocker from it. The build job (`buildvllm.sh`, staged sha256 prefix `15e140d41f44e7c2`) asserted the checkout against the pin BEFORE compiling, printing `PIN CONFIRMED` at `5559679229bc961848b121ccdeaa8fa5d79bec98` and aborting otherwise, took `nvcc` from the toolkit row `MODEL-NEMOTRON-H-ABI-A3-E2E` staged (`NVCC_RC=0`, CUDA `release 13.3, V13.3.73`) and produced `WHEEL_RC=0`, `PERSIST_RC=0` and a 434 MiB `vllm-0.1.dev1+g555967922.cu133-cp312-cp312-linux_aarch64.whl`, sha256 `7c58b339741a288fbb313f4f5196c9c92a9e3b3c3ebe2ea970b0ff50bb9bcba4`. The identity job (`oracleenv.sh`, prefix `6119f5223f5d818c`) asserted from `cd /`, outside any source tree: `vllm.__version__ = 0.1.dev1+g555967922`, `IDENTITY_RC=0`, `cuda True NVIDIA GB10`, `CUDA_RC=0`. SCOPE, and it carries the same weight as the result: RUNNING A MODEL IS UNTESTED. Only build, install, import and `torch.cuda.is_available()` are measured, and [`mtp-k-gt-1.md`](specs/mtp-k-gt-1.md) records that the last time an oracle reached this far it consumed the host in the step AFTER `torch.compile` and REBOOTED the box, at `gpu_memory_utilization` 0.75 and again at 0.30, so the fraction is not the lever. The version string is an OPEN discrepancy: `.agents/upstream-sync.md` records `vllm_runtime_version = 0.23.1rc1.dev1511+g555967922`, the commit segment matches and satisfies the pin's binding `+g` rule, and the prefix differs because a shallow fetch stops `setuptools_scm` counting commits since the last tag, so a full-string gate needs a deeper fetch or a recorded pretend-version. The venv is NOT staged, because that job was killed at a 90-minute ceiling mid-copy and its partial tree was removed, so only the WHEEL is durable. Four staging walls, all artifacts of the NAS rather than of CUDA: `cp -a` preserves `file_mode=0664` so `nvcc` exited 126. CIFS `nounix` stores no symlink so `include` and `lib64` vanished and CMake reported `Could NOT find CUDA (missing: CUDA_INCLUDE_DIRS CUDA_CUDART_LIBRARY) (found version "13.3")`, naming the version and denying the toolkit in one line. 32 library links `libfoo.so` and `libfoo.so.MAJOR` had to be rebuilt because only the `libfoo.so.X.Y.Z` real files survived. And `markupsafe` existed as a dist-info with NO package files from a `pip --target` killed at a 35-minute ceiling, so Marlin codegen died on `ModuleNotFoundError`. The `rc` worker container is REUSED between jobs, so a repair inside a staging branch is skipped on the next run (`nvcc already in place`) and an environment repair must be unconditional and assert its postcondition. CONSEQUENCE for the rows #1129 blocked, [#1003](https://github.com/mudler/vllm.cpp/issues/1003), [#915](https://github.com/mudler/vllm.cpp/issues/915), [#821](https://github.com/mudler/vllm.cpp/issues/821) and [#81](https://github.com/mudler/vllm.cpp/issues/81): UNBLOCKED FOR THE BUILD STEP and STILL BLOCKED FOR A MODEL RUN. None can take a measurement until a model run is demonstrated. Job details, walls and non-claims in [`oracle-wheel-in-lease.md`](specs/oracle-wheel-in-lease.md) | verification | +| [#1197](https://github.com/mudler/vllm.cpp/issues/1197) | `ENG-EXPERT-STREAM` | `Gemma4MoE`'s device-expert LRU tests its slot cap BEFORE its eviction loop, so the eviction opt-in goes inert once the cap is reached. `DevExpertLru::MakeRoom` runs `if (slots.size() >= kMaxSlots) return false;` at `src/vllm/model_executor/models/gemma4_moe.cpp:498` @ `fd64c76ee`, two lines ahead of the `if (allow_evict) { while (used + need > bud && !slots.empty()) EvictOne(d); }` at `:499-500`, and `EvictOne` (`:457`, the DEVICE LRU's — the file carries a host-cache namesake at `:275`) is the only thing that shrinks `slots`. So after 24 admissions every later `MakeRoom` returns false at that first line, the eviction loop is never reached again, and `VT_GEMMA4_EXPERT_EVICT=1` becomes a no-op for the life of the process — the cache degrades permanently to the fill-only mode the opt-in exists to leave. It binds only when the slot cap is reached before the byte budget, i.e. when `24 * expert_bytes < BudgetBytes()` (below ~85.3 MiB per expert at the 2048 MiB default from `BudgetBytes`, `:416-436`, again the device one and not the host cache's at `:262`); above that the byte budget binds first and eviction behaves. Nothing reports which one happened. FILED, NOT FIXED, and not for effort: the one-line repair (move the cap test after the eviction loop, so it caps RESIDENT slots instead of stopping admission forever) wakes more `hipFree` under load, which the code's own comments call a permanent `kfd_wait` hang with the GPU idle, prefill done and no decode tokens (`:459-461` and `:486-488`), so the current ordering may be deliberate. Deciding it needs the dual-RDNA4 lab box of [`gemma4-rocm-fp8-moe.md`](specs/gemma4-rocm-fp8-moe.md); the host that found it has neither a ROCm nor a CUDA device. Found while establishing the facts for [#1126](https://github.com/mudler/vllm.cpp/issues/1126), which required reading `MakeRoom` line by line. Listed under `## Owed` in [`expert-streaming.md`](specs/expert-streaming.md) | bug | +| [#1205](https://github.com/mudler/vllm.cpp/issues/1205) | `ENG-EXPERT-STREAM` | `vt::MatmulBTAlphaBeta` has NO CUDA implementation, and that — not any of the four reasons [#1200](https://github.com/mudler/vllm.cpp/pull/1200) first recorded — is what [#1126](https://github.com/mudler/vllm.cpp/issues/1126) step 1 is blocked on. `src/vt/fused_ops.cpp:117` dispatches to `rocm::MatmulBTAlphaBetaRocm` (`src/vt/rocm/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` — `src/vt/fused_ops.cpp:111-112` is the signature, not the dispatch — and every other device falls through to a throw (`src/vt/fused_ops.cpp:152`). There is no CUDA, Vulkan, Metal or CPU arm: this is a missing kernel, not a missing build flag. #1126 proposes adding `CudaBackend::DeviceMemoryInfo` via `cudaMemGetInfo` so `Gemma4MoE`'s device-expert LRU stops refusing on CUDA; the record said that would wake a slower third path, and it would in fact wake a THROW, mid-decode. The chain: `EnsureGemma4Fp8ExpertOnDevice` (`src/vllm/model_executor/models/gemma4_moe.cpp:548-608`) -> `lru.MakeRoom` at `:587` succeeding once `FreeBytes` can answer -> `true` at `:597` -> the call site at `:1508` -> `ExpertGeGLUDeviceAccum` at `:1509` -> `vt::MatmulBTAlphaBeta` at `:90` -> throw. The `try`/`catch (...)` at `:585-607` wraps only the UPLOAD; the compute at `:1509` is outside it, so the exception leaves the decode step rather than degrading to the host fallback. Latent today only because the other route in, `same_dev` (`:752-753`), needs `ex.gate_up_dev`, assigned nowhere but `src/vt/rocm/rocm_gemma4_experts.hip:207,226` — so the resident arm is UNREACHABLE off ROCm rather than safe. FIXED IN FLOW, PARTLY. (a) The HAZARD is guarded: `EnsureGemma4Fp8ExpertOnDevice` refuses at `src/vllm/model_executor/models/gemma4_moe.cpp:571` when `vt::HasMatmulBTAlphaBeta(d.q)` is false, BEFORE the upload, so the caller takes the host fallback already sitting in the `else` at `:1515-1521` and the decode step answers instead of throwing. The predicate (`src/vt/fused_ops.cpp:102-109`) is the same condition the dispatch at `:117` uses, so the two cannot drift. Gated by `tests/vllm/models/test_gemma4_moe_device_arm_guard.cpp`, which enters through `vllm::RunGemma4Moe` and decorates the registered CPU backend so `DeviceMemoryInfo` answers — the post-#1126 state, and the only one in which the guard binds; deleting the guard turns it RED. (b) The REFUSAL now satisfies the rule that an unimplemented arm names the missing part. "ROCm-only in this build" named neither the device that asked nor where an implementation would go; a non-ROCm queue now gets all three plus this issue (`src/vt/fused_ops.cpp:152`), and a kROCM queue in a build without `-DVLLM_CPP_HIP` gets a different message naming the absent FLAG (`src/vt/fused_ops.cpp:138`), because for that caller the kernel exists. `tests/vt/test_gemma4_rocm_fp8_seams.cpp` gates both on posed queues — CUDA, `kCPU`, `kVULKAN`, `kMETAL` and kROCM — mutation-proven RED by restoring the old message, RED by deleting the refusal outright, and RED by deleting the kROCM branch. That message gate is a UNIT CONTRACT, not observable behaviour: with the guard in place the throw is unreachable off ROCm, which is the point of the guard. The IMPLEMENTATION stays open and is what this issue tracks: a `beta`-accumulating BT GEMM on cuBLASLt is a kernel with its own correctness gate, no per-expert FP8 Gemma-4 checkpoint is pinned to exercise it on, and the host that found this has neither a ROCm nor a CUDA device. The same file refuses six further arms the same way (`MatmulBTFp8Channel`, `DequantFp8ChannelBf16`, `MoeGatherRows`, `MoeWeightedScatterAdd`, `MoeZeroBf16`, `DualRmsNormPlusRes`); only `MatmulBTAlphaBeta` is on the #1126 path, and the others are named so the next reader need not re-derive the list. Listed under `## Owed` in [`expert-streaming.md`](specs/expert-streaming.md) | bug | +| [#1218](https://github.com/mudler/vllm.cpp/issues/1218) | `ENG-EXPERT-STREAM` | `EnsureGemma4Fp8NativeOnDevice` has no arm-existence guard, and it is the DEFAULT Gemma-4 expert arm, so [#1126](https://github.com/mudler/vllm.cpp/issues/1126) step 1 would still throw mid-decode after [#1205](https://github.com/mudler/vllm.cpp/issues/1205)'s guard lands. That guard covers the BF16 device-expert arm (`src/vllm/model_executor/models/gemma4_moe.cpp:571`, inside `:548-608`). The FP8-native twin at `:611` has none, and `VT_GEMMA4_FP8_NATIVE` returns true when unset (`:969-974`), so on a per-expert FP8 checkpoint the expert loop reaches the twin at `:1359` and `:1484` BEFORE the guarded arm. A `true` from it routes into `ExpertGeGLUFp8Native` (`:95-130`), which needs `vt::ExpertGeGLUFp8TopKM1` (`:105`, which merely returns false off ROCm), `vt::DequantFp8ChannelBf16` (`:117`, `:119`, refusing at `src/vt/fused_ops.cpp:194`) and `vt::MatmulBTAlphaBeta` (`gemma4_moe.cpp:128`, refusing at `src/vt/fused_ops.cpp:152`) — so T>1 throws at the first dequant and T==1 falls through the fused kernel and throws too, in both cases outside the upload's own `try`/`catch`. Latent for the same reason and for exactly as long: the twin's `MakeRoom` also needs `vt::Backend::DeviceMemoryInfo`, which only ROCm overrides (`src/vt/rocm/rocm_backend.hip:358-365`). FILED, NOT FIXED, 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 row's review had just corrected in a refusal message. Found while repairing [#1200](https://github.com/mudler/vllm.cpp/pull/1200). Listed under `## Owed` in [`expert-streaming.md`](specs/expert-streaming.md) | bug | | [#1184](https://github.com/mudler/vllm.cpp/issues/1184) | `ENG-CUDAGRAPH-DEDUP` | `VT_CUDA_GRAPH_DEDUP=1` cannot complete a single decode step once a CUDA graph is actually captured. Measured 6/6 deterministic on GB10: `vt graph dedup: captured 1 graphs, deduped to 1 execs`, then `[Qwen3DenseDecodeGraph] captured dense decode graph for padded size S=8`, then `engine-fatal: EngineCore busy loop threw: vt cuda: greedy_argmax launch: invalid device function` -- from a `greedy_argmax` launch that had SUCCEEDED. The OFF and `=0` arms are clean and byte-identical across 7 runs on the identical binary and libraries, which is what makes the asymmetry arm-attributable. `greedy_argmax` is not at fault. `src/vt/graph_dedup_runtime.h` @ `3ed6253c6` walks the retained raw graph with CUDA runtime APIs and is DESIGNED to see several of them fail -- the `cudaGraphExecUpdate` probe refusing a fold is the feature working, not an exception, and the topology walk has five more escapes that degrade the key rather than abort inside a capture -- but it never consumed the runtime's sticky per-thread error: `grep -n 'cudaGetLastError\|hipGetLastError' src/vt/graph_dedup_runtime.h` returned nothing. Twelve sites swallowed a non-success return without clearing (`:98,100,115,130,137,146,154,173,179,281,307,317`); the two that mattered are the probe (`:307,:317`), whose refusal is routine, and `cudaGraphInstantiate` (`:281`). The next unrelated kernel, launched with the ordinary `kernel<<<>>>(); Check(cudaGetLastError())` pattern, then reported our refusal as its own failure. Every symptom follows: it needs BOTH `dedup=1` and a real capture, `CUDA_LAUNCH_BLOCKING=1` does not move it because the latch is host-side and synchronous rather than a deferred async error, and `cudaGraphLaunch` returns success because a return value does not consume the latch. Fixed in flow, structurally rather than site-by-site: twelve hand-placed clears are a fix the thirteenth fallible call misses, so the clear lives in `ScopedLatchClear`'s destructor (`src/vt/graph_dedup_latch.h`, new) installed at the binding's entry points, which are exactly the six `GraphDedupOps` members; `MakeLatchGuardedOps` is the table's only constructor and takes the raw functions as template arguments, so no raw address reaches a field and a seventh operation wired elsewhere leaves its field null, which the registry constructor refuses. One line covers CUDA and HIP because there is one source. The coverage gap that hid it is closed in the same change: the device-free half of the signature walk moves to `src/vt/graph_dedup_signature.h` and is gated by the new `tests/vt/test_graph_dedup_runtime.cpp` (13 cases, 51 assertions, 7/7 negative mutations detected) -- Kahn ordering, topological re-index, sorted edge emission, the depth-4 child bound and the four graph-level escapes, none of which was reached by any test on any tier before. STATED LIMIT: a CPU test drives a fake runtime and cannot observe the real latched error, so it proves the guard's structure and not the device outcome; the device A/B re-run is owed under `## Owed` in [`eng-cudagraph-dedup.md`](specs/eng-cudagraph-dedup.md), which also records that the shipped async serving path never reaches this feature at all | bug | | [#1192](https://github.com/mudler/vllm.cpp/issues/1192) | `ENG-CUDAGRAPH-BREAK` | W1 of the break-point capture seam: `vt::BreakableGraph`, `vt::GraphCaptureScope` and `vt::GraphBreak` land in `include/vt/breakable_graph.h` + `src/vt/breakable_graph.cpp`, with the SGLang unit suite ported case for case and ONE break point registered on ONE model. **The stage's exit criterion was answered first, on a leased GPU, because W0 deliberately did not answer it**: `cudaStreamEndCapture` followed by `cudaStreamBeginCapture` on the SAME stream mid-forward with EAGER work between them is LEGAL under `cudaStreamCaptureModeThreadLocal`, our backend's mode (`src/vt/cuda/cuda_backend.cu:204-206`) — measured on `orin:gpu0` through an `rc` lease, driver 12060, three replays with fresh inputs and 0 mismatches, plus the bare zero-work re-begin (`breakable_cuda_graph.py:370-374`). The first probe REFUSED and that refusal was the probe's, not CUDA's: `dlsym` on the bare symbol name binds the LEGACY v1 entry points, which are not capture-aware, so `cuMemcpyDtoDAsync` returned `CUDA_ERROR_STREAM_CAPTURE_UNSUPPORTED` and `cuStreamBeginCapture` (v1) silently took no capture-mode argument — a broken instrument failing toward a verdict about the code. Bound by exact versioned name, the criterion holds. The break point is the DENSE ATTENTION ENTRY of `Qwen3ForCausalLM` (`src/vllm/model_executor/models/qwen3.cpp`, inside `RunLayer`), which is vLLM's boundary and not one this row invents (`splitting_ops` defaulted to the attention family, `vllm/config/compilation.py:517,764-772,1145` @ `5559679229`) registered in SGLang's form because we have no compiler to match operation names against (`layers/radix_attention.py:256` @ `f63458b5be`). The G2 reachability gate is `tests/vllm/models/test_qwen3_break_point.cpp`: it drives the production `Qwen3DenseModel::Forward` with a scope open and counts `num_hidden_layers + 1` segments, so deleting the one call site takes it to 1 and turns the gate RED — a unit test that constructs the container by hand could not have asked that. It also holds G4 in the same case: the logits are compared BIT FOR BIT against the unscoped forward, 500 values, 0 differing. **STAGED SLICE, named rather than implied:** `GraphCaptureScope` and `BreakableGraph` are not yet ENTERED from a production step, because no driver opens a scope until W2 migrates `Qwen3DenseDecodeGraph`; the break point itself is on the production path on every forward and takes the pass-through arm. Listed under `## Owed` in [`eng-cudagraph-break.md`](specs/eng-cudagraph-break.md). The capture-failure drain is NOT owed: it landed here, as behaviour and as three gated arms. The destructor `catch` alone covered a throwing `EndCaptureGraph` and nothing else, so a break function or ordinary model code throwing mid-capture left a container reporting `captured() == true` over a PARTIAL forward; comparing `std::uncaught_exceptions()` against the depth recorded at scope entry is what tells those apart, and tests 13a, 13b and 13c gate the three arms. Still owed by later stages and NOT claimed here: the auxiliary-stream auto-join before every segment close (spec D10, live at `qwen3_5.cpp:6254-6255,6384` and `laguna.cpp:2572-2576,2612`, which W4 and W5 migrate), the non-capturing-backend arm on ROCm and Tenstorrent (G5), G1 bit-exactness on a real GPU over more than one replay, and one residual the fresh review surfaced: an exception CAUGHT INSIDE the capture scope leaves the segment closed and the rest of the forward uncaptured while `captured()` stays true, which the drain cannot see because nothing is unwinding at scope exit (spec `## Owed`, W2). No throughput claim is made anywhere in this change: the prefill refutation stands (GB10 3.8% host-idle between launches, GPU-busy above 96%, 27B prefill gap 92.5% non-GEMM glue) | feature | | [#1189](https://github.com/mudler/vllm.cpp/issues/1189) | `VT-QUANT-FP8-GROUP` | Block-wise (128x128) FP8 so `Qwen/Qwen3.8-27B-FP8` runs instead of being refused. `weight_block_size` appears nowhere in `src/` or `include/`, so `469f38395` refuses the arm by name (#1166). Six independently landable milestones; M1 lands here. M1 is `vt::QuantFp8Group`, the dynamic per-token per-group activation quant, CPU and CUDA. The numerics mirror the kernel that ACTUALLY EXECUTES on a CUDA-alike platform with a contiguous input, which is the C++ custom op at `csrc/libtorch_stable/quantization/w8a8/fp8/per_token_group_quant.cu:42-96` and NOT the Triton kernel at `fp8_utils.py:95-150`: `fp8_utils.py:635-650` calls the former and returns before the latter. The two arms differ, and the difference is measurable rather than cosmetic. The CUDA kernel divides twice, `local_absmax / max_8bit` at `:68` and `static_cast(src) / y_s` at `:85`; the Triton kernel multiplies by `(1.0 / fp8_max)` at `fp8_utils.py:145` under a comment that names the 1-ULP gap. Upstream's own test tolerates the gap with `rtol=0.15` (`test_block_fp8.py:112-114`), so a value comparison cannot tell the two apart and only a byte comparison against a spelled-out reference can. `eps` is the reduction's INITIAL value (`:47`), not a post-clamp, which is what keeps an all-zero group from dividing by zero. Scope refused here and owed to later milestones: the block-scaled GEMM (M2), `Fp8BlockWeight` and the loader (M3), `Fp8BlockLinearMethod` and the Qwen3.5 wiring (M4), the mainloop-scaled CUTLASS kernel and the column-major/TMA-aligned scale layouts (M5), merged `gate_up`/QKV (M6). M1 lands UNREACHED: no production entry point dispatches `vt::QuantFp8Group` at its merge commit, M4 owns the wiring, and `.agents/specs/vt-quant-fp8-group.md` lists it under `## Owed`. The CUDA arm compiles and its on-hardware leg is owed too, because the row took no GPU lease by design: the CPU arm is the gateable one | feature | diff --git a/.agents/specs/expert-streaming.md b/.agents/specs/expert-streaming.md index f3d74317b..232114652 100644 --- a/.agents/specs/expert-streaming.md +++ b/.agents/specs/expert-streaming.md @@ -1556,7 +1556,7 @@ The probe is added to `CudaPlatform`, which already includes `` and already probes device attributes at registration, and NOT to `Backend::DeviceMemoryInfo`. That seam's comment claimed "ROCm/CUDA override with hipMemGetInfo/cudaMemGetInfo" and only ROCm does -(`src/vt/rocm/rocm_backend.hip:338-345`). The comment is **corrected in this +(`src/vt/rocm/rocm_backend.hip:358-365`). The comment is **corrected in this change**, in the two places that carried it: `include/vt/backend.h:78-93` on the seam, and `gemma4_moe.cpp:440-448` on the only call site — the second copy was found by this round's audit and is why the first correction alone would have left @@ -1606,9 +1606,14 @@ Carried debt for this row. Each item names why it is not closed here. | **`--device cuda` still cannot SERVE a larger-than-pool GGUF; it only refuses by name now.** The device-slot arm is the missing capability: a `DeviceExpertSlotStore` behind `ExpertSlotStore`, a read accessor on that interface (`KqExpertSlice` reads `HostExpertSlotStore::Slot()`, the CONCRETE class, so the seam cannot be swapped today), a device filler that is not `pread`-into-host (`ExpertSlotStore::SlotForWrite` is handed straight to `::pread`, `expert_streamer.cpp:76-94`), and lifting the `is_cpu()` guard at `qwen3_5.cpp:5578`. Sized by the measurement above: 2790 slices per token at 2,490,368 bytes is 6.95 GB per token against a 119.631 GiB pool that already holds the dense remainder. Tracked as [#1124](https://github.com/mudler/vllm.cpp/issues/1124). | It is a campaign, not a fix: W7 (the pluggable backing store) is its declared owner in the work breakdown, and the CPU arm's own I/O rate is still unmeasured on a live cache two rows above. Building a device lane on top of a host lane whose bandwidth number is void would be optimising against a number nobody has. | | **The fit bound omits everything that is not a weight.** KV cache, activations, the scratch pools and the CUDA context are not counted, so a checkpoint at 0.95x of the pool passes the refusal and still dies on the first forward. | A headroom fraction invented here would be exactly the guess the per-tensor bound exists to avoid. The number wants the startup memory profile that `KV-WARMUP-PROFILE` owns (`INVENTORIED`; upstream's is `GPUWorker.determine_available_memory`, `vllm/v1/worker/gpu_worker.py:451-495`, around `profile_run`, `vllm/v1/worker/gpu/model_runner.py:682`), which is a different row. Those two anchors are stated here from the pinned tree because that row's own three anchors are stale at the current pin, and `gguf_device_fit.h` had copied two of them — filed as [#1139](https://github.com/mudler/vllm.cpp/issues/1139), owned by `KV-WARMUP-PROFILE`, blocked here only by the `engine-matrix.md` record lock #1119 holds. | | **The fit bound also counts too MUCH, and that direction can refuse a load that fits.** A tensor present in the file and not staged by THIS load is a positive over-count. On a default load that is the MTP / `nextn` block: 8,940,488,704 bytes, 8.33 GiB, 2.2506 % of the target checkpoint. A budget in that window refuses a weight set that would have fitted. | Not closed here. Closing it means the bound taking a per-tensor staging POLICY as input, which is the caller's knowledge and not the file's, and the exclusion's own failure mode is an under-count to nothing — which restores the 26-minute-then-OOM this row exists to remove, on a device nobody here has to measure it on. So the direction is stated in `gguf_device_fit.h`, pinned executably by `test_gguf_device_fit`, exposed to operators in `docs/USAGE.md`, and tracked as [#1136](https://github.com/mudler/vllm.cpp/issues/1136). `VT_DEVICE_WEIGHT_BUDGET_BYTES` is the way out of the window in the meantime. | -| **`Backend::DeviceMemoryInfo` has no CUDA override.** Only ROCm implements it (`src/vt/rocm/rocm_backend.hip:338-345`), so `Gemma4MoE`'s device-expert LRU refuses on every CUDA device and falls back to host H2D silently. | The CAPABILITY is [#1126](https://github.com/mudler/vllm.cpp/issues/1126) and is not built here, because adding the override wakes a landed residency policy that needs its own measurement. What IS done here is the false COMMENT, corrected in both places that carried it: `include/vt/backend.h:78-93` and `gemma4_moe.cpp:440-448`. The second copy was found by this round's own audit; correcting only the seam would have left the claim in the tree, which is the shape this row keeps hitting. | +| **`Backend::DeviceMemoryInfo` has no CUDA override, and waking it is not the one-line port of the ROCm one that #1126 describes: on CUDA it would wake a THROW.** Only ROCm implements it (`src/vt/rocm/rocm_backend.hip:358-365`), so `Gemma4MoE`'s device-expert LRU refuses on every CUDA device (`gemma4_moe.cpp:506`). Where a per-expert FP8 checkpoint is present, the expert's BF16 bytes are re-copied HOST->DEVICE on every use instead (`ExpertGeGLUHost`, `gemma4_moe.cpp:49-74`, reached at `:1515-1521`; the H2D is `:59-60` and it drains the queue per expert at `:73`), silently and for the life of the process. That H2D cost is CONDITIONAL, not present-tense: the whole device LRU is `ex.is_fp8`-gated (`:991`, `:1506`), and by point (1) below no such Gemma-4 checkpoint is pinned anywhere, so on CUDA today the LRU-MISS fallback is never even asked for. (A BF16 Gemma-4 checkpoint reaches `ExpertGeGLUHost` at `:1525` too, but through the host-weight branch the LRU never governs, so it is not a cost of the missing probe.) The name misleads: it computes on the DEVICE from host-resident weights, so what the dead LRU would cost is bandwidth and a per-expert `Synchronize`, not a wrong answer. | The CAPABILITY is [#1126](https://github.com/mudler/vllm.cpp/issues/1126) and is still not built. Note that the MERGED `#1126` index row cites `rocm_backend.hip:338-345`, which a later commit moved to `:358-365`. That row is append-only and is deliberately left alone: editing a merged row makes a union merge DUPLICATE it rather than merge it, which is a worse outcome than one stale number. This spec row carries the current anchor and is authoritative for it. The false COMMENT was corrected by this row in both places that carried it, `include/vt/backend.h:78-93` and `gemma4_moe.cpp:440-448`; both anchors were re-verified exact against `fd64c76ee`, as were `rocm_backend.hip:358-365`, `gemma4_moe.cpp:506`, `platforms/cuda.cpp:71` and `platforms/rocm.cpp:74`. SCOPE, because the previous headline read as coverage it did not have. The first review repair re-audited every anchor cited by the FOUR `## Owed` rows in this cluster — this one, the #1126 step-3 row, the #1197 row and the #1205 row — against the repaired tree: 53 examined, 53 exact, 0 stale. It audited those four rows and nothing else. The GitHub ISSUE BODIES were never in the audited set, and #1205's body was in fact 11 lines stale at `7beada17c` for exactly that reason. The second review repair widened the set: it re-derived every anchor in those four rows AND in the #1197 and #1205 issue bodies AND in the #1205 index row against the final tree — 71 examined, 71 exact, 0 stale after repair. Anchors from `gemma4_moe.cpp:549` onward moved by 22 lines in that repair, because the arm-existence guard it added sits at `:571`. **Four things were established while re-reading it for #1126, and each one raises the price of the override.** (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, so the woken path cannot be exercised, here or elsewhere, until one is. (2) *The device-resident arm has no CUDA implementation to route into. It has a throw.* `ExpertGeGLUDeviceAccum` (`gemma4_moe.cpp:76-93`) READS as generic — `vt::MatmulBT`, `GeluAndMul`, `vt::MatmulBTAlphaBeta` — and an earlier draft of this row concluded from that reading that it "would run". It does not. `vt::MatmulBTAlphaBeta` (`src/vt/fused_ops.cpp:111-157`, dispatching at `:117`) is guarded on `#if defined(VLLM_CPP_HIP)` AND `q.device.type == kROCM`; its 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`), and every other device falls through to a refusal. So the chain the override WOULD wake is: `EnsureGemma4Fp8ExpertOnDevice` (`gemma4_moe.cpp:548-608`) -> `lru.MakeRoom` (`:587`) succeeding as soon as `FreeBytes` can answer -> `true` at `:597` -> the call site at `:1508` -> `ExpertGeGLUDeviceAccum` at `:1509` -> `vt::MatmulBTAlphaBeta` at `:90` -> THROW, mid-decode. The `try`/`catch (...)` at `:585-607` wraps only the UPLOAD; the compute at `:1509` sits outside it, so the exception would propagate out of the decode step rather than degrading to the host fallback. **That chain is now cut at its first link.** `EnsureGemma4Fp8ExpertOnDevice` refuses at `:571` when `vt::HasMatmulBTAlphaBeta(d.q)` is false, BEFORE the upload, so the caller takes the `else` at `:1515-1521` — `EnsureGemma4Fp8ExpertCached` plus `ExpertGeGLUHost` — and the step answers instead of throwing. The refusal at `:90` stays as the backstop. It is latent today only because the other route into that function, `same_dev` (`:752-753`), needs `ex.gate_up_dev`, which is assigned nowhere but `src/vt/rocm/rocm_gemma4_experts.hip:207,226` — so the resident arm is UNREACHABLE off ROCm rather than safe. This is the actual blocker under #1126, it was recorded nowhere, and it is a stronger argument than the other three: filed as [#1205](https://github.com/mudler/vllm.cpp/issues/1205), and the refusal itself is now gated by `tests/vt/test_gemma4_rocm_fp8_seams.cpp`, which is what a CUDA implementation will have to satisfy. **Two corrections to the earlier draft's supporting claims, both of which overstated the case.** *The HIP-only list was one symbol too long.* Three of the four are genuinely HIP-only stubs inside `gemma4_moe.cpp`'s ONLY `#ifndef VLLM_CPP_HIP` block (`gemma4_moe.cpp:1596-1650`): `RunGemma4FusedTopkExpertGeGLU` (`:1621`), `PeerCopyGemma4Fp8ExpertSlice` (`:1629`) and `RunGemma4Fp8TopKOnExpertDevice` (`:1633`). `ExpertGeGLUDeviceBatched` (`:240`) is NOT: it sits OUTSIDE that block, in an anonymous namespace, with no HIP implementation and no header declaration, and is unconditionally `return false` under its own lab note (`: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 naming it beside the three inflated the ROCm/CUDA asymmetry. *The token-neutrality argument had the wrong mechanism.* The conclusion stands — the swap would not be token-neutral — but not because "the two arms sum the top-k experts in a different order". They do not: both run inside the SAME `for (int i = 0; i < top_k; ++i)` at `:1453`, so the summation order is identical. The real difference 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` (`:1456`, `:1464`, `:90`), so the weight multiplies in the GEMM's own accumulator. The fallback writes the UNWEIGHTED product to a BF16 buffer (`:67`, `:70`) and applies `ww` afterwards with separate BF16 kernels — `vt::MulScalar` at `:1546` on the first expert, `MulScalar` plus `vt::Add` at `:1548-1549` on the rest. Two extra BF16 roundings per expert, with the weight applied post-rounding. (3) *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 a quantity distinct from host RAM. On a GB10 it is not. This row already measured that `cudaMemGetInfo`'s `total` there is EXACTLY `/proc/meminfo MemTotal` times 1024, which is why it reached for that instrument where `nvidia-smi` answers `[N/A]`; its `free` is therefore reported over the same unified pool, and the host BF16 expert cache the device upload exists to relieve (`ex.cached_gu`/`cached_dn` via `EnsureGemma4Fp8ExpertCached`, bounded by the host LRU at `gemma4_moe.cpp:352`) is drawn from that pool too. So the admission test would double-count, on a box whose unified-memory OOM takes the host down with it. The `free` half is an inference from the measured `total`, not a second measurement, and it wants confirming on the device before any override lands. (4) *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` at `fd64c76ee`, `DeviceMemoryInfo` is the ONLY optional `vt::Backend` seam ROCm answers and CUDA does not. | +| **#1126's own closing plan, step 3, must be NARROWED before it is done: the load-time fit check may read the seam's `total`, and never its `free`.** The issue proposes that "the #1123 fit check can then read the budget from the backend seam on every platform that reports one". As written that invites the live half, which would be a defect — but the seam returns BOTH halves (`bool DeviceMemoryInfo(size_t* free_bytes, size_t* total_bytes)`, `include/vt/backend.h:94`), and only one of them is illegitimate here. | The tree holds two answers to "how much device memory", and they answer different questions. `vt::Backend::DeviceMemoryInfo(free, total)` is a LIVE probe that moves with contention; it is authoritative for a RUNTIME admission decision — can this allocation succeed right now — and for nothing else. `vllm::platforms::ResidencyPolicy::device_memory_total_bytes` is a TOTAL probed once at platform registration with `0 == UNKNOWN`; it is authoritative for a LOAD-TIME budget verdict, which has to be reproducible and independent of whatever else the box is doing. Sourcing the load-time verdict from `free` would make the same checkpoint load or be refused depending on the page cache, which is precisely the property **`total`, not `free`** above was chosen to avoid. Reading `total` through the seam is not that. It is contention-independent, it is the same quantity `ResidencyPolicy` already carries, and a seam that answers it on every platform is a defensible place to source it from. What step 3 must additionally preserve is the PROBE-ONCE semantics: `device_memory_total_bytes` is probed at platform registration (`include/vllm/platforms/interface.h:70-72`), and a per-load live call would reintroduce the contention dependence by the back door even reading only `total`, because a load-time verdict has to be reproducible from the record rather than from the moment. So the defensible statement, and the one this row asserts: **step 3 may read only `total`, never `free`, and must keep the value probed once at platform registration.** An earlier draft of this row said "never step 3" outright; that overstated it and would have blocked a legitimate simplification, so it is corrected here rather than quietly narrowed. Both seams already carry the division in prose (`include/vt/backend.h:90-93` and `include/vllm/platforms/interface.h:61-69`); it is restated here because #1126 is the record a reader of that issue will act on, and as filed it points the other way. What #1126 owes is its steps 1 and 2 together — the override AND the Gemma4 measurement, with [#1205](https://github.com/mudler/vllm.cpp/issues/1205) ahead of both — never step 1 alone, and step 3 only in the narrowed form above. | +| **The device-expert LRU's slot cap makes its own eviction opt-in inert.** `MakeRoom` tests `slots.size() >= kMaxSlots` (`gemma4_moe.cpp:498`) BEFORE the eviction loop (`:499-500`), and `EvictOne` (`:457`, the device LRU's — a host-cache namesake sits at `:275`) is the only thing that SHRINKS `slots`. The one other statement that touches its size, `slots.clear()` in `DevExpertLru::Note` (`:522`), is a device-index RESET rather than an eviction: it drops bookkeeping when `dev != d.q.device.index` and frees nothing, and it is unreachable in a single-device process. It is named here so the next reader does not conclude the #1197 sweep missed it. So once 24 slots are resident `VT_GEMMA4_EXPERT_EVICT=1` never runs again and the cache degrades permanently to fill-only. It binds only when `24 * expert_bytes < BudgetBytes()`, so it is condition-dependent and silent either way. Tracked as [#1197](https://github.com/mudler/vllm.cpp/issues/1197). | Filed, not fixed, and for the same reason as the row above rather than for effort: the one-line repair wakes more `hipFree` under load, which the surrounding comments say has been observed as a permanent `kfd_wait` hang with the GPU idle and no decode tokens. The current ordering may well be deliberate belt-and-braces. Deciding that needs the dual-RDNA4 box `.agents/specs/gemma4-rocm-fp8-moe.md` describes; this host has neither a ROCm nor a CUDA device. It closes when the cap moves after the eviction loop and a run stays hang-free, or when the comment says the cap is by design — one of the two, not silence. | +| **`vt::MatmulBTAlphaBeta` is ROCm-only and has no CUDA implementation at all, which is what #1126 step 1 is actually blocked on.** `src/vt/fused_ops.cpp:117` dispatches to `rocm::MatmulBTAlphaBetaRocm` (`src/vt/rocm/rocm_matmul_hipblaslt.hip:516`) under `#if defined(VLLM_CPP_HIP)` and `q.device.type == kROCM` — `src/vt/fused_ops.cpp:111-112` is the signature, not the dispatch — and every other device falls through to the refusal at `src/vt/fused_ops.cpp:152`. There is no CUDA, Vulkan, Metal or CPU arm. The full chain from the missing `DeviceMemoryInfo` override to that refusal is traced in the first row above. Tracked as [#1205](https://github.com/mudler/vllm.cpp/issues/1205). | The REFUSAL is fixed in flow, because a bare `std::runtime_error` reading "ROCm-only in this build" does not satisfy the standing rule that an unimplemented arm refuses with a message NAMING the missing part: a caller who hits it on CUDA cannot tell a missing kernel from a missing build flag. It now names the device that asked, names the one arm that exists, and names the issue (`:152`), and a kROCM queue — which reaches the same line in a build configured without `-DVLLM_CPP_HIP` — gets a DIFFERENT message naming the absent build flag (`:138`), because for that caller the kernel exists and telling them to write one would send them to fix the wrong thing. `tests/vt/test_gemma4_rocm_fp8_seams.cpp` gates both messages on a posed CUDA queue, on `kCPU`/`kVULKAN`/`kMETAL`, and on kROCM — mutation-proven by restoring the old message (RED), by deleting the refusal outright (RED), and by deleting the kROCM branch so that case falls to the generic message (RED). **Say plainly what that message change does and does not pin: a contract in a unit test, not observable behaviour.** The throw is unreachable off ROCm in any shipped configuration, so no production run can print either string today; what the test fixes is what a CUDA implementation has to satisfy when someone writes one. **The reachable half of this row is the GUARD.** `EnsureGemma4Fp8ExpertOnDevice` refuses at `gemma4_moe.cpp:571` when `vt::HasMatmulBTAlphaBeta(d.q)` is false, BEFORE the upload rather than after it, which converts the mid-decode exception traced above into the host fallback that was already sitting in the `else` at `:1515-1521`: slower, two extra BF16 roundings per expert, and correct. The predicate (`include/vt/fused_ops.h`, defined `src/vt/fused_ops.cpp:102-109`) is the same condition the dispatch at `:117` uses rather than a second copy of it, so the two cannot drift and writing the CUDA kernel wakes the device arm with no edit 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 `DeviceMemoryInfo` ANSWERS, which is the post-#1126 state and the only state in which the guard binds at all. Deleting the guard makes that test RED with the exact `no implementation for device 'cpu'` throw; forcing `HasMatmulBTAlphaBeta` to `true` makes it RED too. A test that constructed the `Dev` or the LRU by hand would have stayed green under both. The IMPLEMENTATION stays owed and is what [#1205](https://github.com/mudler/vllm.cpp/issues/1205) tracks. It is not written here: a `beta`-accumulating BT GEMM on cuBLASLt is a kernel with its own correctness gate, the `DeviceMemoryInfo` row's 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. | | **`model_loader.cpp` is cited by absolute line number from 109 sites in 45 files, and this row's change moved them.** Measured between `e7d0a1f7c` and the repaired head: 203 moved line references over 109 citing sites, 10 unmoved. The file is ~1640 lines and almost every engine and model row edits it, so any edit near its top invalidates citations in files the editing change never opens. | Not swept here, deliberately, and the reason is not effort: several of the 109 were ALREADY stale (`model-matrix.md:197` cites `:184-223` as the "live loader"; line 184 at `e7d0a1f7c` is `static const bool once = [] {`), and rewriting all of them from the current tree would launder pre-existing debt into a clean-looking record. What IS fixed here is the two anchors this change authored itself, checked against the final tree. Tracked as [#1143](https://github.com/mudler/vllm.cpp/issues/1143), which lists the three candidate fixes; it needs a row of its own and is parked here because this row is what measured it. | | **The budget knob is an environment variable, not a config key.** `VT_DEVICE_WEIGHT_BUDGET_BYTES`. | `ENG-RESIDENCY-CONFIG` ([#1110](https://github.com/mudler/vllm.cpp/issues/1110), PR #1119) is in flight and adds exactly the `vllm_cpp` namespace inside `--offload-config` this key belongs in. Landing a second, competing config surface while that one is unmerged would create the conflict both changes then have to resolve. Migrate once #1119 lands; tracked as [#1127](https://github.com/mudler/vllm.cpp/issues/1127). | +| **`EnsureGemma4Fp8NativeOnDevice` has the same missing-arm shape and no guard, and it is the DEFAULT arm.** The guard this row added covers the BF16 device-expert arm (`gemma4_moe.cpp:571`). Its FP8-native twin at `:611` does not have one, and `VT_GEMMA4_FP8_NATIVE` defaults to TRUE (`:969-974`), so on a per-expert FP8 checkpoint the expert loop reaches the twin at `:1359` and `:1484` FIRST. A `true` from it routes into `ExpertGeGLUFp8Native` (`:95-130`), which needs `vt::DequantFp8ChannelBf16` (`:117`, `:119`; refuses at `src/vt/fused_ops.cpp:194`) and `vt::MatmulBTAlphaBeta` (`gemma4_moe.cpp:128`; refuses at `src/vt/fused_ops.cpp:152`). Latent for the same reason and for exactly as long: its `MakeRoom` also needs `Backend::DeviceMemoryInfo`, so #1126 step 1 wakes this arm BEFORE it wakes the guarded one. Tracked as [#1218](https://github.com/mudler/vllm.cpp/issues/1218). | Not fixed in flow, and not for effort. The BF16 guard keys on ONE predicate that is the same condition its dispatch uses, which is what makes it honest. The twin depends on three different ops, so an honest guard for it needs a predicate per op; reusing `HasMatmulBTAlphaBeta` there would be a guard naming the wrong arm, which is the defect this row's own review just corrected in a refusal message. That is a distinct change with its own gate. Recording it is what stops the default arm being discovered by whoever lands #1126. | +| **A production-entered gate for the guard exists; a production-entered gate for the REFUSAL MESSAGE does not, and cannot be built here.** `test_gemma4_moe_device_arm_guard.cpp` drives `vllm::RunGemma4Moe`, so the guard is measured as a capability. The message itself is only reachable when the guard is absent, which is precisely what that test forbids, so the message's own gate is a unit contract on a posed `vt::Queue`. | This is a property of the fix, not a gap in the test. A refusal that a correct program never reaches has no production path by construction; the alternative would be to leave the hazard unguarded so the string could be observed. Naming it here so no later reader reads the seams suite as a reachability proof. Closed when a CUDA `MatmulBTAlphaBeta` lands under [#1205](https://github.com/mudler/vllm.cpp/issues/1205) and the message stops being the answer at all. | ## Risks/decisions diff --git a/include/vt/fused_ops.h b/include/vt/fused_ops.h index ceb31e8d0..c051b70c3 100644 --- a/include/vt/fused_ops.h +++ b/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); diff --git a/src/vllm/model_executor/models/gemma4_moe.cpp b/src/vllm/model_executor/models/gemma4_moe.cpp index ede647835..31c711cb9 100644 --- a/src/vllm/model_executor/models/gemma4_moe.cpp +++ b/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; diff --git a/src/vt/fused_ops.cpp b/src/vt/fused_ops.cpp index 236c97e40..209cef360 100644 --- a/src/vt/fused_ops.cpp +++ b/src/vt/fused_ops.cpp @@ -1,6 +1,7 @@ #include "vt/fused_ops.h" #include +#include #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, diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt index 8dd1b3e7c..cccec93f6 100644 --- a/tests/CMakeLists.txt +++ b/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) diff --git a/tests/vllm/models/test_gemma4_moe_device_arm_guard.cpp b/tests/vllm/models/test_gemma4_moe_device_arm_guard.cpp new file mode 100644 index 000000000..1c6a92f27 --- /dev/null +++ b/tests/vllm/models/test_gemma4_moe_device_arm_guard.cpp @@ -0,0 +1,312 @@ +// #1205 / #1126 step 1: `EnsureGemma4Fp8ExpertOnDevice` must REFUSE the device +// upload on a device whose down-projection GEMM does not exist, instead of +// admitting and then throwing out of the middle of a decode step. +// +// WHAT MAKES THIS A REACHABILITY GATE AND NOT A UNIT TEST. The guard is entered +// through `vllm::RunGemma4Moe` — the Gemma-4 MoE layer entry point that +// `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 here constructs the +// LRU, the `Dev`, or `vt::MatmulBTAlphaBeta` by hand. Delete the `if +// (!vt::HasMatmulBTAlphaBeta(d.q)) return false;` line from +// `EnsureGemma4Fp8ExpertOnDevice` and this file goes red, because the layer then +// throws `vt::MatmulBTAlphaBeta: no implementation for device 'cpu'` out of +// `ExpertGeGLUTopKFusedGelu`. +// +// WHY THE BACKEND IS DECORATED. The hazard is latent on every machine that can +// run this suite, and for exactly one reason: the LRU's `MakeRoom` refuses +// unless `vt::Backend::DeviceMemoryInfo` answers, and only ROCm overrides it +// (`src/vt/rocm/rocm_backend.hip`). So on a stock CPU or CUDA build the upload +// is refused by the memory probe and the guard never binds — an +// unmoved-mutation trap, not a proof. #1126 step 1 IS the change that makes +// `DeviceMemoryInfo` answer on a device with no `MatmulBTAlphaBeta` arm, so this +// test puts the tree in exactly that state: the registered CPU backend, wrapped, +// with `DeviceMemoryInfo` answering and nothing else altered. Every op still +// runs its real CPU kernel; `vt::MatmulBTAlphaBeta` still has no CPU arm. That +// combination is the one the guard exists for, and it is the only combination +// under which deleting the guard is observable. +// +// The decoration is a real `vt::Backend` registered through the public +// `vt::RegisterBackend`, restored in a destructor, and the test asserts BOTH +// directions of the probe (real backend: false; wrapped: true) before it +// concludes anything — an instrument that silently failed to arm would otherwise +// report the guard as proven when the run never reached it. +// +// `vt::RegisterBackend` documents that all registration completes before +// `main()`, and this call is after it. What that contract buys is the absence of +// synchronization on the read path: the registry is a plain pointer store and +// lock-free reads thereafter. This binary is single-threaded across the swap — +// one doctest case, no queue with a worker, no other case touching the CPU +// backend — so there is no concurrent reader to race, and the pointer is put +// back before the case returns. It is the narrowest way to reach the state the +// guard exists for; the alternative is a ROCm device, which no host running this +// suite has. +#include "vllm/model_executor/models/gemma4_moe.h" + +#include + +#include +#include +#include +#include +#include +#include + +#include "vt/backend.h" +#include "vt/device.h" +#include "vt/dtype.h" +#include "vt/fused_ops.h" +#include "vt/tensor.h" + +namespace { + +using vt::DType; + +// Every knob `RunGemma4Moe` reads lands in a function-local `static const`, so +// the FIRST call in the process freezes it. Set them during static +// initialization rather than inside a TEST_CASE, so no ordering between cases +// can decide which arm the suite measured. +struct EnvOnce { + EnvOnce() { + // The BF16 device-expert LRU arm is the one #1205 traces and the one this + // guard covers. Its FP8-native twin (`EnsureGemma4Fp8NativeOnDevice`, the + // default) carries the same latent shape through `vt::DequantFp8ChannelBf16` + // and is NOT guarded here — filed as its own issue and recorded under + // `## Owed` in .agents/specs/expert-streaming.md. + ::setenv("VT_GEMMA4_FP8_NATIVE", "0", 1); + // A definite budget: the default is 2048 MiB, but reading it from a default + // makes the precondition invisible in the log. + ::setenv("VT_GEMMA4_EXPERT_VRAM_MB", "256", 1); + // Everything else must stay off — each of these reroutes the expert loop. + ::unsetenv("VT_GEMMA4_HOST_AXPY"); + ::unsetenv("VT_GEMMA4_BATCH_EXPERTS"); + ::unsetenv("VT_GEMMA4_CUSTOM_EXPERT"); + ::unsetenv("VT_GEMMA4_EXPERT_EVICT"); + ::unsetenv("VT_GEMMA4_PREFILL_BATCH_MOE"); + ::unsetenv("VT_GEMMA4_PROFILE"); + } +}; +const EnvOnce g_env_once; + +// The registered CPU backend with ONE behaviour changed: the optional device +// memory probe answers. Everything else forwards, so the allocations, copies and +// kernels under test are the production CPU ones. +class ProbingBackend final : public vt::Backend { + public: + explicit ProbingBackend(vt::Backend& inner) : inner_(inner) {} + + void* Alloc(size_t bytes) override { return inner_.Alloc(bytes); } + void Free(void* p) override { inner_.Free(p); } + void Memset(vt::Queue& q, void* p, int value, size_t bytes) override { + inner_.Memset(q, p, value, bytes); + } + void Copy(vt::Queue& q, void* dst, const void* src, size_t bytes) override { + inner_.Copy(q, dst, src, bytes); + } + vt::Queue CreateQueue() override { return inner_.CreateQueue(); } + void DestroyQueue(vt::Queue& q) override { inner_.DestroyQueue(q); } + void Synchronize(vt::Queue& q) override { inner_.Synchronize(q); } + void FlushPending() override { inner_.FlushPending(); } + bool UnifiedMemory() const override { return inner_.UnifiedMemory(); } + bool DeviceMemoryIsHostAddressable() const override { + return inner_.DeviceMemoryIsHostAddressable(); + } + bool SupportsAuxStream() const override { return inner_.SupportsAuxStream(); } + bool SupportsAsyncSampledTokenReadback() const override { + return inner_.SupportsAsyncSampledTokenReadback(); + } + bool SupportsCompressedConvState() const override { + return inner_.SupportsCompressedConvState(); + } + bool SupportsCompressedGdnState() const override { + return inner_.SupportsCompressedGdnState(); + } + bool SupportsGraphCapture() const override { return inner_.SupportsGraphCapture(); } + int DeviceCapabilityMajor() const override { return inner_.DeviceCapabilityMajor(); } + int DeviceCapabilityMinor() const override { return inner_.DeviceCapabilityMinor(); } + + // The one override. Numbers chosen well above `MakeRoom`'s 1.5 GiB headroom so + // the admission decision turns on the guard and not on arithmetic. + bool DeviceMemoryInfo(size_t* free_bytes, size_t* total_bytes) const override { + if (free_bytes != nullptr) *free_bytes = 8ull << 30; + if (total_bytes != nullptr) *total_bytes = 16ull << 30; + return true; + } + + private: + vt::Backend& inner_; +}; + +struct BackendSwap { + vt::Backend* prev; + explicit BackendSwap(vt::Backend& replacement) + : prev(&vt::GetBackend(vt::DeviceType::kCPU)) { + vt::RegisterBackend(vt::DeviceType::kCPU, &replacement); + } + ~BackendSwap() { vt::RegisterBackend(vt::DeviceType::kCPU, prev); } + BackendSwap(const BackendSwap&) = delete; + BackendSwap& operator=(const BackendSwap&) = delete; +}; + +constexpr int64_t kE = 4; // experts +constexpr int64_t kI = 16; // expert intermediate +constexpr int64_t kH = 32; // hidden +constexpr int kTopK = 2; + +// Finite, well-formed F8_E4M3 codes. Random bytes would put the e4m3 NaN +// encodings (0x7F/0xFF) into the weights, which propagate through the GEMM and +// make every later comparison vacuous. +constexpr uint8_t kF8Codes[] = {0x38, 0x3C, 0x30, 0xB8, 0x34, 0x40, 0xB0, 0xBC}; + +uint8_t F8Code(uint64_t i) { return kF8Codes[i % (sizeof(kF8Codes) / sizeof(kF8Codes[0]))]; } + +vllm::OwnedTensor MakeF8(int64_t n_rows, int64_t n_cols, uint64_t seed) { + vllm::OwnedTensor t; + t.dtype = DType::kI8; // F8_E4M3 carried as I8, as the loader does + t.rank = 2; + t.shape[0] = n_rows; + t.shape[1] = n_cols; + std::vector b(static_cast(n_rows * n_cols)); + for (size_t i = 0; i < b.size(); ++i) b[i] = F8Code(seed + i); + t.bytes = vllm::OwnedBytes(std::move(b)); + return t; +} + +vllm::OwnedTensor MakeBf16(const std::vector& shape, float base, float step) { + vllm::OwnedTensor t; + t.dtype = DType::kBF16; + t.rank = static_cast(shape.size()); + int64_t n = 1; + for (int i = 0; i < t.rank; ++i) { + t.shape[i] = shape[static_cast(i)]; + n *= t.shape[i]; + } + std::vector b(static_cast(n) * 2); + auto* p = reinterpret_cast(b.data()); + for (int64_t i = 0; i < n; ++i) + p[i] = vt::F32ToBF16(base + step * static_cast(i % 7)); + t.bytes = vllm::OwnedBytes(std::move(b)); + return t; +} + +vllm::Gemma4MoeLayerWeights MakeMoeWeights() { + vllm::Gemma4MoeLayerWeights moe; + moe.enabled = true; + moe.top_k = kTopK; + moe.moe_intermediate = kI; + // `nk = true` [E, H]: the orientation RunGemma4Moe's VT_CHECK requires. + moe.router_proj = MakeBf16({kE, kH}, -0.25f, 0.1f); + moe.router_proj.nk = true; + + auto& ex = moe.experts; + ex.is_fp8 = true; + ex.num_experts = kE; + ex.intermediate = kI; + ex.hidden = kH; + ex.fp8.resize(static_cast(kE)); + for (int64_t e = 0; e < kE; ++e) { + auto& fex = ex.fp8[static_cast(e)]; + const uint64_t s = static_cast(e) * 131u; + fex.gate_w = MakeF8(kI, kH, s + 1); + fex.up_w = MakeF8(kI, kH, s + 2); + fex.down_w = MakeF8(kH, kI, s + 3); + fex.gate_s = MakeBf16({kI}, 0.05f, 0.01f); + fex.up_s = MakeBf16({kI}, 0.04f, 0.01f); + fex.down_s = MakeBf16({kH}, 0.03f, 0.01f); + } + return moe; +} + +// One decode token through the real layer entry point. Returns the [1, H] bf16 +// result copied out immediately: for T == 1 the layer hands back a view of a +// thread_local buffer that the NEXT call overwrites. +std::vector RunOneToken(vt::Queue& q, const vllm::Gemma4MoeLayerWeights& moe, + const std::vector& x) { + vt::Tensor xin = vt::Tensor::Contiguous(const_cast(x.data()), DType::kBF16, + q.device, {1, kH}); + vllm::Gemma4MoeScratch out = + vllm::RunGemma4Moe(q, moe, /*router_in=*/xin, /*expert_in=*/xin, /*T=*/1, kH, + /*rms_eps=*/1e-6f); + std::vector got(static_cast(kH)); + std::memcpy(got.data(), out.tensor.data, got.size() * sizeof(uint16_t)); + return got; +} + +bool AnyNonZero(const std::vector& v) { + for (const uint16_t h : v) + if (h != 0) return true; + return false; +} + +} // namespace + +TEST_CASE("gemma4 moe: the device-expert upload refuses where MatmulBTAlphaBeta has no arm") { + vt::Queue q; + q.device.type = vt::DeviceType::kCPU; + q.device.index = 0; + + // Precondition 1 — the arm really is absent for this queue. True on every + // build, HIP included: the only arm is kROCM's. + REQUIRE_FALSE(vt::HasMatmulBTAlphaBeta(q)); + + // Precondition 2 — the STOCK backend does not answer the memory probe, which + // is why the hazard is latent and why the decoration below is needed at all. + { + size_t free_b = 1, total_b = 1; + REQUIRE_FALSE(vt::GetBackend(vt::DeviceType::kCPU).DeviceMemoryInfo(&free_b, &total_b)); + } + + const vllm::Gemma4MoeLayerWeights moe = MakeMoeWeights(); + std::vector x(static_cast(kH)); + for (size_t i = 0; i < x.size(); ++i) + x[i] = vt::F32ToBF16(0.1f + 0.05f * static_cast(i % 5)); + + // Baseline: the layer as it runs today, host-fallback expert path. + const std::vector want = RunOneToken(q, moe, x); + REQUIRE(want.size() == static_cast(kH)); + // A layer that produced all zeros would make the equality below vacuous. + REQUIRE(AnyNonZero(want)); + + // Now put the tree in the post-#1126 state: the memory probe answers, so the + // LRU's `MakeRoom` will admit, and only the guard stands between the upload + // and a `vt::MatmulBTAlphaBeta` that has no arm for this device. + std::vector got; + { + vt::Backend& real = vt::GetBackend(vt::DeviceType::kCPU); + ProbingBackend probing(real); + const BackendSwap swap(probing); + + // Precondition 3 — the instrument is ARMED. Without this, a decoration that + // failed to take would leave the run on the same refuse-on-unknown path as + // the baseline and report the guard proven by a run that never reached it. + size_t free_b = 0, total_b = 0; + REQUIRE(vt::GetBackend(vt::DeviceType::kCPU).DeviceMemoryInfo(&free_b, &total_b)); + REQUIRE(free_b > (2ull << 30)); + + // THE ASSERTION. With the guard: the upload is refused, the expert loop + // takes `EnsureGemma4Fp8ExpertCached` + `ExpertGeGLUHost`, and the layer + // answers. Without it: `ExpertGeGLUTopKFusedGelu` reaches + // `vt::MatmulBTAlphaBeta` and this throws out of the decode step. + REQUIRE_NOTHROW(got = RunOneToken(q, moe, x)); + } + + // Degraded, but CORRECT: refusing the device arm must not change the answer. + REQUIRE(got.size() == want.size()); + int same = 0; + for (size_t i = 0; i < want.size(); ++i) { + if (got[i] == want[i]) ++same; + } + CHECK(same == static_cast(want.size())); + + // The refusal must not have leaked device pointers into the expert records: + // a `true` return is what assigns them, so they stay null. + int null_dev = 0; + for (const auto& fex : moe.experts.fp8) { + if (fex.dev_gu == nullptr && fex.dev_dn == nullptr) ++null_dev; + } + CHECK(null_dev == static_cast(kE)); + + // The backend registry is back to what it was. + size_t free_b = 1, total_b = 1; + CHECK_FALSE(vt::GetBackend(vt::DeviceType::kCPU).DeviceMemoryInfo(&free_b, &total_b)); +} diff --git a/tests/vt/test_gemma4_rocm_fp8_seams.cpp b/tests/vt/test_gemma4_rocm_fp8_seams.cpp index 12abaf6d2..dc698877c 100644 --- a/tests/vt/test_gemma4_rocm_fp8_seams.cpp +++ b/tests/vt/test_gemma4_rocm_fp8_seams.cpp @@ -4,8 +4,11 @@ #include #include +#include #include +#include "vt/device.h" +#include "vt/dtype.h" #include "vt/fused_ops.h" namespace { @@ -76,3 +79,121 @@ TEST_CASE("gemma4 rocm fp8 seams: recipe env knobs parse inert defaults") { CHECK(EnvInt("VT_GEMMA4_FP8_HW_CVT", 1) == 0); } } + +// #1205: `vt::MatmulBTAlphaBeta` has NO CUDA implementation. The only one in the +// tree is `rocm::MatmulBTAlphaBetaRocm` (rocm_matmul_hipblaslt.hip:516), reached +// through a `#if defined(VLLM_CPP_HIP)` + `kROCM` guard in fused_ops.cpp; every +// other device falls through to a throw. That is the actual blocker under #1126: +// adding `CudaBackend::DeviceMemoryInfo` alone would have let the Gemma4 expert +// LRU admit (gemma4_moe.cpp:587,597), take ExpertGeGLUDeviceAccum at :1509, and +// reach this call at :90 — outside the upload's try/catch (:585-607), so +// mid-decode. `EnsureGemma4Fp8ExpertOnDevice` now refuses that upload at :571 +// when `vt::HasMatmulBTAlphaBeta` is false, so the throw below is the backstop +// and not the guard; the capability gate for the refusal-before-upload is +// tests/vllm/models/test_gemma4_moe_device_arm_guard.cpp, which enters through +// `vllm::RunGemma4Moe`. What THIS file pins is the message contract a CUDA +// implementation has to satisfy — a unit contract, since with the guard in place +// no production run off ROCm reaches the throw. +// +// AGENTS.md requires an unimplemented arm to refuse with a message that NAMES the +// missing part. "ROCm-only in this build" names neither the device that asked nor +// where the implementation would go, so a caller who hits it on CUDA cannot tell a +// missing kernel from a missing build flag. This pins the contract; when the CUDA +// arm is written, this is what it has to satisfy. +TEST_CASE("gemma4 rocm fp8 seams: MatmulBTAlphaBeta refuses a non-ROCm queue by name") { + // The refusal reads `q.device.type` and nothing else, so a CPU build can pose + // as any device here. Buffers are never dereferenced on the refusing path. + auto refusal_for = [](vt::DeviceType type) -> std::string { + vt::Queue q; + q.device.type = type; + q.device.index = 0; + try { + vt::MatmulBTAlphaBeta(q, nullptr, nullptr, nullptr, /*M=*/1, /*N=*/1, /*K=*/1, + /*alpha=*/1.f, /*beta=*/0.f, vt::DType::kBF16); + } catch (const std::runtime_error& e) { + return e.what(); + } + return std::string(); + }; + + // CUDA is the device #1126 would wake, so it is the one the message must name. + const std::string cuda_msg = refusal_for(vt::DeviceType::kCUDA); + REQUIRE_FALSE(cuda_msg.empty()); + INFO("cuda refusal: " << cuda_msg); + CHECK(cuda_msg.find("MatmulBTAlphaBeta") != std::string::npos); + CHECK(cuda_msg.find("cuda") != std::string::npos); + CHECK(cuda_msg.find("rocm") != std::string::npos); + CHECK(cuda_msg.find("1205") != std::string::npos); + + // Not a CUDA special case: every non-ROCm device refuses, naming itself. + for (const auto type : {vt::DeviceType::kCPU, vt::DeviceType::kVULKAN, + vt::DeviceType::kMETAL}) { + const std::string msg = refusal_for(type); + // std::string, not the char* — doctest stringifies a bare char* as a bool. + INFO("refusal for " << std::string(vt::DeviceTypeName(type)) << ": " << msg); + REQUIRE_FALSE(msg.empty()); + CHECK(msg.find(vt::DeviceTypeName(type)) != std::string::npos); + CHECK(msg.find("1205") != std::string::npos); + } + + // kROCM is the one device the four above make the message look right FOR and + // wrong ABOUT. A kROCM queue in a build without -DVLLM_CPP_HIP falls through + // the same `#if` and refuses too — but the 'rocm' kernel EXISTS and was + // compiled out, so telling that caller "no implementation ... see #1205" would + // send them to write a kernel the tree already has. The two absences get two + // messages, and this pins the difference in both directions. + { + vt::Queue rocm_q; + rocm_q.device.type = vt::DeviceType::kROCM; + rocm_q.device.index = 0; + // Build-agnostic: on a HIP build the arm is present and the call would run a + // real GEMM over these null pointers, so ask the same predicate the dispatch + // asks rather than repeating the `#if` here. + if (!vt::HasMatmulBTAlphaBeta(rocm_q)) { + const std::string rocm_msg = refusal_for(vt::DeviceType::kROCM); + REQUIRE_FALSE(rocm_msg.empty()); + INFO("rocm refusal: " << rocm_msg); + CHECK(rocm_msg.find("MatmulBTAlphaBeta") != std::string::npos); + CHECK(rocm_msg.find("rocm") != std::string::npos); + // Names the ABSENT BUILD FLAG, which is what is actually missing here. + CHECK(rocm_msg.find("VLLM_CPP_HIP") != std::string::npos); + // And does NOT claim the kernel is unwritten, nor point at the CUDA issue. + CHECK(rocm_msg.find("1205") == std::string::npos); + CHECK(rocm_msg.find("no implementation for device") == std::string::npos); + } + } +} + +// The predicate the guard in `EnsureGemma4Fp8ExpertOnDevice` keys on. It has to +// agree with the dispatch EXACTLY — a caller that trusts `true` and then throws +// is worse than no predicate — so this asserts the agreement over every device +// rather than reading the two conditions and calling them equal. +TEST_CASE("gemma4 rocm fp8 seams: HasMatmulBTAlphaBeta agrees with the dispatch") { + int checked = 0; + for (const auto type : {vt::DeviceType::kCPU, vt::DeviceType::kCUDA, vt::DeviceType::kROCM, + vt::DeviceType::kVULKAN, vt::DeviceType::kMETAL, + vt::DeviceType::kXPU}) { + vt::Queue q; + q.device.type = type; + q.device.index = 0; + const bool has = vt::HasMatmulBTAlphaBeta(q); + INFO("device " << std::string(vt::DeviceTypeName(type)) << " has=" << has); + if (has) { + // Do NOT call through: a present arm would run a real GEMM over null + // pointers. The absence direction is the one this guard depends on. + ++checked; + continue; + } + bool threw = false; + try { + vt::MatmulBTAlphaBeta(q, nullptr, nullptr, nullptr, /*M=*/1, /*N=*/1, /*K=*/1, + /*alpha=*/1.f, /*beta=*/0.f, vt::DType::kBF16); + } catch (const std::runtime_error&) { + threw = true; + } + CHECK(threw); + ++checked; + } + // A loop that ran zero times would report SUCCESS with nothing examined. + CHECK(checked == 6); +}