feat(ENG-RESIDENCY-CONFIG): the device budget is a config key, and --offload-config reaches vllm-cli and the pooling server (#1127, #1135)

.agents/specs/expert-streaming.md

@@ -1611,7 +1611,7 @@ Carried debt for this row. Each item names why it is not closed here.
 | **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). |
+| **The budget knob is an environment variable, not a config key.** `VT_DEVICE_WEIGHT_BUDGET_BYTES`. | CLOSED. It waited for `ENG-RESIDENCY-CONFIG` ([#1110](https://github.com/mudler/vllm.cpp/issues/1110), PR #1119) to land the `vllm_cpp` namespace inside `--offload-config`, because landing a second, competing config surface while that one was unmerged would have created the conflict both changes then had to resolve. `ENG-RESIDENCY-CONFIG` W2 then added `vllm_cpp.device_fit.weight_budget_bytes`, and `DeviceWeightBudgetBytes` now resolves environment variable > config > device probe. `0` still suppresses the refusal from either input. [#1127](https://github.com/mudler/vllm.cpp/issues/1127); the key is specified in [`weight-residency-config.md`](weight-residency-config.md). |
 | **`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. |