diff --git a/.agents/issue-index.md b/.agents/issue-index.md index 5c1c465cc..f78020d76 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 | +| [#1195](https://github.com/mudler/vllm.cpp/issues/1195) | `ARCH-ONE-SURFACE` | The transcription-only server path has no seam either half of an `--offload-config` document could reach, so [#1135](https://github.com/mudler/vllm.cpp/issues/1135) REFUSED the flag there instead of wiring it. `ParakeetTranscriber::FromDir` (`src/vllm/multimodal/parakeet_transcription.cpp:49`) builds no `EngineParams` and calls no `LoadedEngine::FromModelDir`; it reads its weights through `LoadParakeetForCTC` / `LoadParakeetTransducer`, so that path has no `SetWeightResidencyConfig` call, no `CreateWeightOffloader` call, no GGUF mapping and no expert slot store — no field of either half has a reader on it. Wiring it means first giving the transcription stack a loader seam that consults the process-global offloader and the residency config, which is larger than #1135 and belongs to whoever adds one. WHAT LANDED INSTEAD: the server parses `--offload-config` ONCE, ahead of the architecture branch, so a typo is refused at startup on every path, and a NON-EMPTY document on the transcription-only path aborts at startup naming the missing seam rather than being dropped in silence. Listed under `## Owed` in [`weight-residency-config.md`](specs/weight-residency-config.md) | gap | +| [#1196](https://github.com/mudler/vllm.cpp/issues/1196) | `SERVE-POOLING-ENDPOINTS` | The server's pooling path builds `EngineParams` from 8 of the engine flags and drops the rest, `--device` included. Found while closing [#1135](https://github.com/mudler/vllm.cpp/issues/1135), which is the same shape over one flag. The `if (pooling_model)` block in `src/vllm/entrypoints/openai/server_main.cpp` sets `block_size`, `num_blocks`, `gpu_memory_utilization`, `kv_cache_memory_bytes`, `max_model_len`, `max_num_seqs`, `max_num_batched_tokens` and `enable_prefix_caching`; #1135 adds `offload_config` and `weight_residency`. STILL DROPPED: `--device` (an embedding server started with `--device cuda` runs the accelerator-first probe instead of the named device, so an explicitly named ABSENT device does not fail loudly there — the one with a user-visible consequence today), `--scheduling-policy`, `--kv-transfer-config`, `--speculative-config`, `--enable-jump-forward`/`--disable-jump-forward`, and the multimodal per-modality limits. Each is accepted on the command line, reported by nothing and honoured by nothing; the shape predates [#1110](https://github.com/mudler/vllm.cpp/issues/1110). NOT fixed in the #1135 flow: each dropped flag is a separate behaviour with its own dispatch test surface, and the pooling dispatch block is `ARCH-ONE-SURFACE ROW 6`'s surface rather than `ENG-RESIDENCY-CONFIG`'s. Closing it means one construction of `EngineParams` shared by both branches, or an explicit refusal per flag the pooling path cannot honour. Listed under `## Owed` in [`weight-residency-config.md`](specs/weight-residency-config.md) | bug | +| [#1206](https://github.com/mudler/vllm.cpp/issues/1206) | `ENG-RESIDENCY-CONFIG` | `docs/USAGE.md` invoked `./build/vllm-server` and `./build/vllm-cli` in THREE command lines (`grep -c 'build/vllm-' docs/USAGE.md` == 3 at this branch's base `fd64c76ee`, and the same at the merge base; lines 2354, 3835 and 3848), and `examples/CMakeLists.txt` builds both under `examples/`, so the binaries are at `build/examples/vllm-server` and `build/examples/vllm-cli`. The same document already used the correct spelling elsewhere (`build/examples/vllm-cli` in "Running inference (CLI)"), so a reader got two answers and one of them failed with "No such file or directory". Verified on a fresh `cmake -S . -B build -G Ninja && cmake --build build`: `ls build/vllm-server build/vllm-cli` reports no such file, and `find build -maxdepth 2` finds both under `build/examples/`. FOUND while closing [#1127](https://github.com/mudler/vllm.cpp/issues/1127) and [#1135](https://github.com/mudler/vllm.cpp/issues/1135), which add command lines to the same two sections, and FIXED IN THAT FLOW: all three spellings now name `build/examples/`, and the count is 0 at head. The row first said FIVE, which was never measured; the reviewer of [#1216](https://github.com/mudler/vllm.cpp/pull/1216) counted three and the row was corrected while that pull request was still open. The FIX was always complete — only the count was wrong. No checker covers a command line in a document, so this is a reading rather than a gate | bug | | [#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 | diff --git a/.agents/specs/expert-streaming.md b/.agents/specs/expert-streaming.md index 232114652..92e7cce03 100644 --- a/.agents/specs/expert-streaming.md +++ b/.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. | diff --git a/.agents/specs/weight-residency-config.md b/.agents/specs/weight-residency-config.md index fe3e30968..981ccea82 100644 --- a/.agents/specs/weight-residency-config.md +++ b/.agents/specs/weight-residency-config.md @@ -82,10 +82,10 @@ Three findings shape the change: | Field | Content | |---|---| | Row ID | `ENG-RESIDENCY-CONFIG` (engine-matrix, KV cache and memory). Issue [#1110](https://github.com/mudler/vllm.cpp/issues/1110); fixes [#1109](https://github.com/mudler/vllm.cpp/issues/1109) in flow | -| In | A vllm.cpp-original `WeightResidencyConfig` under the `vllm_cpp` key of the existing `--offload-config` document; its parser, which refuses an unknown key at every level of the document (the four legal top-level keys included) and a wrong-typed or non-positive field; a process-global install/resolve seam with a defined config-vs-env precedence and a late-install refusal; the three call sites that resolve these knobs today (`GgufLoadPolicy::FromEnv` for `mmap`, `PrefaultBorrowedSpan` for `prefault`, `Qwen35ExpertStreamRequested` + the `Qwen35ExpertStream` constructor for the streaming lane); the flag→`EngineParams`→install chain through both production entry points (`server_main.cpp` and the C ABI's `offload_config`); `docs/USAGE.md` and `docs/ENVIRONMENT.md` | -| Out | Any change to `OffloadConfig`, `UVAOffloadConfig`, `PrefetchOffloadConfig` or their validator — the mirror stays byte-faithful. Any change to what the knobs *do*: this row moves where their value comes from and nothing else. A new flag. `VT_MOE_EXPERT_STREAM_STATS_EVERY` (see below). `VT_GGUF_KEEP_QUANT`, `VT_CPU_REF`, `VT_GGUF_KEEP_F16` and the rest of the load-transform family — they are a different tier and a different row | -| Supported modes | `{"vllm_cpp":{"mmap":{"enabled":bool,"prefault":bool},"expert_stream":{"enabled":bool,"slots":int,"slot_bytes":int}}}`. Every field is optional and every absent field means "unchanged", so an absent `vllm_cpp` key is byte-identical to today | -| Dispatch behavior | Resolved from **env var if set, else config if set, else the built-in default**, at each read. `mmap` and `prefault` are read per load and per span; `expert_stream` is cached on first read and the two sizes are fixed when the slot store is built. Nothing is resolved when neither input is set, so the default engine path is byte-identical | +| In | W2 adds the sixth knob and the three unreached entry points; W1 is everything below it. A vllm.cpp-original `WeightResidencyConfig` under the `vllm_cpp` key of the existing `--offload-config` document; its parser, which refuses an unknown key at every level of the document (the four legal top-level keys included) and a wrong-typed or non-positive field; a process-global install/resolve seam with a defined config-vs-env precedence and a late-install refusal; the three call sites that resolve these knobs today (`GgufLoadPolicy::FromEnv` for `mmap`, `PrefaultBorrowedSpan` for `prefault`, `Qwen35ExpertStreamRequested` + the `Qwen35ExpertStream` constructor for the streaming lane); the flag→`EngineParams`→install chain through both production entry points (`server_main.cpp` and the C ABI's `offload_config`); `docs/USAGE.md` and `docs/ENVIRONMENT.md`. **W2** ([#1127](https://github.com/mudler/vllm.cpp/issues/1127), [#1135](https://github.com/mudler/vllm.cpp/issues/1135)): a sixth key, `vllm_cpp.device_fit.weight_budget_bytes`, which takes over `VT_DEVICE_WEIGHT_BUDGET_BYTES`'s middle tier at `DeviceWeightBudgetBytes`; and the `--offload-config` document reaching `vllm-cli` and the server's pooling path, with the transcription-only path refusing the flag by name instead of dropping it | +| Out | Any change to `OffloadConfig`, `UVAOffloadConfig`, `PrefetchOffloadConfig` or their validator — the mirror stays byte-faithful. Any change to what the knobs *do*: this row moves where their value comes from and nothing else. A new flag NAME: W2 gives `vllm-cli` the `--offload-config` spelling the server already uses, and adds no second flag for either tier. `VT_MOE_EXPERT_STREAM_STATS_EVERY` (see below). Honouring a weight-residency document on the transcription-only server path, which has no loader to honour it with (see `## Owed`). The six other `EngineParams` fields the pooling path also drops (see `## Owed`). `VT_GGUF_KEEP_QUANT`, `VT_CPU_REF`, `VT_GGUF_KEEP_F16` and the rest of the load-transform family — they are a different tier and a different row | +| Supported modes | `{"vllm_cpp":{"mmap":{"enabled":bool,"prefault":bool},"expert_stream":{"enabled":bool,"slots":int,"slot_bytes":int},"device_fit":{"weight_budget_bytes":int}}}`. Every field is optional and every absent field means "unchanged", so an absent `vllm_cpp` key is byte-identical to today. `slots` and `slot_bytes` must be positive; `weight_budget_bytes` must be non-negative, because `0` is the documented spelling of "suppress the device-fit refusal" | +| Dispatch behavior | Resolved from **env var if set, else config if set, else the built-in default**, at each read. `mmap` and `prefault` are read per load and per span; `expert_stream` is cached on first read and the two sizes are fixed when the slot store is built; `weight_budget_bytes` is read once per GGUF load, at the fit check, and caches nothing. Nothing is resolved when neither input is set, so the default engine path is byte-identical | | Regimes served | A checkpoint larger than host RAM on a single box: the mmap-borrowed weight tower plus the bounded expert slot cache. CPU keep-quant expert towers today; a device platform serves the slice device-resident and is unaffected | ## Upstream chain @@ -157,6 +157,10 @@ resolve sites have no input but `getenv`. | `expert_stream`, `slots`, `slot_bytes` resolve | `src/vllm/model_executor/models/qwen3_5.cpp`, `Qwen35ExpertStreamRequested` and the `Qwen35ExpertStream` constructor | | Docs | `docs/USAGE.md` (the streaming section gains the config form), `docs/ENVIRONMENT.md` (precedence note + the #1109 default correction) | | Named resolvers | `ResolveGgufMmap`, `ResolveGgufPrefault`, `ResolveExpertStreamRequested` (+ its pure `ExpertStreamRequestedFrom`), `ResolveExpertStreamSlots`, `ResolveExpertStreamSlotBytes` — one per knob, each the sole reader of its variable | +| W2: the budget key | `device_weight_budget_bytes` in `WeightResidencyConfig`, parsed from `vllm_cpp.device_fit.weight_budget_bytes`, resolved by `ResolveDeviceWeightBudgetBytes` in the same pair | +| W2: the budget call site | `DeviceWeightBudgetBytes` in `src/vllm/model_executor/model_loader/gguf_device_fit.cpp`, which keeps its name and its digits-only environment grammar and delegates the precedence to the resolver | +| W2: the server flag, once | `src/vllm/entrypoints/openai/server_main.cpp` — the parse moves ahead of the architecture branch; the pooling path takes both halves; the transcription-only path refuses a non-empty document | +| W2: `vllm-cli` | `examples/cli/main.cpp` — a `--offload-config` flag assigned to `vllm_model_params.offload_config`, which `src/capi/vllm_c.cpp` already parses for both halves | **Five knobs, five named resolvers, and the polarities are not the same.** Each knob gets one function in the new header that owns its environment NAME and its @@ -264,6 +268,35 @@ New tests, each red before its implementation: stderr for a `--offload-config` document carrying only a `vllm_cpp` key. This is the test the reachability mutation deletes the call site under. +W2 adds these, each red before its implementation: + +- `tests/vllm/config/test_weight_residency_config.cpp` gains the budget key: the + parse of `vllm_cpp.device_fit.weight_budget_bytes`, `0` ACCEPTED (it is the + documented spelling of "suppress the refusal", so the sibling rule that refuses + a non-positive value would be wrong here), a negative value refused, a + misspelled `device_fit` and a misspelled `weight_budget_bytes` each refused by + name, the three precedence cases, "absent means unchanged" across two partial + documents, and the absence of a latch — a late install of the budget after a + load is accepted. +- `tests/vllm/entrypoints/test_gguf_device_fit_reach.cpp` gains the budget key's + REACHABILITY: the same synthetic GGUF and fake staging platform the environment + cases use, with NO variable set and the budget arriving through + `EngineParams::weight_residency`, asserting the refusal message and its byte + count. It is the case the reachability mutation deletes the install call site + under, and it also carries the environment-beats-config direction through the + loader. +- `tests/vllm/entrypoints/openai/test_serve_residency_config.cpp` gains the two + server entry points: a model directory whose `config.json` names `LlamaModel` + takes the pooling branch AND prints the install line, and one that names + `ParakeetForCTC` takes the transcription branch and ABORTS when + `--offload-config` is non-empty while starting normally when it is absent. +- `tests/vllm/entrypoints/test_cli_offload_config.cpp` — `vllm-cli`'s own + reachability, by running the built `vllm-cli` binary on a nonexistent model + with `--offload-config` and reading the install line off its stderr. A test + that called the C ABI directly would prove the ABI, which + `test_weight_residency_reach.cpp` already does; only running the binary proves + the FLAG arrives. + **Two existing suites gain the observable their knob never had.** Both were found by the mutation pass, not by reading, and both are gaps that predate this row — the old inline `getenv` calls were equally unwatched: @@ -307,6 +340,15 @@ claims none. ./build/tests/test_gguf_keep_quant ./build/tests/test_expert_stream_mixed_slot ``` + + W2 adds four more to that list: + + ```sh + ./build/tests/test_gguf_device_fit + ./build/tests/test_gguf_device_fit_reach + ./build/tests/test_cli_offload_config + ./build/tests/test_offload_config + ``` 3. Every guarantee mutation-proven: for each added test, delete or invert the behavior it names, rebuild, require its suite red with a non-zero case count, restore by byte copy and verify by sha256. A mutation that fails to compile is @@ -338,8 +380,8 @@ claims none. ## Work breakdown -One wave. The change is one field wide at every hop, and splitting it would land -a parser nothing reaches. +W1 was one wave. The change is one field wide at every hop, and splitting it +would have landed a parser nothing reaches. | Step | Content | |---|---| @@ -348,8 +390,60 @@ a parser nothing reaches. | W1c | `EngineParams`, `server_main.cpp`, the C ABI, the install site, and both reachability suites | | W1d | `docs/USAGE.md`, `docs/ENVIRONMENT.md` (including the #1109 correction), the records | +W2 closes two of W1's own `## Owed` entries in one change, because they edit the +same two files. Splitting them would put two branches on `server_main.cpp`'s +argument block and `weight_residency.cpp`'s parser at the same time. + +| Step | Content | +|---|---| +| W2a | `device_fit.weight_budget_bytes` ([#1127](https://github.com/mudler/vllm.cpp/issues/1127)): the parser key, the merge, the announce line, `ResolveDeviceWeightBudgetBytes`, and `DeviceWeightBudgetBytes` delegating to it | +| W2b | The three entry points ([#1135](https://github.com/mudler/vllm.cpp/issues/1135)): the server's offload parse hoisted ahead of the architecture branch, the pooling path taking both halves, the transcription-only path refusing the flag, and `vllm-cli` gaining `--offload-config` | +| W2c | `docs/USAGE.md`, `docs/ENVIRONMENT.md`, and this spec. Also [#1206](https://github.com/mudler/vllm.cpp/issues/1206) in flow: THREE `docs/USAGE.md` command lines named `./build/vllm-server` or `./build/vllm-cli` (lines 2354, 3835, 3848 at the branch base `fd64c76ee`; `grep -c 'build/vllm-' docs/USAGE.md` is 3 there and at the merge base, and 0 at head), and both binaries are built under `build/examples/`. W2 first wrote FIVE here and in three other places; the number was never counted, and the review of PR #1216 corrected it | + ## Risks and decisions +**W2 decisions.** + +- **The budget key is `vllm_cpp.device_fit.weight_budget_bytes`, not a scalar + beside `mmap` and `expert_stream`.** Two reasons. `vllm_cpp` maps a name to an + OBJECT today, and the parser's `ExtObject` walk depends on that; a bare scalar + sibling makes the level heterogeneous for one field. And the budget is a third + KNOB FAMILY — the load-time device-fit check — rather than a member of either + existing one, so it gets its own object exactly as they do. The field name drops + the family prefix, which is the transformation `VT_GGUF_MMAP` -> `mmap.enabled` + and `VT_MOE_EXPERT_STREAM_SLOTS` -> `expert_stream.slots` already use. +- **`0` is legal for this key and illegal for its two integer siblings.** `slots` + and `slot_bytes` are sizes, and a zero size that silently became 64 is a cache + the operator does not have, so `ExtPositiveInt` refuses it. `0` on the budget is + the DOCUMENTED spelling of "suppress the device-fit refusal", identical to + `VT_DEVICE_WEIGHT_BUDGET_BYTES=0`, and `CheckDeviceWeightFit` reads a zero budget + as UNKNOWN and decides nothing. Refusing it would remove the escape hatch this + key exists to give. So the budget parses through `ExtNonNegativeInt` and a + negative value is still refused. +- **The environment grammar does not change, and the config sits UNDER it.** + `DeviceWeightBudgetBytes` accepts decimal digits only: a sign, a space or + trailing garbage is ignored and the next tier stands. That tier used to be the + probe and is now the config, so a run with no config resolves byte-for-byte as + before. The rule is transcribed once, in `ResolveDeviceWeightBudgetBytes`, and + `DescribeEnvOverrides` asks the SAME predicate rather than reporting presence — + the #1122 L7 shape, where `SLOTS=banana` was announced as an override the + resolver ignored. +- **The budget latches nothing.** It is read once per GGUF load, at the fit check, + through no static. A late install that sets it is therefore accepted, and + `ResidencyLatch` gains no enumerator. That is pinned rather than asserted: a case + installs a budget after a completed load and requires no throw. +- **The transcription-only path refuses the flag rather than accepting it.** + AGENTS.md: refuse an unimplemented arm with a message that names the missing + part, and record the arm as owed. A warning would leave a server running while + it holds a placement instruction it does not follow, which is the shape #1135 + was filed about. The cost is a launcher script that passes one flag to every + model and now has to stop passing it to an ASR model. See `## Owed` and #1195. +- **The server parses `--offload-config` ONCE, ahead of the architecture branch.** + Three consumers over one parse, rather than three parses. It also moves the + typo refusal earlier than the `server: loading model from` line, so a mistyped + document now aborts before that line prints instead of after it. + + - **A namespaced key inside a mirrored flag can read as "vLLM takes this".** It does not. The key is literally `vllm_cpp`, which is the cheapest possible signal that the contents are not upstream, and the docs say so in the same @@ -424,6 +518,262 @@ a parser nothing reaches. ## Evidence +### W2 (#1127, #1135) + +CPU host, documented recipe (`cmake -S . -B build -G Ninja`, no build type, so +asserts are live), 20 cores, shared with three concurrent agent builds. Every +build below reports its own exit status and an ENOSPC count, because +`ld: final link failed: No space left on device` reads exactly like broken code +and free disk moved between 42 GB and 7.3 GB during the run. **No build in this +wave hit ENOSPC: the count is 0 on every one.** + +**Red first, for the budget key.** The field and `ResolveDeviceWeightBudgetBytes` +were declared and the resolver stubbed to `return probed_total_bytes;`, then the +suite was built (rc 0) and run: `test_weight_residency_config` rc **1**, 24 cases +/ 17 passed / **7 failed**, 271 assertions / 15 failed. The seven are exactly the +seven cases W2 adds. With the real implementation: rc **0**, 24 cases / 24 passed, +315 assertions. + +**Red first, for the three entry points, is the M8/M10/M11 row of the mutation +table.** Each of those deletes the assignment the wave adds, which is the state of +the tree before the wave, and each turns its suite red. + +**The three `test_gguf_device_fit_reach` cases were written AFTER the resolver, so +their red is a mutation rather than a pre-implementation run — and it takes THREE +mutations, not one.** This first said "their red is M7", which is true of ONE case +out of the three. M7 restores `DeviceWeightBudgetBytes`'s pre-#1127 body — the +environment grammar and the probe, with no config tier — and it kills the case that +carries the budget in as a config key with no variable set. The other two are pinned +elsewhere. Measured in the review repair, whole binaries, the failing case names read +off the doctest output rather than inferred: + +| case | killed by | and NOT by, because | +|---|---|---| +| the budget arrives as a CONFIG KEY, with no variable set | M7, M12 | M4 — the case sets no variable, so inverting the precedence changes nothing | +| a config budget of ZERO suppresses the refusal | M12 alone | M7 — it touches neither the install nor the read-back; M4 — there is no variable for the config to beat | +| the VARIABLE beats the config key, through the loader | M4 alone | M7 and M12 — both leave the variable in force, and the variable winning is what this case asserts | + +**And the ZERO case's "no refusal" half is VACUOUS on that fixture.** The fake +staging platform probes 0 and `CheckDeviceWeightFit` reads a zero budget as +UNKNOWN, so no refusal fires whether the configured zero arrives or not. That half +distinguishes nothing there. The case's one discriminating assertion is the +read-back `installed.device_weight_budget_bytes == 0`, which is exactly why M12 — +the mutation that removes the install — is the only one of the three that kills it. +The case comment in the file says this; the spec did not, and a reader of the spec +alone would have counted a guarantee the fixture cannot give. The positive control +for the refusal direction is the CONFIG KEY case above it, on the same file and the +same platform. + +**Mutations.** Twelve, each applied alone from a pristine byte copy, with the +file's sha256 printed before and after so a never-applied edit cannot read as a +pass, the build's exit status printed beside every result so a non-building +mutation is INVALID rather than a pass, a NON-ZERO doctest case count required, the +LAST `test cases:` line taken, and the tree restored by byte copy with the sha256 +compared against the pre-mutation value. All twelve built (rc 0) and all twelve +were killed. + +| id | what it breaks | suite | rc | cases | +|---|---|---|---|---| +| M1 | `device_fit` drops out of the `vllm_cpp` enumeration | `test_weight_residency_config` | 1 | 24 | +| M2 | the parsed budget is never stored on the config | `test_weight_residency_config` | 1 | 24 | +| M3 | `ExtNonNegativeInt` refuses zero, losing the suppression spelling | `test_weight_residency_config` | 1 | 24 | +| M4 | the config beats the environment, inverting the precedence | `test_weight_residency_config`, `test_gguf_device_fit_reach` | 1, 1 | 24, 8 | +| M5 | the merge assigns wholesale, so an absent budget CLEARS the installed one | `test_weight_residency_config` | 1 | 24 | +| M6 | the override note reports PRESENCE, announcing a value the resolver ignores | `test_weight_residency_config` | 1 | 24 | +| M7 | **the #1127 call site**: `DeviceWeightBudgetBytes` reverts to environment-only | `test_gguf_device_fit_reach` | 1 | 8 | +| M8 | **the #1135 pooling call site**: `embed_params` drops both halves again | `test_serve_residency_config` | 1 | 11 | +| M9 | the transcription-only refusal never fires | `test_serve_residency_config` | 1 | 11 | +| M10 | **the #1135 `vllm-cli` call site**: the flag is parsed and never passed on | `test_cli_offload_config` | 1 | 4 | +| M11 | the TEXT path drops the residency half of the hoisted parse | `test_serve_residency_config` | 1 | 11 | +| M12 | **the reachability mutation**: `SetWeightResidencyConfig` is never called | `test_serve_residency_config`, `test_cli_offload_config`, `test_gguf_device_fit_reach` | 1, 1, 1 | 11, 4, 8 | + +**Focused suites, green, with case counts:** `test_weight_residency_config` 24, +`test_gguf_device_fit` 8, `test_gguf_device_fit_reach` 8, +`test_serve_residency_config` 11, `test_cli_offload_config` 4 — each rc 0. + +**Full gate.** `cmake --build build -j 6` rc 0, ENOSPC 0. +`ctest --test-dir build -j 6` rc 0: **100% tests passed, 0 failed out of 516**, +735.62 s, with `test_modelopt_mixed_precision_checkpoint` and `test_voxtral_e2e` +skipped as they are on this host. + +**A doctest filter trap, recorded because it cost a reading.** Two W2 case names +contain a comma, and `--test-case=` splits its argument on commas, so filtering to +one of them ran 0 cases and printed `Status: SUCCESS!`. Every result above comes +from running a whole binary and reading its case count, never from a filtered run. + +**Not measured, and not claimed.** This wave moves where a value comes from and +which entry points carry a document. It changes no kernel, no dtype and no +allocation, so it has no throughput axis, and the 370 GiB reproduction stays owed +above. + +### W2 review repair (PR #1216) + +Same host and same recipe: `cmake -S . -B build -G Ninja`, no build type so asserts +are live, `-j 6`, 20 cores, shared with several concurrent agent builds. Free disk +moved between 30 GB and 17 GB during the run, so every build below reports an ENOSPC +count beside its exit status; **every one is 0**. `ctest -N` reports 523 +tests on this branch, not the 516 W2 recorded, because `origin/main` added test +binaries between the two runs. + +**Two results that were previously reported as one run, separated.** The pull +request body said "100% tests passed, 0 failed out of 516" and, two paragraphs +later, an A/B with "3 failures against this branch's 2". Those are different runs, +different instruments, different times and different lanes. +`test_cpu_x86_llamacpp_floor` is not a ctest test at all: `ctest -N | grep -i floor` +matches nothing. It is `tests/scripts/test_cpu_x86_llamacpp_floor.py`, a Python +harness contract run by `ci.yml:210` and `scripts/agent-preflight.sh:145`, and it can +never appear in a ctest count. + +**The concurrent A/B measured the A/B, not either tree.** Running this branch and a +clean `origin/main` worktree at the same moment makes each arm the other's +contention, and the failure lands on whichever loses the race. The review's own rerun +of that A/B put the failure on the CONTROL arm, inverting the 3-versus-2 recorded +here. Both readings are artefacts of the method and neither is evidence. **The +3-versus-2 claim is withdrawn.** + +**What does isolate it, because neither half races.** First, byte identity of the +three files the harness reads, on this branch and on `origin/main`: + +| file | blob at `aed3aa5e9` | blob at `origin/main` `65d6cdaed` | +|---|---|---| +| `tests/scripts/test_cpu_x86_llamacpp_floor.py` | `358927d17` | `358927d17` | +| `scripts/cpu-x86-llamacpp-floor.sh` | `bd7a51925` | `bd7a51925` | +| `docs/bench-evidence/cpu-x86-llamacpp-20260811.md` | `94f8509c4` | `94f8509c4` | + +A branch that changes none of the three cannot be why the harness reads them +differently. Second, a SINGLE-arm sweep on this branch alone, one run at a time: + +| 1-minute load average | rc | what fired | +|---|---|---| +| 172.06 before, 194.47 after | 1 | the harness's own quiet gate: `waiting for quiet: 15s busy=920% builders=0 load=194.47`, then `NO_QUIET_WINDOW after 15s`. 2 of its 10 unittest cases failed and the run took 114.7 s | +| 92.56 before, 107.42 after | 1 | the same gate, later in the session: `NO_QUIET_WINDOW after 30s (busy=104% builders=0 load=103.79 ...)`. 1 of 10 failed, 86.0 s | +| 53.41 before, 49.45 after | **0** | nothing: `Ran 10 tests in 3.338s` / `OK` — the same three files, the same tree, 34x faster | + +Three single-arm runs on ONE tree, and the verdict tracks the box's load rather +than anything about the branch. `scripts/agent-preflight.sh` is rc 1 exactly when +this harness is rc 1, and rc 0 when it is rc 0. + +The assertion that fires is the harness's own quiet gate, not a figure it computed, +which is the signature `.agents/environment.md` already records for this box. + +**A gate that claimed to SKIP actually FAILED, and the mutation that proves the +repair.** `CMakeLists.txt` said `test_cli_offload_config` "is SKIPPED when examples +are not built". With `VLLM_CLI_BINARY` undefined the file defined it as `""` and all +four cases ran `REQUIRE(std::string(kCliBinary) != "")`, which is a failure. The +suite now exits 77 — the `SKIP_RETURN_CODE` `vllm_cpp_add_test` already sets on every +test — and CTest reports **Skipped**. `cmake -DVLLM_CPP_BUILD_TESTS=ON +-DVLLM_CPP_BUILD_EXAMPLES=OFF` throughout the first two rows: + +| | build rc | run rc | doctest / ctest | +|---|---|---|---| +| repaired | 0 | 77 | `ctest -R '^test_cli_offload_config$'` reports `***Skipped` and itself exits 0 | +| **mutated** back to the PR-head file | 0 | 1 | `test cases: 4` / `0 passed` / `4 failed` / `0 skipped` | +| repaired, examples ON (the declared recipe) | 0 | 0 | `test cases: 4` / `4 passed` / `0 failed`, 31 assertions | + +The mutation is the whole pre-repair file taken from the PR head, because a PARTIAL +revert — swapping only the four `RequireCliBinary()` calls back to `REQUIRE` — does +**not compile**: `SkipGate` then trips `-Werror=unused-function`, build rc 1. That +attempt is recorded rather than dropped, because a mutation that fails to build is +INVALID and reads exactly like a passing test if only the suite's exit code is +printed. It also means the repaired shape cannot be half-reverted by accident. + +**`include/vllm.h`.** The `vllm_cpp` block documented the pre-#1127 schema in the one +file AGENTS.md names for a shipped capability, and `vllm-cli` — this wave's own new +entry point and a pure ABI client — passes exactly the document that block called +invalid. Measured: `git grep device_fit aed3aa5e9 -- include/vllm.h src/capi/` matches +nothing, while `include/vllm/config/weight_residency.h` has seven hits at the same +commit, so the key existed everywhere except the one file that is the public surface. +Three statements were wrong and all three are fixed: the schema omitted `device_fit`, +the precedence list omitted `VT_DEVICE_WEIGHT_BUDGET_BYTES`, and the refusal list +never said that `weight_budget_bytes` refuses a negative while accepting `0`. + +**`docs/USAGE.md` quoted a refusal message the parser had stopped printing.** W2 added +`device_fit` to the `vllm_cpp` level, so `RejectUnknownKeys` enumerates three names; +the sample output still read `(expected one of: mmap expert_stream)`, in the section +whose subject is that a typo is named rather than ignored. Found by +`scripts/check-doc-checkpoint.py` refusing the repair commit for touching +`include/vllm.h` and `CMakeLists.txt` without touching `docs/USAGE.md`. The gate was +owed a real edit and this is it. No separate issue: the defect was introduced by this +unlanded pull request and is repaired inside it. + +**The setter now refuses what the parser refuses.** `SetWeightResidencyConfig` is +declared in a public header and takes the struct, so a hand-built config reaches the +process-global having run no parser, and `ResolveDeviceWeightBudgetBytes` casts its +`int64_t` to `size_t`: an installed `-1` resolved to `SIZE_MAX`, an effectively +infinite budget that switched the load-time device-fit refusal off in silence — the +failure that refusal's own text warns about. The resolver's comment justified the +cast by trusting a parser this door does not run, so the trust is made true rather +than narrated. Red first, whole binary, `test_weight_residency_config`: rc +1, 25 cases / 1 failed, 324 assertions / 7 failed, the failing +message logged as `ACCEPTED (no throw)`. With the guard: rc 0, +25 cases, 324 assertions / 324 passed. The suite's case count moves 24 → 25, which +is the check that the case was added rather than silently skipped. `0` still installs, +because it is this field's suppression spelling; `slots` and `slot_bytes` get the same +rule at the same door, so the two doors into one struct state one thing. + +**Counts.** #1206 said five `docs/USAGE.md` command lines and it is three +(`grep -c 'build/vllm-' docs/USAGE.md` is 3 at the branch base `fd64c76ee`, 3 at the +merge base `5af6e763`, at lines 2354, 3835 and 3848, and 0 at head); corrected in the +index row, in `## Work breakdown` W2c, in the pull request body and on the issue, +while the pull request is still open, because an append-only index row cannot be +corrected after it. The sixth knob also invalidated six "five" statements the wave did +not repair — `weight_residency.h:62`, `:261`, `:312`, `:315`, +`weight_residency.cpp:618` and `test_serve_residency_config.cpp:58` — and left "These +two" above three cases in `test_gguf_device_fit_reach.cpp:321`. +`weight_residency.h:92` is historical and `weight_residency.cpp:325` is relative to +the budget's own siblings; both are still true and both are untouched. + +**The repair's own mutations**, run with the same harness the wave used, plus one +correction to it: the restore is `cp` + `touch` + a rebuild, never `cp -p`, because +`cp -p` preserves the mtime and ninja then skips the rebuild — a stale mutated binary +then produces a kill it has not earned. The baseline was re-run green between +mutations, and every restore was verified byte-identical by sha256. + +| id | what it breaks | applied | build rc | suite | rc | cases | cases it killed | +|---|---|---|---|---|---|---|---| +| M7 | `DeviceWeightBudgetBytes` reverts to environment-only | YES | 0 | `test_gguf_device_fit_reach` | 1 | 8 | the CONFIG KEY case, alone | +| M4 | the config beats the environment | YES | 0 | `test_gguf_device_fit_reach`, `test_weight_residency_config` | 1, 1 | 8, 25 | the VARIABLE-beats-config case, alone; and `the budget resolves env > config > probed total` | +| M12 | the install call site is deleted | YES | 0 | `test_gguf_device_fit_reach`, `test_cli_offload_config`, `test_serve_residency_config` | 1, 1, 1 | 8, 4, 11 | the CONFIG KEY case and the ZERO case; 1 of 4; 4 of 11 | +| F2 | only the four `RequireCliBinary()` calls reverted | YES | **1** | — | — | — | **INVALID: did not build** (`-Werror=unused-function` on `SkipGate`) | +| F2b | the whole pre-repair `test_cli_offload_config.cpp` | YES | 0 | `test_cli_offload_config` (EXAMPLES=OFF) | 1 | 4 | all 4 | + +**Focused suites after the repair, whole binaries, each with its case count:** +`test_weight_residency_config` 25 cases / 324 assertions, +`test_gguf_device_fit` 8 / 61, `test_gguf_device_fit_reach` 8 / 38, +`test_serve_residency_config` 11 / 126, `test_cli_offload_config` 4 / 31, +`test_weight_residency_reach` 7 / 76, `test_expert_stream_latch` 1 / 9 — each rc 0. + +**Full gate, twice, because `origin/main` moved between them.** Both runs are the +declared recipe on the same host. + +| tree | build rc | ENOSPC | `ctest --test-dir build -j 6` | +|---|---|---|---| +| the repair on `65d6cdaed` | 0 | 0 | rc 0, **100% tests passed, 0 failed out of 523**, 466.09 s | +| the repair merged onto `c20018f8d` | 0 | 0 | rc 0, **100% tests passed, 0 failed out of 523**, 587.09 s | +| the same, merged onto `727163997` | 0 (`ninja: no work to do`) | 0 | rc 0, **100% tests passed, 0 failed out of 523**, 2218.79 s | + +`test_modelopt_mixed_precision_checkpoint` and `test_voxtral_e2e` are Skipped in +all three, as they are on this host. **The LAST run is the one that counts**: a +clean merge is not a merge that builds, so the gate is rerun on each merged tree +rather than inherited. The third run's build reported `ninja: no work to do`, +because `727163997` changes only records and checkers, so its binaries are the +second run's binaries; it was rerun anyway rather than argued. + +The three wall times differ for contention only, and the difference is the box +rather than the tree: the 1-minute load average ranged from about 40 to above 220 +across them, and `test_ltx2_video` alone took 461.46 s, 542.97 s and 2147.17 s — +three worktrees on this host were running that same test at once. + +`scripts/check-symbol-anchors.py`, which `origin/main` added in the interval, +reports OK: 618 citations, 93 in-repo checked, 93 fresh, 0 stale. +`scripts/agent-preflight.sh` is rc 0 with no failing gate. It was rc 1 before the +last merge, on `test_check_gate_commands`; the spec that assertion reads is +byte-identical on this branch and on `origin/main` (`eebe438b1` both sides), so it +was measuring `origin/main`, and `origin/main`'s own `FIX-GATE-COMMANDS-PROSE-PIN` +(#1229) is the repair. Merging it was the whole fix. + +### W1 (#1110, #1109, #1122, #1133) + CPU host, documented recipe (`cmake -S . -B build -G Ninja`, no build type, so asserts are live), 20 cores. Executables are linked with `-Wl,-s`; that strips symbol tables only, and it was needed because the box had 8 GB free at the time. @@ -702,23 +1052,35 @@ count from the last `test cases:` match. `mincore()` over a host mirror, and the same instrument would apply here. Owned by `ENG-RESIDENCY-CONFIG`, issue [#1110](https://github.com/mudler/vllm.cpp/issues/1110). -- **The config form does not reach three entry points, and two of them are - server-side.** `--offload-config` is parsed once, after the architecture - resolution, so the server's POOLING/embedding path - (`server_main.cpp`, the `if (pooling_model)` block) and its transcription-only - path build their `EngineParams` without it — the MIRRORED `uva`/`prefetch` half is - dropped there too, and has been since before this key existed, so this is a - pre-existing gap that the new key inherits rather than a regression. `vllm-cli` - has no such flag at all, which is a deliberate scope line (this row adds no new - flag) but is unrecorded. Both are documented in `docs/USAGE.md` beside the config - form so a reader is not left to discover it. Fixing the server half means moving - the offload parse ahead of the architecture branch, which is - `ENG-WEIGHT-OFFLOAD`'s surface as much as this one's. Owned by - `ENG-RESIDENCY-CONFIG`, issue - [#1135](https://github.com/mudler/vllm.cpp/issues/1135) — filed for this gap - specifically, because it was previously listed against - [#1122](https://github.com/mudler/vllm.cpp/issues/1122), the review issue this pull - request closes, so on landing the gap would have had no open issue (#1133 L8). +- **The transcription-only server path cannot honour a weight-residency + document, so W2 refuses the flag there instead of wiring it.** + `ParakeetTranscriber::FromDir` (`src/vllm/multimodal/parakeet_transcription.cpp:49`) + builds no `EngineParams` and calls no `LoadedEngine`. It reads its weights + through `LoadParakeetForCTC` and `LoadParakeetTransducer`, so that path has no + `SetWeightResidencyConfig` call, no `CreateWeightOffloader` call, no GGUF + mapping and no expert slot store. There is no field of either half of the + document that any code on it could read. Wiring it means giving the + transcription stack a loader seam, which is a bigger change than #1135 and + belongs to whoever adds one. What W2 lands instead: the parse still runs, so a + typo is refused at startup exactly as on every other path, and a non-empty + `--offload-config` on that path aborts with a message naming the missing seam + rather than being dropped in silence. Owned by + `ARCH-ONE-SURFACE`, issue + [#1195](https://github.com/mudler/vllm.cpp/issues/1195). +- **The pooling path drops six of the server's `EngineParams` fields.** W2 adds + `offload_config` and `weight_residency` to that block, because those two are + #1135's subject; it built eight fields before and builds ten now. The same + block still drops `device`, `policy`, `kv_transfer_config`, + `speculative_config`, `enable_jump_forward` and `multimodal`, so the flags + behind them — `--device`, + `--scheduling-policy`, `--kv-transfer-config`, `--speculative-config`, + `--enable-jump-forward` / `--disable-jump-forward` and the multimodal limits — + are accepted and honoured by nothing. `vllm-server --device cuda` on an + embedding model therefore runs the accelerator-first probe instead of the + named device. That is the same shape as #1135 over a wider set of flags, it + predates this row, and the pooling dispatch block is `ARCH-ONE-SURFACE ROW 6`'s + surface rather than this row's. Owned by `SERVE-POOLING-ENDPOINTS`, issue + [#1196](https://github.com/mudler/vllm.cpp/issues/1196). ## Now @@ -745,7 +1107,23 @@ enumeration reaches the mirrored `uva`/`prefetch` sub-objects, which makes the A "anywhere in the document" true. `#1135` now owns the unreached-entry-point gap that was listed against the closing review issue. +W2 closes #1127 and #1135. The document carries a sixth knob, +`vllm_cpp.device_fit.weight_budget_bytes`, so the escape hatch for the load-time +device-fit refusal is a config key on the same terms as its five siblings: +`VT_DEVICE_WEIGHT_BUDGET_BYTES` still wins over it, `0` still suppresses the +refusal, and a misspelling anywhere in the new level is refused by name. It is +the only key here whose legal range includes zero, and the parser has a separate +helper for that. + +The document also reaches two more entry points. `vllm-cli` takes +`--offload-config` and passes it to the C ABI field that already parses both +halves. The server parses the flag ONCE, ahead of the architecture branch, and +the pooling path now takes both halves. The transcription-only path takes +neither, because it has no loader seam that could read either — it REFUSES a +non-empty document at startup and names what is missing, which is #1195. + What keeps it `ACTIVE` rather than `DONE` is what is above under `## Owed`: nobody has yet driven the 370 GiB checkpoint through the JSON form on the box that can -hold it, and the config form still does not reach the pooling, transcription or -`vllm-cli` entry points. +hold it, the transcription-only path refuses the document rather than honouring +it (#1195), and the pooling path still drops the server's other engine flags +(#1196). diff --git a/CMakeLists.txt b/CMakeLists.txt index f425fb5f1..f2a8b82cc 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -2383,3 +2383,17 @@ endif() if(VLLM_CPP_BUILD_EXAMPLES) add_subdirectory(examples) endif() + +# `vllm-cli`'s own reachability gate (#1135) runs the BINARY, so it needs the +# binary built first. The dependency is declared HERE rather than in +# tests/CMakeLists.txt because `add_subdirectory(tests)` runs before +# `add_subdirectory(examples)`, so the `vllm-cli` target does not exist yet at +# that point. The `$` compile definition in tests/ is fine +# there, because a generator expression is evaluated after every directory has +# been processed. With examples off there is no `vllm-cli` target, no +# `VLLM_CLI_BINARY` define, and the test EXITS 77 so CTest reports it Skipped +# (see tests/vllm/entrypoints/test_cli_offload_config.cpp); this guard is the only +# thing that has to know about the ordering. +if(TARGET test_cli_offload_config AND TARGET vllm-cli) + add_dependencies(test_cli_offload_config vllm-cli) +endif() diff --git a/docs/ENVIRONMENT.md b/docs/ENVIRONMENT.md index d2fcd9393..3ea7e2043 100644 --- a/docs/ENVIRONMENT.md +++ b/docs/ENVIRONMENT.md @@ -11,13 +11,14 @@ Unless stated otherwise, a flag-style knob is read as on when its value is a non-empty, non-`0`/`false`/`off` string, and the listed default applies when the variable is unset. None of these are required to run the engine. -## Five of these are also config keys, and the environment wins +## Six of these are also config keys, and the environment wins The weight-residency knobs `VT_GGUF_MMAP`, `VT_GGUF_PREFAULT`, -`VT_MOE_EXPERT_STREAM`, `VT_MOE_EXPERT_STREAM_SLOTS` and -`VT_MOE_EXPERT_STREAM_SLOT_BYTES` can also be set in the server's JSON config, -under the `vllm_cpp` key of `--offload-config`. That is the **documented -deployment surface**; these variables are the **override**: +`VT_MOE_EXPERT_STREAM`, `VT_MOE_EXPERT_STREAM_SLOTS`, +`VT_MOE_EXPERT_STREAM_SLOT_BYTES` and `VT_DEVICE_WEIGHT_BUDGET_BYTES` can also be +set in the server's JSON config, under the `vllm_cpp` key of `--offload-config`. +That is the **documented deployment surface**; these variables are the +**override**: ```text environment variable > --offload-config's vllm_cpp key > built-in default @@ -35,6 +36,15 @@ empty, garbage or non-positive value falls through to the config under the tolerant parse those readers have always used, so `VT_MOE_EXPERT_STREAM_SLOTS=banana` overrides nothing and is not reported as an override. +`VT_DEVICE_WEIGHT_BUDGET_BYTES` is narrower again, and in a different direction. +Its grammar is decimal digits only — no sign, no space, no trailing text — so a +value outside it falls through to +`vllm_cpp.device_fit.weight_budget_bytes` and then to the device probe. Inside +that grammar every value wins, `0` included: `0` is this knob's suppression +spelling rather than an absent one, so it is an override and is reported as one. +It joined the config surface after the other five +([#1127](https://github.com/mudler/vllm.cpp/issues/1127)). + See [USAGE.md](USAGE.md#streaming-routed-experts-from-disk-capacity-mode) for the config form and [`.agents/specs/weight-residency-config.md`](../.agents/specs/weight-residency-config.md) @@ -118,7 +128,7 @@ allocated up front and never grown — the engine prints the resolved values as | `VT_MOE_EXPERT_STREAM_SLOTS` | `64` | How many expert slices stay resident. Parsed as a decimal integer; unset, empty, zero, negative and unparseable values all fall THROUGH to `--offload-config`'s `vllm_cpp.expert_stream.slots` if that is set, and to `64` otherwise. So this variable overrides the config only when its value parses to a positive integer, and the startup override line says so on the same terms. Every slot acquired during a step is protected from eviction until the step ends, so a budget smaller than one step's working set exhausts the cache: those slices fall back to reading the tower directly, which is correct but slow, and is counted. Sized against a real model this wants to be large — the measured run used `8000`. The config key `vllm_cpp.expert_stream.slots`, unlike this variable, REFUSES a zero or negative value at startup instead of silently keeping `64` | | `VT_MOE_EXPERT_STREAM_SLOT_BYTES` | the LARGEST of the gate/up/down slices of the first MoE layer reached | Bytes reserved per slot, fixed for the process's life. Parsed as a decimal integer; unset, empty, zero, negative and unparseable values all fall THROUGH to `--offload-config`'s `vllm_cpp.expert_stream.slot_bytes` if that is set, and to the default otherwise — this variable overrides the config only when its value parses to a positive integer. The default is the largest of the three slices rather than the first one taken, because a dynamic (UD) quant keeps `down_proj` at a higher precision than the gate/up pair and sizing from a gate slice then refuses the first down slice mid-decode. A slice that still does not fit is refused BY NAME (`vt: expert stream: a slice of N bytes exceeds the slot budget of M; raise VT_MOE_EXPERT_STREAM_SLOT_BYTES`) rather than truncated or silently routed back to the mmap path, so a streaming benchmark cannot quietly measure the mmap path instead. The config key `vllm_cpp.expert_stream.slot_bytes`, unlike this variable, refuses a non-positive value at startup | | `VT_MOE_EXPERT_STREAM_STATS_EVERY` | `16` | How many decode steps between the PERIODIC expert-stream statistics line on stderr; `0` silences the periodic line only. Parsed as a decimal integer; unset, empty, negative and unparseable values all keep `16`. The line is `[expert-stream] steps=N hits=H misses=M evictions=E fills=F bytes=B exhausted=X advised=A`. **Exactly one FINAL line is printed when the process ends**, whatever this is set to and whatever the run did, including `steps=0`, for as long as the lane built a store. That is the line to read, and it exists because the row's first published decode figure was measured on a cache that had switched itself off partway through the third token while nothing in the run could say so. **`steps == 0` or `exhausted > 0` means the lane is not streaming**, whatever the startup line claimed. Absence of the final line means either that no store was ever built — in which case the `[expert-stream] ON ...` banner is absent too, and the lane was never reached — or that the process did not run its static destructors (a crash, a signal, `_exit`). A fourth shape exists but no shipped command can produce it: the line is printed once per process, and an internal test seam that flushes it mid-run takes that one print. `docs/USAGE.md` tabulates all four. Environment-ONLY by decision: it changes a diagnostic cadence rather than what the process reserves, so it is the instrument and not the configuration, and `--offload-config` refuses it as an unknown key rather than accepting and dropping it | -| `VT_DEVICE_WEIGHT_BUDGET_BYTES` | the device's own probe (`cudaMemGetInfo` total on CUDA; UNKNOWN, i.e. no check, everywhere else) | Overrides the device memory pool that a GGUF's staged weight bytes are compared against at LOAD time (issue #1123). A GGUF whose weights cannot fit is refused by name during the load instead of dying on the first forward with `vt cuda: cudaMalloc: out of memory` — `Qwen3.8-2.4T-A95B UD-Q1_0` (369.96 GiB) reached a serving state on a 119.631 GiB GB10 after 26 minutes and then died mid-stream, because the larger-than-memory lane that makes it fit is HOST-ONLY. Set this LOWER when something else lives in the pool, or HIGHER (or `0`) to suppress the refusal and get the late failure back — it does not make the model fit. Parsed as decimal digits only: a value with a sign, a space or trailing garbage is IGNORED and the probe stands, because reading a typo as `0` would silently disable the guard. Compared against the pool TOTAL, not the free bytes, so the verdict does not move with contention. The bound counts WEIGHTS only, never the KV cache, activations or the driver context, so a checkpoint just under the pool still passes and can still fail later. It can also count a little too MUCH: a tensor present in the file that this load will not stage — the MTP / `nextn` block on a load with no speculator, 8.33 GiB of the measured 369.96 GiB checkpoint — is still in the sum, so a budget in that narrow window refuses a weight set that would have fitted; raise this value if you land in it ([#1136](https://github.com/mudler/vllm.cpp/issues/1136)). Inert on every platform that does not stage weights, which today means everything except CUDA — including every `--device cpu` load, and including ROCm, Vulkan and Metal, which read the mapping where it lies and have no staging allocation to fail | +| `VT_DEVICE_WEIGHT_BUDGET_BYTES` | the device's own probe (`cudaMemGetInfo` total on CUDA; UNKNOWN, i.e. no check, everywhere else) | Also settable as `--offload-config`'s `vllm_cpp.device_fit.weight_budget_bytes`; this variable overrides it, and a value outside this variable's digits-only grammar falls THROUGH to that key before it falls through to the probe. Overrides the device memory pool that a GGUF's staged weight bytes are compared against at LOAD time (issue #1123). A GGUF whose weights cannot fit is refused by name during the load instead of dying on the first forward with `vt cuda: cudaMalloc: out of memory` — `Qwen3.8-2.4T-A95B UD-Q1_0` (369.96 GiB) reached a serving state on a 119.631 GiB GB10 after 26 minutes and then died mid-stream, because the larger-than-memory lane that makes it fit is HOST-ONLY. Set this LOWER when something else lives in the pool, or HIGHER (or `0`) to suppress the refusal and get the late failure back — it does not make the model fit. Parsed as decimal digits only: a value with a sign, a space or trailing garbage is IGNORED and the next tier decides — the config key if it is set, otherwise the probe — because reading a typo as `0` would silently disable the guard. Compared against the pool TOTAL, not the free bytes, so the verdict does not move with contention. The bound counts WEIGHTS only, never the KV cache, activations or the driver context, so a checkpoint just under the pool still passes and can still fail later. It can also count a little too MUCH: a tensor present in the file that this load will not stage — the MTP / `nextn` block on a load with no speculator, 8.33 GiB of the measured 369.96 GiB checkpoint — is still in the sum, so a budget in that narrow window refuses a weight set that would have fitted; raise this value if you land in it ([#1136](https://github.com/mudler/vllm.cpp/issues/1136)). Inert on every platform that does not stage weights, which today means everything except CUDA — including every `--device cpu` load, and including ROCm, Vulkan and Metal, which read the mapping where it lies and have no staging allocation to fail | ## Rollback and bisect switches diff --git a/docs/USAGE.md b/docs/USAGE.md index 68ba7b5fa..f6cd42837 100644 --- a/docs/USAGE.md +++ b/docs/USAGE.md @@ -320,6 +320,7 @@ build/examples/vllm-cli \ | `--seed S` | (unset) | RNG seed (enables seeded sampling) | | `--stream` | off | Stream token deltas to stdout | | `--speculative-config ''` | (unset) | Speculative decoding, same JSON as vLLM's flag. Every key is checked and none is dropped: an unknown or misspelled name is refused at startup by name, and a real vLLM key this engine does not implement is refused as such ([#1160](https://github.com/mudler/vllm.cpp/issues/1160)). See [docs/SPECULATIVE-DECODING.md](SPECULATIVE-DECODING.md) | +| `--offload-config ''` | (unset) | Weight placement, the same JSON document `vllm-server` takes and the same C ABI field. Both halves: vLLM's mirrored `uva`/`prefetch` device-to-host weight offload, and vllm.cpp's `vllm_cpp` key for the host-to-disk residency tier that makes a checkpoint larger than host RAM loadable. An unknown key at any level of the document is refused at startup by name. Added by [#1135](https://github.com/mudler/vllm.cpp/issues/1135); see [Streaming routed experts from disk](#streaming-routed-experts-from-disk-capacity-mode) | | `--max-num-seqs N` | engine default (32) | Max concurrent sequences. Under speculative decoding on a GDN model the recurrent state is `max-num-seqs x (k+1)` per slot, so this is the knob to lower when a run is refused for state budget | | `--repeat N` | `1` | Load once, then run N blocking completions. Use it to read a warm decode tok/s without paying model load each time. Not supported with `--stream`, which falls back to 1 | | `-h`, `--help` | | Print usage and exit | @@ -2341,7 +2342,7 @@ a stop token early. | `--tool-call-parser ` | `hermes` | Tool-call dialect (42 names over 38 families). `auto` detects from the chat template, `none` disables. For `gemma4`, OpenAI chat uses the text-seam parser (wrapped `<\|tool_call>` **or** bare `call:NAME{ARGS}`) so free-form / detokenized tool bodies still become `tool_calls`. **`inkling` needs `"skip_special_tokens": false` on the request today** — its whole grammar is special tokens and we have no `adjust_request` seam to force the flag off for you, so at the `true` default the detokenizer strips the markers before the parser runs ([#695](https://github.com/mudler/vllm.cpp/issues/695)). `--reasoning-parser inkling` is not registered at all ([#703](https://github.com/mudler/vllm.cpp/issues/703)) | | `--reasoning-parser ` | `none` | Reasoning parser (`think_auto`, `deepseek_r1`, `deepseek_v3`, `holo2`, `mistral`, `minimax_m2`, `minimax_m2_append_think`, `step3`, `olmo3`, `muse_glimmer`, `qwen3`, `mimo`). `auto` detects, `none` disables. `qwen3` and its `mimo` alias are the engine-backed adapter (one upstream class, two registry names): thinking is ON, so a marker-less stream is reasoning and a `` ends reasoning with no ``. `auto` never selects it — a generic `` template resolves to `think_auto`, which is the right default for hybrid-thinking models that may answer with no think block at all | | `--kv-transfer-config ''` | (unset) | External KV connector, same JSON as vLLM's flag. See [docs/KV-OFFLOAD.md](KV-OFFLOAD.md) | -| `--offload-config ''` | (unset) | Weight offload, the same JSON vLLM's `OffloadConfig` takes (distinct from `--kv-transfer-config`, which offloads KV blocks). Parsed and validated at startup, so a malformed document, an unknown backend, an unknown TOP-LEVEL key (the four legal ones are `offload_backend`, `uva`, `prefetch` and `vllm_cpp`) or a validator violation is refused before any model I/O; a backend/field mismatch is a warning, as upstream. **Enabling it fails startup on every model today**: no loader consults the offloader, so the engine refuses the configuration by architecture name rather than accept a budget that frees nothing. A config that leaves offloading disabled still parses and reports normally. On unified memory such as GB10 offload cannot help at all, because host and device share one pool. See [docs/WEIGHT-OFFLOAD.md](WEIGHT-OFFLOAD.md). The same document also carries the **`vllm_cpp` key**, which governs the tier BELOW this one — weights borrowed out of the file mapping rather than moved to host RAM — and which is live rather than refused: see [Streaming routed experts from disk](#streaming-routed-experts-from-disk-capacity-mode). A `vllm_cpp`-only document does not enable vLLM's offload backends and is not subject to the refusal above | +| `--offload-config ''` | (unset) | Weight offload, the same JSON vLLM's `OffloadConfig` takes (distinct from `--kv-transfer-config`, which offloads KV blocks). Parsed and validated at startup, so a malformed document, an unknown backend, an unknown TOP-LEVEL key (the four legal ones are `offload_backend`, `uva`, `prefetch` and `vllm_cpp`) or a validator violation is refused before any model I/O; a backend/field mismatch is a warning, as upstream. **Enabling it fails startup on every model today**: no loader consults the offloader, so the engine refuses the configuration by architecture name rather than accept a budget that frees nothing. A config that leaves offloading disabled still parses and reports normally. On unified memory such as GB10 offload cannot help at all, because host and device share one pool. See [docs/WEIGHT-OFFLOAD.md](WEIGHT-OFFLOAD.md). The same document also carries the **`vllm_cpp` key**, which governs the tier BELOW this one — weights borrowed out of the file mapping rather than moved to host RAM — and which is live rather than refused: see [Streaming routed experts from disk](#streaming-routed-experts-from-disk-capacity-mode). A `vllm_cpp`-only document does not enable vLLM's offload backends and is not subject to the refusal above. The flag is accepted by `vllm-server` (the generate/chat and the pooling/embedding paths), by `vllm-cli`, and by the C ABI; the server's transcription-only path REFUSES it by name, because that path builds no engine and could only accept the document and ignore it ([#1195](https://github.com/mudler/vllm.cpp/issues/1195)) | | `--speculative-config ''` | (unset) | Speculative decoding (`mtp`, `dflash`, `ngram`), same JSON as vLLM's flag. For `mtp`, `num_speculative_tokens` sets the draft DEPTH and defaults to the checkpoint's `mtp_num_hidden_layers`, which is 1 on both gate checkpoints, so the default is unchanged. A value above it must be a multiple of it, mirroring vLLM. Depth cannot move the emitted tokens under greedy decoding, and no speed number is claimed above k=1 yet ([#81](https://github.com/mudler/vllm.cpp/issues/81)). What is gated on CPU at k=1..4 is that the propose runs `k-1` draft decode forwards per propose call, that k drafts reach the verify path, and that the drafts DELIVERED to the verify path vary with depth rather than repeating the first one. That last one is counted over a RUN and never per call, because a correct drafter may resample the same token and this fixture does. Two things are NOT gated there. A draft is never accepted at depth, because acceptance is zero at every depth on the synthetic gate model. And nothing here proves the draft at depth j came from the j-th forward. Both are owed to the GPU gate, which must close the second by comparing the per-depth acceptance RATE against a PADDED control rather than by asserting a non-zero acceptance count, because a padded drafter earns acceptance at depth whenever the target's own greedy continuation repeats a token. `dspark` speculates on the Qwen3.6 gate models (native + Speculators drafts), token-identically to speculative-off, but is not gated on speed: the cross-engine ratio is UNSETTLED, with a matched-and-warm paired measurement of 0.834x against the pinned oracle and the earlier 0.957x-0.989x figures taken against a single COLD oracle invocation on a machine that has since been reimaged. A GGUF target, or a target with no aux multi-tap, is refused by name (`SPEC-DSPARK`). Its sequential Markov sampling runs on device by default; `VT_DSPARK_DEVICE_SAMPLE=0` restores the host loop (token-identical, cost only). The speculative verify runs from a captured CUDA graph, worth +12.2%/+3.5% on the 35B cells; `VT_SPEC_DECODE_GRAPH=0` restores the eager verify (also token-identical). The object is admitted key by key and NOTHING is dropped ([#1160](https://github.com/mudler/vllm.cpp/issues/1160)): the honoured keys are `method`, `num_speculative_tokens`, `model`, `prompt_lookup_min` and `prompt_lookup_max`, plus `draft_sample_method` and `rejection_sample_method` at their upstream defaults `greedy` and `standard`, which are what this engine implements. Any other value of those two names row `SPEC-ACCEPT-VARIANTS` and is refused. A name vLLM's `SpeculativeConfig` declares but this engine does not implement, such as `quantization`, is refused as exactly that, and any other name is refused as unknown with the accepted list. Before this the extra key was discarded, so `draft_sample_method=probabilistic` ran GREEDY and a misspelled `num_speculatve_tokens` took the default, both silently and both at exit 0. For `dspark`, `num_speculative_tokens` may no longer sit BELOW the draft checkpoint's block: DSpark drafts a block, our block is sized from this value alone, and a shorter one drafted a structurally wrong block in silence. It is refused now, before any weight is loaded, naming the block, the config key the block was read from, and the value given ([#1225](https://github.com/mudler/vllm.cpp/issues/1225)). The block is read from the draft config's `dspark_block_size`, or from `block_size` when that key is absent, which is the case on every published Qwen3 draft (`deepseek-ai/dspark_qwen3_4b_block7` and `RadixArk/Qwen3.8-27B-DSpark` both carry `block_size: 7`, so k must be at least 7). vLLM reads only the first key and accepts the shorter value. vLLM also builds its model config BEFORE its speculative config, so a command that names both a target directory it cannot open and a short `k` hears about the target there and about the `k` here. Those are the two recorded divergences, both argued in `.agents/specs/dspark-block-size-guard.md`. A k at or above the block behaves exactly as before. See [docs/SPECULATIVE-DECODING.md](SPECULATIVE-DECODING.md) | | `--language-model-only` / `--no-language-model-only` | off | Disable all multimodal input by setting **every** modality limit to 0, mirroring vLLM's flag of the same name. It is not a "skip the encoder" switch: the server then **refuses** a multimodal request with ``400 At most 0 image(s) may be provided in one prompt. Set `--limit-mm-per-prompt` to increase this limit.`` It does **not** free VRAM yet — nothing gates tower construction on it ([#607](https://github.com/mudler/vllm.cpp/issues/607) wave L3) | | `--limit-mm-per-prompt ''` | (unset ⇒ 999 per modality) | Maximum multimodal input items per prompt, per modality, as the same JSON object vLLM's flag takes: `'{"image": 2, "video": 0}'`, or with profiling options `'{"video": {"count": 1, "num_frames": 32}}'` (the options are validated and ignored — they size dummy inputs for memory profiling, which this engine does not do). A limit can only **lower** what the model/seam supports, never raise it. Malformed JSON, a negative count, or an unknown option on `image` / `video` / `audio` is refused at startup rather than defaulted. An unknown option on any other modality name is dropped rather than refused, mirroring upstream, whose fallback `BaseDummyOptions` is the one such dataclass without `extra="forbid"`. Upstream's dotted spelling (`--limit-mm-per-prompt.image 2`) is not accepted here, as for `--kv-transfer-config` and `--speculative-config` | @@ -2460,7 +2461,7 @@ straight at the file; the config comes from the GGUF's own metadata, so no `config.json` is needed: ```sh -./build/vllm-server --model /path/to/muse-glimmer-30B-kquant-17gb.gguf +./build/examples/vllm-server --model /path/to/muse-glimmer-30B-kquant-17gb.gguf ``` Both published k-quants load (`muse-glimmer-30B-kquant-17gb.gguf` and the mixed @@ -3941,7 +3942,7 @@ save. ```sh VT_MOE_EXPERT_STREAM=1 \ VT_MOE_EXPERT_STREAM_SLOTS=8000 \ - ./build/vllm-cli --model /models/Qwen3.8-2.4T-A95B-UD-Q1_0-00001-of-00008.gguf \ + ./build/examples/vllm-cli --model /models/Qwen3.8-2.4T-A95B-UD-Q1_0-00001-of-00008.gguf \ --prompt "The capital of France is" --max-tokens 16 ``` @@ -3949,12 +3950,13 @@ VT_MOE_EXPERT_STREAM_SLOTS=8000 \ The residency knobs are also config keys, under the `vllm_cpp` key of `--offload-config` — the flag that already carries vLLM's weight-offload -document. One flag covers both tiers: vLLM's own `uva`/`prefetch` keys move +document. `vllm-cli` takes the same flag, so the two recipes here differ only in +which binary they start, not in what each one can express. One flag covers both tiers: vLLM's own `uva`/`prefetch` keys move weights from the device to host RAM, and the `vllm_cpp` key governs the tier below that, where weights stay borrowed out of the file mapping. ```sh -./build/vllm-server --model /models/Qwen3.8-2.4T-A95B-UD-Q1_0-00001-of-00008.gguf \ +./build/examples/vllm-server --model /models/Qwen3.8-2.4T-A95B-UD-Q1_0-00001-of-00008.gguf \ --offload-config '{"vllm_cpp":{"mmap":{"enabled":true,"prefault":false}, "expert_stream":{"enabled":true,"slots":8000}}}' ``` @@ -3966,6 +3968,7 @@ below that, where weights stay borrowed out of the file mapping. | `vllm_cpp.expert_stream.enabled` | `VT_MOE_EXPERT_STREAM` | off | | `vllm_cpp.expert_stream.slots` | `VT_MOE_EXPERT_STREAM_SLOTS` | `64`; a real model wants thousands | | `vllm_cpp.expert_stream.slot_bytes` | `VT_MOE_EXPERT_STREAM_SLOT_BYTES` | the largest gate/up/down slice of the first MoE layer reached | +| `vllm_cpp.device_fit.weight_budget_bytes` | `VT_DEVICE_WEIGHT_BUDGET_BYTES` | the device's own probe (`cudaMemGetInfo` total on CUDA; no check elsewhere). `0` suppresses the load-time device-fit refusal; it is the only key here that accepts `0`, and a negative value is refused | Every field is optional, and an absent field means unchanged, so an `--offload-config` without a `vllm_cpp` key behaves exactly as it did before this @@ -4000,16 +4003,30 @@ kind of thing rather than a mixture. Read the two lines together: `expert_stream beside `VT_MOE_EXPERT_STREAM (expert_stream) OVERRIDES` means the document said on and the variable decides. -**Where the config form reaches, and where it does not.** It reaches the -generate/chat server path (`vllm-server`) and the C ABI's -`vllm_model_params.offload_config`, which is the whole of the library surface. It -does NOT reach `vllm-cli`, nor the server's pooling/embedding and -transcription-only paths, which build their engine parameters without the offload -document at all — the mirrored `uva`/`prefetch` half is dropped there too, and has -been since before this key existed. On those three, use the environment form -above. Recorded under `## Owed` in +**Where the config form reaches, and where it does not.** It reaches +`vllm-server`'s generate/chat path, `vllm-server`'s pooling/embedding path, +`vllm-cli`, and the C ABI's `vllm_model_params.offload_config`, which is the whole +of the library surface. All four take BOTH halves of the document, and the server +parses it once, before it reads the model's architecture, so a typo is refused at +startup whichever path the model then takes. + +It does NOT reach the server's **transcription-only** path, and that path +**refuses the flag** rather than accepting it and doing nothing: + +```text +server: fatal: --offload-config is not supported on a transcription-only model +(ParakeetForCTC). THE MISSING PART: this path serves /v1/audio/transcriptions +through ParakeetTranscriber, which loads its own weights and never builds an +engine, so neither vLLM's uva/prefetch weight offload nor vllm.cpp's vllm_cpp +weight-residency tier has a call site on it. ... +``` + +Use the environment form above on that path, or serve a text-generation or +embedding model. Recorded under `## Owed` in [`.agents/specs/weight-residency-config.md`](../.agents/specs/weight-residency-config.md) -with [#1135](https://github.com/mudler/vllm.cpp/issues/1135). +with [#1195](https://github.com/mudler/vllm.cpp/issues/1195). +[#1135](https://github.com/mudler/vllm.cpp/issues/1135) is the issue this section +answered for the other three. **A misspelled key is refused at startup, not ignored — at every level of the document.** vLLM's own parser ignores a key it does not recognise, which is what @@ -4020,7 +4037,7 @@ spelling is the likeliest typo of all, because every flag around it is hyphenate So the whole document is enumerated and the offender is named: ```text -offload config: unknown key "vllm_cpp.mmapp" (expected one of: mmap expert_stream) +offload config: unknown key "vllm_cpp.mmapp" (expected one of: mmap expert_stream device_fit) offload config: unknown key "vllm-cpp" (expected one of: offload_backend uva prefetch vllm_cpp) offload config: unknown key "uva.cpu_offload_GB" (expected one of: cpu_offload_gb cpu_offload_params) ``` @@ -4148,12 +4165,27 @@ things it deliberately does not do: load will not stage — the MTP / `nextn` block on a load with no speculator, 8.33 GiB of the measured 369.96 GiB checkpoint — is still in the sum, so a budget in that narrow window refuses a weight set that would have fitted. Raise - `VT_DEVICE_WEIGHT_BUDGET_BYTES` if you land in it + the budget if you land in it ([#1136](https://github.com/mudler/vllm.cpp/issues/1136)). -`VT_DEVICE_WEIGHT_BUDGET_BYTES` moves the budget: lower it when something else -lives in the pool, or raise it (or set `0`) to suppress the refusal and get the -late failure back. It does not make the model fit. +**Moving the budget.** Lower it when something else lives in the pool, or raise +it (or set `0`) to suppress the refusal and get the late failure back. It does +not make the model fit. Two ways to say it, and the first beats the second: + +```sh +VT_DEVICE_WEIGHT_BUDGET_BYTES=68719476736 ./build/examples/vllm-server --model ... +./build/examples/vllm-server --model ... \ + --offload-config '{"vllm_cpp":{"device_fit":{"weight_budget_bytes":68719476736}}}' +``` + +The config key is the same `--offload-config` document the residency knobs use, +so one flag still covers weight placement +([#1127](https://github.com/mudler/vllm.cpp/issues/1127)). `0` from either input +suppresses the refusal. The environment variable takes decimal digits only: a +value with a sign, a space or trailing garbage is ignored and falls through to +the config, then to the probe, because reading a typo as `0` would silently +disable the guard. A malformed config value cannot get that far, because the +parser refuses it at startup. **The instrument matters here.** `nvidia-smi --query-gpu=memory.total,memory.free,memory.used` answers `[N/A], [N/A], [N/A]` diff --git a/examples/cli/main.cpp b/examples/cli/main.cpp index b45685319..01a06e989 100644 --- a/examples/cli/main.cpp +++ b/examples/cli/main.cpp @@ -41,6 +41,13 @@ struct Args { bool stream = false; int repeat = 1; // load once, complete N times (warm tok/s) std::string speculative_config; // vLLM --speculative-config JSON; "" => off. + // --offload-config (#1135): the same JSON document `vllm-server` takes, and + // the same ABI field. It carries TWO halves — vLLM's mirrored `uva`/`prefetch` + // weight offload, and vllm.cpp's `vllm_cpp` weight-residency tier, which is + // what makes a checkpoint larger than host RAM loadable. `vllm-cli` had no + // such flag, so the only way to reach the residency tier from it was the + // `VT_*` environment variables. "" => NULL => the byte-identical default. + std::string offload_config; // --device (ABI v14): "auto" (default probe), "cpu", or "cuda" — the names // of vLLM's DeviceConfig.device this build serves. Mapped to the int the ABI // takes (0/1/2) in ParseArgs; an unknown name is rejected there. @@ -72,7 +79,7 @@ void Usage(const char* argv0, std::FILE* out) { " [--seed S] [--stream] [--repeat N]\n" " [--gpu-memory-utilization F] [--kv-cache-memory BYTES]\n" " [--max-num-seqs N]\n" - " [--speculative-config '']\n" + " [--speculative-config ''] [--offload-config '']\n" "\n" "Runs completion(s) over the vllm.cpp C ABI (libvllm). holds\n" "config.json, tokenizer.json and the *.safetensors shards.\n" @@ -121,6 +128,8 @@ bool ParseArgs(int argc, char** argv, Args& a, int& exit_code) { if (a.repeat < 1) a.repeat = 1; } else if (flag == "--speculative-config") { a.speculative_config = NextArg(argc, argv, i); + } else if (flag == "--offload-config") { + a.offload_config = NextArg(argc, argv, i); } else if (flag == "--gpu-memory-utilization") { a.gpu_memory_utilization = std::atof(NextArg(argc, argv, i)); } else if (flag == "--kv-cache-memory") { @@ -207,6 +216,15 @@ int main(int argc, char** argv) { if (!args.speculative_config.empty()) { mp.speculative_config = args.speculative_config.c_str(); } + // --offload-config (#1135): one string, both halves, parsed by the library at + // `vllm_engine_load` (src/capi/vllm_c.cpp) exactly as `vllm-server` parses it. + // This example stays a thin ABI client and adds no parsing of its own, so a + // malformed document, an unknown key at any level of it, or a residency + // document that a decision has already fixed all fail the load below with the + // library's own message. + if (!args.offload_config.empty()) { + mp.offload_config = args.offload_config.c_str(); + } // --device: explicit device selection (ABI v14). 0 (the default) keeps the // accelerator-first probe; an explicitly named absent device fails the load // below with the library's message (never a silent fallback). diff --git a/include/vllm.h b/include/vllm.h index 099ace301..e77f6679f 100644 --- a/include/vllm.h +++ b/include/vllm.h @@ -446,11 +446,13 @@ typedef struct vllm_model_params { * meaning unchanged: * {"vllm_cpp":{"mmap":{"enabled":bool,"prefault":bool}, * "expert_stream":{"enabled":bool,"slots":int, - * "slot_bytes":int}}} + * "slot_bytes":int}, + * "device_fit":{"weight_budget_bytes":int}}} * Precedence per field is environment variable > this document > built-in * default, so an exported VT_GGUF_MMAP / VT_GGUF_PREFAULT / VT_MOE_EXPERT_STREAM - * / VT_MOE_EXPERT_STREAM_SLOTS / VT_MOE_EXPERT_STREAM_SLOT_BYTES still wins; the - * engine prints one line on stderr naming what it installed, plus a second line + * / VT_MOE_EXPERT_STREAM_SLOTS / VT_MOE_EXPERT_STREAM_SLOT_BYTES / + * VT_DEVICE_WEIGHT_BUDGET_BYTES still wins; the engine prints one line on + * stderr naming what it installed, plus a second line * naming the variables that override it when there are any. The engine acts on it * during weight load, so it must be installed before then, which vllm_engine_load * does. Loading a SECOND engine in one process is legal: an absent field means @@ -461,14 +463,27 @@ typedef struct vllm_model_params { * REFUSALS ADDED WITH THAT KEY, all VLLM_ERR_INVALID_ARGUMENT before any model * I/O: an UNKNOWN key anywhere in the document — a misspelled top-level key * (`{"vllm-cpp":...}` with a hyphen, `{"uvaa":...}`), a misspelled key inside - * `vllm_cpp` (`{"vllm_cpp":{"mmapp":...}}`), and a misspelled key inside the - * mirrored `uva` or `prefetch` object (`{"uva":{"cpu_offload_GB":10}}`); a - * wrong-typed field; and a non-positive `slots` or `slot_bytes`. A typo is refused - * rather than defaulted because a silently disabled residency tier, or a budget - * the operator believes is set, is met as an out-of-memory kill rather than as an - * error. Upstream refuses one too: every vLLM config dataclass carries - * `extra="forbid"`. The four legal top-level keys are `offload_backend`, `uva`, - * `prefetch` and `vllm_cpp`. See docs/USAGE.md. + * `vllm_cpp` or inside any of its THREE objects + * (`{"vllm_cpp":{"mmapp":...}}`, + * `{"vllm_cpp":{"device_fit":{"weight_budget":0}}}`), and a misspelled key + * inside the mirrored `uva` or `prefetch` object + * (`{"uva":{"cpu_offload_GB":10}}`); a wrong-typed field; a non-positive + * `slots` or `slot_bytes`; and a NEGATIVE `weight_budget_bytes`. A typo is + * refused rather than defaulted because a silently disabled residency tier, or + * a budget the operator believes is set, is met as an out-of-memory kill rather + * than as an error. Upstream refuses one too: every vLLM config dataclass + * carries `extra="forbid"`. The four legal top-level keys are + * `offload_backend`, `uva`, `prefetch` and `vllm_cpp`. + * + * `weight_budget_bytes` is the ONE field of the six that ACCEPTS `0`, and the + * asymmetry is the reason the key exists. It is a BUDGET, not a size: `0` is + * the documented spelling of "suppress the load-time device-fit refusal and get + * the late failure back", because the fit check reads a zero budget as UNKNOWN + * and decides nothing — exactly what `VT_DEVICE_WEIGHT_BUDGET_BYTES=0` already + * means. `slots` and `slot_bytes` are sizes, and a slot count that silently + * became its default is a cache the operator does not have, so those two keep + * refusing `0`. Only a NEGATIVE budget is refused, and the message says "must + * not be negative" rather than "must be positive". See docs/USAGE.md. * Borrowed for the call only. */ const char* offload_config; /* ── Jump-forward decoding (ABI v10) ─────────────────────────────────────── diff --git a/include/vllm/config/weight_residency.h b/include/vllm/config/weight_residency.h index c4135ec21..11029951a 100644 --- a/include/vllm/config/weight_residency.h +++ b/include/vllm/config/weight_residency.h @@ -23,7 +23,8 @@ // // {"vllm_cpp": {"mmap": {"enabled": bool, "prefault": bool}, // "expert_stream": {"enabled": bool, "slots": int, -// "slot_bytes": int}}} +// "slot_bytes": int}, +// "device_fit": {"weight_budget_bytes": int}}} // // Every field is optional and an absent field means UNCHANGED. That sentence // binds TWO pieces of code, and the first two shapes of this row broke it in both @@ -58,8 +59,8 @@ // and it is the thing an operator flips while watching a run. Recorded so a // later reader sees a decision rather than an omission. // -// THE LATCH, which is the one real hazard here — and it covers TWO of the five -// knobs, not all five. What genuinely freezes: +// THE LATCH, which is the one real hazard here — and it covers THREE of the six +// knobs, not all six. What genuinely freezes: // // * `expert_stream`, because `ResolveExpertStreamRequested` below caches the // answer in a function-local static. (`Qwen35ExpertStreamRequested` is the @@ -72,9 +73,11 @@ // two sizes freezes nothing; building the store does. // // What does NOT freeze: `mmap`, because `GgufLoadPolicy::FromEnv()` is called per -// load and always has been, and `prefault`, because this row deliberately removed -// that site's function-local static (see `ResolveGgufPrefault` below). A second -// engine in one process may therefore still set either of them. +// load and always has been; `prefault`, because this row deliberately removed +// that site's function-local static (see `ResolveGgufPrefault` below); and +// `device_fit.weight_budget_bytes`, which one GGUF load reads once at its fit +// check through no static at all. A second engine in one process may therefore +// still set any of the three. // // So the refusal is scoped to what it can actually justify, and it took three // shapes to get there. `SetWeightResidencyConfig` THROWS only when the incoming @@ -100,6 +103,7 @@ #ifndef VLLM_CONFIG_WEIGHT_RESIDENCY_H_ #define VLLM_CONFIG_WEIGHT_RESIDENCY_H_ +#include #include #include #include @@ -137,6 +141,20 @@ struct WeightResidencyConfig { // case where the computed default is wrong. std::optional expert_stream_slot_bytes; + // `device_fit.weight_budget_bytes` -> VT_DEVICE_WEIGHT_BUDGET_BYTES. The device + // memory pool that a GGUF's staged weight bytes are compared against at LOAD + // time (issue #1123), overriding the platform's own probe. Set it LOWER when + // something else lives in the pool. + // + // ZERO IS LEGAL HERE AND NOWHERE ELSE IN THIS STRUCT, and it is not a + // degenerate size. `CheckDeviceWeightFit` reads a zero budget as UNKNOWN and + // decides nothing, so `0` is the documented spelling of "suppress the refusal + // and get the late failure back" — exactly what + // `VT_DEVICE_WEIGHT_BUDGET_BYTES=0` already means. `slots` and `slot_bytes` are + // sizes and refuse a zero, because a slot count that silently became 64 is a + // cache the operator does not have. A NEGATIVE budget is still refused. + std::optional device_weight_budget_bytes; + // True when the operator set nothing, i.e. the byte-identical default path. bool empty() const; @@ -175,10 +193,15 @@ struct WeightResidencyConfig { // // Throws std::invalid_argument on a malformed document, a non-object document, a // `vllm_cpp` that is not an object, an UNKNOWN key at ANY level of the document, a -// field of the wrong type, or a non-positive `slots` / `slot_bytes`. "Any level" -// means the top level, the inside of `vllm_cpp`, the inside of `vllm_cpp.mmap` and -// `vllm_cpp.expert_stream`, and the inside of the two MIRRORED sub-objects `uva` and -// `prefetch`. The last of those was missing while this comment already claimed it, +// field of the wrong type, a non-positive `slots` / `slot_bytes`, or a NEGATIVE +// `device_fit.weight_budget_bytes`. "Any level" means the top level, the inside of +// `vllm_cpp`, the inside of `vllm_cpp.mmap`, `vllm_cpp.expert_stream` and +// `vllm_cpp.device_fit`, and the inside of the two MIRRORED sub-objects `uva` and +// `prefetch`. A `weight_budget_bytes` of ZERO is accepted, and it is the only zero +// this parser accepts in a field it OWNS: it is the documented spelling of +// "suppress the device-fit refusal", which `CheckDeviceWeightFit` implements by +// reading a zero budget as UNKNOWN. (Values inside `uva` and `prefetch` are not +// this parser's to accept or refuse; it reads their NAMES only.) The last of those was missing while this comment already claimed it, // so `{"uva":{"cpu_offload_GB":10}}` started a server with a 0 GiB budget the // operator believed was set (#1133 H3). TYPE checking inside `uva`/`prefetch` stays // with the mirrored parser that owns those fields; this one only enumerates names. @@ -195,6 +218,10 @@ struct WeightResidencyConfig { // tier keeping a 370 GiB model in 119 GB is worse than a startup error, so it is an // error. // +// `{"vllm_cpp":{"device-fit":...}}` and `{"vllm_cpp":{"device_fit":{"weight_budget_byte":1}}}` +// are refused on the same terms, and for the same consequence one level down: a +// budget the operator believes is set, with the platform probe silently in force. +// // The four legal top-level keys are `offload_backend`, `uva`, `prefetch` and // `vllm_cpp`: the three the mirrored parser reads by name, plus this extension. // Refusing the rest is the MIRROR-FAITHFUL polarity rather than a local invention. @@ -231,7 +258,7 @@ void SetWeightResidencyConfig(const WeightResidencyConfig& config); // The installed config, BY VALUE. Empty until something installs one. A reference // would be read after the lock was released, which is an unsynchronised read behind -// a lock that looks like it covers one; the copy is five optionals. +// a lock that looks like it covers one; the copy is six optionals. WeightResidencyConfig ActiveWeightResidencyConfig(); // The decisions that genuinely freeze, one enumerator each. There is no `kMmap` or @@ -282,10 +309,10 @@ int64_t ResolveResidencyCount(const char* env_name, std::optional configured, int64_t builtin_default); -// ── The five knobs, one named resolver each ─────────────────────────────────── +// ── The six knobs, one named resolver each ──────────────────────────────────── // // Each one owns its environment NAME and its exact historical POLARITY, and each -// is the SOLE reader of its variable after this row. Two reasons this is five +// is the SOLE reader of its variable after this row. Two reasons this is six // functions and not one call at each site with a string literal. // // First, the polarities are NOT the same and one of them is deliberately odd. @@ -381,6 +408,28 @@ int64_t ResolveExpertStreamSlots(); // take). int64_t ResolveExpertStreamSlotBytes(int64_t computed_default); +// `VT_DEVICE_WEIGHT_BUDGET_BYTES` > `vllm_cpp.device_fit.weight_budget_bytes` > +// `probed_total_bytes` (the platform's own `device_memory_total_bytes`). +// +// THE SIXTH KNOB, and the only one of the three integers that does NOT go through +// `ResolveResidencyCount`. That helper ignores a non-positive ENVIRONMENT value, +// which is right for a slot count and wrong here: `0` is this knob's suppression +// spelling and has to survive as a budget of zero, which `CheckDeviceWeightFit` +// then reads as UNKNOWN. (Neither helper filters the CONFIGURED value, because +// the parser already refused what each field considers out of range.) So this function transcribes the environment grammar +// `DeviceWeightBudgetBytes` has used since #1123 — one or more decimal digits and +// nothing else, no sign and no space — and inserts the config UNDER it. A value +// outside that grammar is ignored and falls through to the config, where it used +// to fall through to the probe; a run with no config therefore resolves +// byte-for-byte as before. `strtoull` alone is not enough for the grammar: it +// skips leading whitespace and it ACCEPTS a leading '-' and wraps it to +// ULLONG_MAX, so "-1" would parse as an effectively infinite budget. +// +// NOT LATCHING. It is read once per GGUF load, at the fit check in +// `LoadedEngine::FromModelDir`, through no static, so `ResidencyLatch` gains no +// enumerator and a second engine may still set it. +size_t ResolveDeviceWeightBudgetBytes(size_t probed_total_bytes); + } // namespace vllm #endif // VLLM_CONFIG_WEIGHT_RESIDENCY_H_ diff --git a/include/vllm/model_executor/model_loader/gguf_device_fit.h b/include/vllm/model_executor/model_loader/gguf_device_fit.h index c0bff5b16..b672d0555 100644 --- a/include/vllm/model_executor/model_loader/gguf_device_fit.h +++ b/include/vllm/model_executor/model_loader/gguf_device_fit.h @@ -100,12 +100,22 @@ GgufStagedFootprint GgufStagedWeightFootprint(const GgufFile& gguf, // // `device_memory_total_bytes` is the platform's own probe // (`ResidencyPolicy::device_memory_total_bytes`), which is 0 on every platform -// that does not probe one. `VT_DEVICE_WEIGHT_BUDGET_BYTES` overrides it, for an -// operator whose pool is smaller than the probe reports because something else -// lives in it, and for an operator who wants to attempt the load anyway. A -// value of 0 in the environment means "unknown", i.e. disables the check, and -// an unparseable value is ignored rather than treated as 0, because silently -// disabling a guard on a typo is the failure shape this tree refuses. +// that does not probe one. Two things override it, for an operator whose pool is +// smaller than the probe reports because something else lives in it, and for an +// operator who wants to attempt the load anyway: +// `--offload-config '{"vllm_cpp":{"device_fit":{"weight_budget_bytes":N}}}'`, and +// `VT_DEVICE_WEIGHT_BUDGET_BYTES`, which beats the config. A value of 0 from +// either means "unknown", i.e. disables the check, and an unparseable +// ENVIRONMENT value is ignored rather than treated as 0, because silently +// disabling a guard on a typo is the failure shape this tree refuses; a +// malformed CONFIG value cannot get this far, because the parser refuses it at +// startup. +// +// THE RULE ITSELF LIVES IN `vllm/config/weight_residency.h` +// (`ResolveDeviceWeightBudgetBytes`), and this function is a delegation to it +// (issue #1127). That keeps one reader for the variable, which is what stops the +// install-time override announcement from drifting away from what the resolver +// does with the value. // // TOTAL rather than FREE on purpose: `free` at load time carries the page cache // and whatever else the box is doing, which would make the verdict a function diff --git a/src/vllm/config/weight_residency.cpp b/src/vllm/config/weight_residency.cpp index a14d87c46..c6d994d90 100644 --- a/src/vllm/config/weight_residency.cpp +++ b/src/vllm/config/weight_residency.cpp @@ -4,9 +4,11 @@ #include "vllm/config/weight_residency.h" #include +#include #include #include #include +#include #include #include @@ -41,6 +43,33 @@ std::optional EnvCountThatWins(const char* env_name) { return static_cast(parsed); } +// `VT_DEVICE_WEIGHT_BUDGET_BYTES`'s value, if and only if it would beat a +// configured budget. The grammar is transcribed verbatim from the reader this row +// took over (`DeviceWeightBudgetBytes`, gguf_device_fit.cpp @ #1123): one or more +// DECIMAL DIGITS and nothing else, no sign and no space. Anything else is +// IGNORED, so an environment-only run resolves exactly as before. +// +// It is NOT `EnvCountThatWins`. That helper drops a non-positive value, and `0` is +// this knob's suppression spelling — the operator asking for the refusal to be +// switched off — so dropping it would silently restore the guard the operator +// turned off. `strtoull` alone is not enough for the grammar either: it skips +// leading whitespace, and it ACCEPTS a leading '-' and wraps it to ULLONG_MAX, so +// "-1" would parse as an effectively infinite budget. +// +// ONE rule with two callers, for the same reason the count rule has two: the +// resolver below, and `DescribeEnvOverrides`, which must not announce a value the +// resolver ignores (#1122 L7). +std::optional DeviceWeightBudgetEnvThatWins() { + const char* v = std::getenv("VT_DEVICE_WEIGHT_BUDGET_BYTES"); + if (v == nullptr || v[0] < '0' || v[0] > '9') return std::nullopt; + errno = 0; + char* end = nullptr; + const unsigned long long parsed = // NOLINT(runtime/int) strtoull's type + std::strtoull(v, &end, 10); + if (*end != '\0' || errno != 0) return std::nullopt; + return static_cast(parsed); +} + struct Global { std::mutex mu; WeightResidencyConfig config; @@ -226,12 +255,42 @@ std::optional ExtPositiveInt(const nlohmann::json& obj, const char* key return v; } +// The SAME type check, with ZERO accepted. It exists for exactly one field, and +// the difference from `ExtPositiveInt` is the field's meaning rather than a +// looser rule. `slots` and `slot_bytes` are sizes: a zero there is a degenerate +// cache, so the parser refuses it where the operator can still read the message. +// `device_fit.weight_budget_bytes` is a BUDGET that `CheckDeviceWeightFit` reads +// as UNKNOWN when it is zero, so `0` means "make no fit decision" — the same +// suppression `VT_DEVICE_WEIGHT_BUDGET_BYTES=0` has meant since #1123. Refusing +// it would delete the escape hatch the key exists to give. A NEGATIVE budget is +// still refused, and the message says "must not be negative" rather than "must be +// positive", because the second sentence would be false about a value this parser +// accepts. +std::optional ExtNonNegativeInt(const nlohmann::json& obj, + const char* key, const char* path) { + auto it = obj.find(key); + if (it == obj.end() || it->is_null()) return std::nullopt; + if (!it->is_number_integer()) { + throw std::invalid_argument(std::string("offload config: \"") + path + "." + + key + "\" must be an integer"); + } + const int64_t v = it->get(); + if (v < 0) { + throw std::invalid_argument(std::string("offload config: \"") + path + "." + + key + "\" must not be negative (got " + + std::to_string(v) + + "); 0 means \"suppress the device-fit refusal\""); + } + return v; +} + } // namespace bool WeightResidencyConfig::empty() const { return !mmap.has_value() && !prefault.has_value() && !expert_stream.has_value() && !expert_stream_slots.has_value() && - !expert_stream_slot_bytes.has_value(); + !expert_stream_slot_bytes.has_value() && + !device_weight_budget_bytes.has_value(); } bool WeightResidencyConfig::operator==( @@ -239,7 +298,8 @@ bool WeightResidencyConfig::operator==( return mmap == other.mmap && prefault == other.prefault && expert_stream == other.expert_stream && expert_stream_slots == other.expert_stream_slots && - expert_stream_slot_bytes == other.expert_stream_slot_bytes; + expert_stream_slot_bytes == other.expert_stream_slot_bytes && + device_weight_budget_bytes == other.device_weight_budget_bytes; } std::string WeightResidencyConfig::Describe() const { @@ -263,6 +323,11 @@ std::string WeightResidencyConfig::Describe() const { add_bool("expert_stream", expert_stream); add_int("expert_stream_slots", expert_stream_slots); add_int("expert_stream_slot_bytes", expert_stream_slot_bytes); + // Under the FIELD name rather than the document path, like its five siblings, + // so the line reads as one list of knobs. A budget of 0 prints as + // `device_weight_budget_bytes=0`, which is what the operator asked for and is + // distinguishable from the field being absent, where nothing prints at all. + add_int("device_weight_budget_bytes", device_weight_budget_bytes); return out; } @@ -302,6 +367,13 @@ std::string WeightResidencyConfig::DescribeEnvOverrides() const { {expert_stream_slot_bytes.has_value(), count_wins("VT_MOE_EXPERT_STREAM_SLOT_BYTES"), "VT_MOE_EXPERT_STREAM_SLOT_BYTES", "expert_stream_slot_bytes"}, + // The budget asks its OWN predicate, not `count_wins`: `0` wins here and is + // dropped there, and `0` is the one value an operator is most surprised to + // have inherited from an exported variable, because it switches the + // device-fit refusal off entirely. + {device_weight_budget_bytes.has_value(), + DeviceWeightBudgetEnvThatWins().has_value(), + "VT_DEVICE_WEIGHT_BUDGET_BYTES", "device_weight_budget_bytes"}, }; std::string out; for (const Pair& p : pairs) { @@ -378,7 +450,7 @@ WeightResidencyConfig parse_weight_residency_extension_json( const nlohmann::json* ext = ExtObject(doc, "vllm_cpp", ""); if (ext == nullptr) return cfg; - RejectUnknownKeys(*ext, "vllm_cpp", {"mmap", "expert_stream"}); + RejectUnknownKeys(*ext, "vllm_cpp", {"mmap", "expert_stream", "device_fit"}); if (const nlohmann::json* m = ExtObject(*ext, "mmap", "vllm_cpp")) { RejectUnknownKeys(*m, "vllm_cpp.mmap", {"enabled", "prefault"}); @@ -394,6 +466,19 @@ WeightResidencyConfig parse_weight_residency_extension_json( cfg.expert_stream_slot_bytes = ExtPositiveInt(*s, "slot_bytes", "vllm_cpp.expert_stream"); } + // A THIRD knob family: the load-time device-fit check (#1123), whose budget + // used to be reachable only as `VT_DEVICE_WEIGHT_BUDGET_BYTES` (#1127). It is an + // OBJECT rather than a scalar beside `mmap` and `expert_stream` for two reasons. + // `vllm_cpp` maps a name to an object today and `ExtObject` walks it on that + // assumption, so one scalar sibling would make the level heterogeneous. And this + // is a family, not a member of either existing one — the field name drops the + // family prefix exactly as `VT_GGUF_MMAP` -> `mmap.enabled` and + // `VT_MOE_EXPERT_STREAM_SLOTS` -> `expert_stream.slots` already do. + if (const nlohmann::json* d = ExtObject(*ext, "device_fit", "vllm_cpp")) { + RejectUnknownKeys(*d, "vllm_cpp.device_fit", {"weight_budget_bytes"}); + cfg.device_weight_budget_bytes = + ExtNonNegativeInt(*d, "weight_budget_bytes", "vllm_cpp.device_fit"); + } return cfg; } @@ -475,9 +560,52 @@ std::string DecisionSummary(const Global& g) { return out; } +// THE SETTER REFUSES WHAT THE PARSER REFUSES, because there are TWO doors into +// this struct and only one of them had range rules. +// `parse_weight_residency_extension_json` refuses a non-positive `slots` or +// `slot_bytes` and a negative budget, and every production caller goes through +// it. But `SetWeightResidencyConfig` is declared in a PUBLIC header and takes the +// struct, so a hand-built config reaches the process-global having passed no +// parser at all. +// +// The budget is the dangerous one rather than merely the wrong one. +// `ResolveDeviceWeightBudgetBytes` casts the configured `int64_t` to `size_t`, so +// an installed `-1` resolves to SIZE_MAX: an effectively infinite budget that +// switches the load-time device-fit refusal OFF and says nothing — precisely the +// "a budget the operator believes is set" failure that refusal's own text names. +// The resolver's comment justified that cast by trusting a parser that this door +// does not run, so the trust is made true here instead of being narrated there. +// +// `0` still installs. It is this field's suppression spelling, and a guard that +// refused it would delete the escape hatch the key exists to give. +void RejectOutOfRangeFields(const WeightResidencyConfig& c) { + const auto positive = [](const char* name, std::optional v) { + if (v.has_value() && *v <= 0) { + throw std::invalid_argument(std::string("weight residency config: ") + + name + " must be positive (got " + + std::to_string(*v) + ")"); + } + }; + positive("expert_stream_slots", c.expert_stream_slots); + positive("expert_stream_slot_bytes", c.expert_stream_slot_bytes); + if (c.device_weight_budget_bytes.has_value() && + *c.device_weight_budget_bytes < 0) { + throw std::invalid_argument( + "weight residency config: device_weight_budget_bytes must not be " + "negative (got " + + std::to_string(*c.device_weight_budget_bytes) + + "); 0 means \"suppress the device-fit refusal\", and a negative value " + "would resolve to SIZE_MAX and switch that refusal off silently"); + } +} + } // namespace void SetWeightResidencyConfig(const WeightResidencyConfig& config) { + // BEFORE the lock and before the latch check: a value this struct may not hold + // is refused whatever the process has already decided, and the refusal reads the + // argument only. + RejectOutOfRangeFields(config); Global& g = State(); std::lock_guard lk(g.mu); // REFUSE ONLY WHAT CANNOT BE HONOURED. A document that sets a decided field to a @@ -520,11 +648,18 @@ void SetWeightResidencyConfig(const WeightResidencyConfig& config) { if (config.expert_stream_slot_bytes.has_value()) { g.config.expert_stream_slot_bytes = config.expert_stream_slot_bytes; } + // `has_value()`, not a truth test on the number: `0` is a value the operator + // SET, and it is this field's suppression spelling. A merge that skipped it + // would keep an older budget in force while the operator believed the refusal + // was off. + if (config.device_weight_budget_bytes.has_value()) { + g.config.device_weight_budget_bytes = config.device_weight_budget_bytes; + } } // BY VALUE, and copied under the lock. Returning a reference and then releasing // the mutex gave the caller an unsynchronised read behind a lock that looked like -// it covered one (#1122 L3). The copy is five optionals; the callers are per load, +// it covered one (#1122 L3). The copy is six optionals; the callers are per load, // per prefaulted span (against megabytes of pages) and once per store. WeightResidencyConfig ActiveWeightResidencyConfig() { Global& g = State(); @@ -670,4 +805,17 @@ int64_t ResolveExpertStreamSlotBytes(int64_t computed_default) { ActiveWeightResidencyConfig().expert_stream_slot_bytes, computed_default); } +size_t ResolveDeviceWeightBudgetBytes(size_t probed_total_bytes) { + if (const std::optional winner = DeviceWeightBudgetEnvThatWins()) { + return *winner; + } + const std::optional configured = + ActiveWeightResidencyConfig().device_weight_budget_bytes; + // `has_value()`, not `> 0`. A configured `0` is the operator suppressing the + // refusal, and `CheckDeviceWeightFit` reads a zero budget as UNKNOWN. The parser + // already refused a negative value at startup, so the cast is in range. + if (configured.has_value()) return static_cast(*configured); + return probed_total_bytes; +} + } // namespace vllm diff --git a/src/vllm/entrypoints/openai/server_main.cpp b/src/vllm/entrypoints/openai/server_main.cpp index dfbb53f77..f9a951f59 100644 --- a/src/vllm/entrypoints/openai/server_main.cpp +++ b/src/vllm/entrypoints/openai/server_main.cpp @@ -921,6 +921,55 @@ int VllmServerMain(int argc, char** argv) { : dir.parent_path().filename()) : args.served_model_name; + // ── --offload-config, PARSED ONCE, AHEAD OF THE TASK BRANCH (#1135) ─────── + // + // It used to be parsed after the branch below, in the text path only, so the + // pooling and transcription-only paths built their engine parameters without + // it. An embedding server started with `--offload-config` placed its weights + // as though the flag were absent and said nothing. That dropped BOTH halves + // of the document — vLLM's mirrored `uva`/`prefetch` weight offload and the + // vllm.cpp `vllm_cpp` residency extension — and had done so since before the + // extension existed. + // + // ONE PARSE, and every branch below now answers for the document: the text + // path and the pooling path take both halves, and the transcription-only + // path refuses a non-empty flag rather than dropping it. + // + // Both halves come out of the SAME string: + // `parse_offload_config_json` reads vLLM's three keys and + // `parse_weight_residency_extension_json` reads the extension and closes the + // document against a typo at every level. Two parsers over one string is what + // keeps `include/vllm/config/offload.h` a byte-faithful transcription of + // `vllm/config/offload.py` (which has no disk tier) while the operator still + // types one flag for one concept. + // + // Parsing HERE also refuses a malformed document, an unknown key or a + // validator violation BEFORE the architecture is peeked and before the + // `server: loading model from` line, rather than after both. Upstream's three + // backend/field mismatches are WARNINGS in vLLM and stay warnings here. + std::optional parsed_offload_config; + std::optional parsed_weight_residency; + if (!args.offload_config.empty()) { + vllm::OffloadConfig off_cfg = + vllm::parse_offload_config_json(args.offload_config); + off_cfg.Validate(); + for (const std::string& w : off_cfg.warnings) { + std::fprintf(stderr, "[vllm.cpp] offload_config: %s\n", w.c_str()); + } + parsed_offload_config = std::move(off_cfg); + // ENG-RESIDENCY-CONFIG (#1110): the SAME document also carries the + // vllm.cpp-original `vllm_cpp` key, which governs the tier BELOW vLLM's — + // weights borrowed out of the file mapping rather than moved to host RAM. + // + // Its parser REFUSES a key it does not know, which the mirrored parser does + // not do, and that refusal is load-bearing: a silently ignored + // `{"vllm_cpp":{"mmapp":...}}` starts a server running this tier at its + // defaults and is met as an out-of-memory kill instead of an error. + vllm::WeightResidencyConfig res_cfg = + vllm::parse_weight_residency_extension_json(args.offload_config); + if (!res_cfg.empty()) parsed_weight_residency = std::move(res_cfg); + } + // ── TASK DISPATCH (ARCH-ONE-SURFACE ROW 1): a model dir whose // architectures resolve to a SupportsTranscription-ONLY registration // (Parakeet CTC/RNNT/TDT) serves /v1/audio/transcriptions through the ONE @@ -970,6 +1019,18 @@ int VllmServerMain(int argc, char** argv) { embed_params.max_num_seqs = args.max_num_seqs; embed_params.max_num_batched_tokens = args.max_num_batched_tokens; embed_params.enable_prefix_caching = args.enable_prefix_caching; + // #1135: BOTH halves of `--offload-config`, on the same terms as the text + // path. This branch reaches `LoadedEngine::FromModelDir`, which installs + // the residency document and the weight offloader ahead of every path and + // weight operation it performs — so an embedding model is loaded by the + // same loader and nothing about pooling makes either half inapplicable. + // Before this, both were dropped here in silence. + // + // The other engine flags this block still drops, `--device` among them, + // are #1196. They are the same shape over a wider set and belong to the + // row that owns this dispatch. + embed_params.offload_config = parsed_offload_config; + embed_params.weight_residency = parsed_weight_residency; auto loaded_embed = std::shared_ptr( vllm::entrypoints::LoadedEngine::FromModelDir(args.model_dir, embed_params)); @@ -1017,6 +1078,33 @@ int VllmServerMain(int argc, char** argv) { } if (transcription_only) { + // #1135, and this arm is a REFUSAL rather than a wiring. The + // transcription stack has no seam either half of `--offload-config` + // could reach: `ParakeetTranscriber::FromDir` builds no `EngineParams` + // and calls no `LoadedEngine::FromModelDir`, so this path runs no + // `SetWeightResidencyConfig`, no `CreateWeightOffloader`, no GGUF + // mapping and no expert slot store. There is no field of either half + // that any code here could read. + // + // AGENTS.md: refuse an unimplemented arm with a message that names the + // missing part, and record the arm as owed. Accepting the flag and + // warning would leave a server running while it holds a placement + // instruction it does not follow, which is the failure #1135 was filed + // about. The document is still PARSED above, so a typo in it is refused + // here on the same terms as everywhere else. + if (!args.offload_config.empty()) { + throw std::invalid_argument( + "--offload-config is not supported on a transcription-only model " + "(" + archs[0] + + "). THE MISSING PART: this path serves /v1/audio/transcriptions " + "through ParakeetTranscriber, which loads its own weights and " + "never builds an engine, so neither vLLM's uva/prefetch weight " + "offload nor vllm.cpp's vllm_cpp weight-residency tier has a call " + "site on it. Both halves of the document would be accepted and " + "then ignored, which is issue #1135. Remove the flag, or serve a " + "text-generation or embedding model, which honour it. Tracked as " + "issue #1195"); + } std::cerr << "server: transcription-only model (" << archs[0] << "); serving /v1/audio/transcriptions\n"; auto transcriber = @@ -1108,37 +1196,14 @@ int VllmServerMain(int argc, char** argv) { } engine_params.kv_transfer_config = std::move(kv_cfg); } - // --offload-config: WEIGHT offload (ENG-WEIGHT-OFFLOAD W0b). Empty (default) - // leaves the optional unset — the byte-identical no-offload path. Parsed and - // VALIDATED here so a malformed document or a validator violation is refused - // at startup rather than surfacing later. Upstream's three backend/field - // mismatches are WARNINGS in vLLM and stay warnings here. - if (!args.offload_config.empty()) { - vllm::OffloadConfig off_cfg = - vllm::parse_offload_config_json(args.offload_config); - off_cfg.Validate(); - for (const std::string& w : off_cfg.warnings) { - std::fprintf(stderr, "[vllm.cpp] offload_config: %s\n", w.c_str()); - } - engine_params.offload_config = std::move(off_cfg); - // ENG-RESIDENCY-CONFIG (#1110): the SAME document also carries the - // vllm.cpp-original `vllm_cpp` key, which governs the tier BELOW vLLM's — - // weights borrowed out of the file mapping rather than moved to host RAM. - // Two parsers over one string, each reading only its own half, is what keeps - // `include/vllm/config/offload.h` a byte-faithful transcription of - // `vllm/config/offload.py` (which has no disk tier) while still giving the - // operator one flag for one concept. - // - // Parsed HERE, beside the mirrored half, for the same reason: a mistyped key - // costs a second rather than a full load. The extension REFUSES a key it - // does not know, which the mirrored parser does not do — and that refusal is - // load-bearing, because a silently ignored `{"vllm_cpp":{"mmapp":...}}` - // starts a server running this tier at its defaults and is discovered as an - // out-of-memory kill instead of an error. - vllm::WeightResidencyConfig res_cfg = - vllm::parse_weight_residency_extension_json(args.offload_config); - if (!res_cfg.empty()) engine_params.weight_residency = std::move(res_cfg); - } + // --offload-config: WEIGHT offload (ENG-WEIGHT-OFFLOAD W0b) and the + // vllm.cpp `vllm_cpp` residency tier (ENG-RESIDENCY-CONFIG). Both halves were + // parsed and validated above, ahead of the task branch, so that the pooling + // path gets them too and a malformed document is refused before the + // architecture peek (#1135). Empty (default) leaves both optionals unset — + // the byte-identical no-offload path. + engine_params.offload_config = parsed_offload_config; + engine_params.weight_residency = parsed_weight_residency; // --speculative-config: speculative decoding (SPEC-MTP I5d). Absent (default) // leaves the optional unset — the byte-identical no-speculation path. The // parse validates method/k here; n_predict + the resolved k are finalized in diff --git a/src/vllm/model_executor/model_loader/gguf_device_fit.cpp b/src/vllm/model_executor/model_loader/gguf_device_fit.cpp index 56f152667..2451294be 100644 --- a/src/vllm/model_executor/model_loader/gguf_device_fit.cpp +++ b/src/vllm/model_executor/model_loader/gguf_device_fit.cpp @@ -1,10 +1,10 @@ // ENG-EXPERT-STREAM, issue #1123. See the header for what this decides and why. #include "vllm/model_executor/model_loader/gguf_device_fit.h" -#include -#include #include +#include "vllm/config/weight_residency.h" + namespace vllm { namespace { @@ -49,23 +49,20 @@ GgufStagedFootprint GgufStagedWeightFootprint(const GgufFile& gguf, } size_t DeviceWeightBudgetBytes(size_t device_memory_total_bytes) { - const char* override_env = std::getenv("VT_DEVICE_WEIGHT_BUDGET_BYTES"); - // A malformed value is IGNORED, never read as 0. Reading it as 0 would - // silently disable the guard on a typo, which is the invisible-fallback shape - // this tree refuses elsewhere. `strtoull` alone is not enough for that: it - // skips leading whitespace, and it ACCEPTS a leading '-' and wraps it to - // ULLONG_MAX, so "-1" would parse as an effectively infinite budget. The - // accepted grammar is therefore explicit: one or more decimal digits, nothing - // else, no sign and no space. - if (override_env != nullptr && override_env[0] >= '0' && - override_env[0] <= '9') { - errno = 0; - char* end = nullptr; - const unsigned long long parsed = // NOLINT(runtime/int) strtoull's type - std::strtoull(override_env, &end, 10); - if (*end == '\0' && errno == 0) return static_cast(parsed); - } - return device_memory_total_bytes; + // THE PRECEDENCE LIVES IN ONE PLACE, and since #1127 that place is + // `vllm/config/weight_residency.h`: `VT_DEVICE_WEIGHT_BUDGET_BYTES` > + // `vllm_cpp.device_fit.weight_budget_bytes` > this probe. The environment + // grammar is unchanged — decimal digits only, so a malformed value is IGNORED + // rather than read as 0, which would silently disable the guard on a typo — + // and what it now falls through to is the config rather than straight to the + // probe. A run with no config therefore resolves byte-for-byte as it did + // before the key existed. + // + // This function keeps its name and its callers. Moving the rule rather than + // the entry point is what lets `ResolveDeviceWeightBudgetBytes` be the SOLE + // reader of the variable, which is the contract every knob in that header has + // and the reason `DescribeEnvOverrides` cannot drift from the resolver. + return ResolveDeviceWeightBudgetBytes(device_memory_total_bytes); } DeviceWeightFit CheckDeviceWeightFit(const GgufFile& gguf, @@ -109,7 +106,9 @@ DeviceWeightFit CheckDeviceWeightFit(const GgufFile& gguf, "checkpoint that fits the pool. This is refused at LOAD on purpose: " "before this check the load succeeded and the FIRST forward died with " "'vt cuda: cudaMalloc: out of memory'. Setting " - "VT_DEVICE_WEIGHT_BUDGET_BYTES higher (or to 0) suppresses this refusal " + "VT_DEVICE_WEIGHT_BUDGET_BYTES higher (or to 0), or " + "--offload-config '{\"vllm_cpp\":{\"device_fit\":" + "{\"weight_budget_bytes\":0}}}', suppresses this refusal " "and restores that late failure; it does not make the model fit."; return fit; } diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt index 863c3dc3f..36555d942 100644 --- a/tests/CMakeLists.txt +++ b/tests/CMakeLists.txt @@ -1168,6 +1168,22 @@ vllm_cpp_add_test(test_expert_stream_latch vllm_cpp_add_test(test_weight_residency_reach vllm/entrypoints/test_weight_residency_reach.cpp) target_include_directories(test_weight_residency_reach PRIVATE ${CMAKE_SOURCE_DIR}/src) +# `vllm-cli`'s `--offload-config` (#1135). It runs the BINARY, because that is the +# only thing that proves the FLAG arrives: `test_weight_residency_reach` already +# drives `vllm_engine_load` with the same JSON string, so a test that called the +# ABI here would re-prove the ABI. `$` resolves at generate +# time, after every directory is processed; the build ORDER dependency is declared +# in the top-level CMakeLists, where the `vllm-cli` target already exists. +# Registered UNCONDITIONALLY, and with examples off it leaves `VLLM_CLI_BINARY` +# undefined: the suite then exits 77 and CTest reports it Skipped, which is what +# `SKIP_RETURN_CODE` above is for. Registering it only under the option would hide +# the gate's absence instead of reporting it. +vllm_cpp_add_test(test_cli_offload_config + vllm/entrypoints/test_cli_offload_config.cpp) +if(TARGET vllm-cli OR VLLM_CPP_BUILD_EXAMPLES) + target_compile_definitions(test_cli_offload_config + PRIVATE VLLM_CLI_BINARY="$") +endif() vllm_cpp_add_test(test_weight_offloader vllm/model_executor/test_weight_offloader.cpp) vllm_cpp_add_test(test_weight_offload_policy vllm/model_executor/test_weight_offload_policy.cpp) vllm_cpp_add_test(test_expert_slot_cache vllm/model_executor/test_expert_slot_cache.cpp) diff --git a/tests/vllm/config/test_weight_residency_config.cpp b/tests/vllm/config/test_weight_residency_config.cpp index 024e3aa7c..463c9baaa 100644 --- a/tests/vllm/config/test_weight_residency_config.cpp +++ b/tests/vllm/config/test_weight_residency_config.cpp @@ -882,3 +882,311 @@ TEST_CASE("residency config: reading mmap or prefault latches NOTHING, so a seco CHECK(vllm::ResolveGgufMmap(/*builtin_default=*/true) == false); CHECK(vllm::ResolveGgufPrefault() == false); } + +// ─── W2: `vllm_cpp.device_fit.weight_budget_bytes` (issue #1127) ────────────── +// +// The sixth knob, and the only one whose legal range includes ZERO. `slots` and +// `slot_bytes` are sizes, so a zero there is refused: a slot count that silently +// became 64 is a cache the operator does not have. `0` on the budget is the +// DOCUMENTED spelling of "suppress the device-fit refusal", identical to +// `VT_DEVICE_WEIGHT_BUDGET_BYTES=0`, because `CheckDeviceWeightFit` reads a zero +// budget as UNKNOWN and decides nothing. Refusing it would remove the escape hatch +// the key exists to give, so it parses through its own non-negative helper. + +TEST_CASE("residency config: the device-fit budget parses, and ZERO is legal") { + const vllm::WeightResidencyConfig c = + vllm::parse_weight_residency_extension_json( + R"({"vllm_cpp":{"device_fit":{"weight_budget_bytes":137438953472}}})"); + REQUIRE(c.device_weight_budget_bytes.has_value()); + CHECK(*c.device_weight_budget_bytes == 137438953472LL); + CHECK_FALSE(c.empty()); + // It reaches the install line under its own name, so an operator reading the + // line can tell a budget was set from a budget that fell through to the probe. + CHECK(Mentions(c.Describe(), "device_weight_budget_bytes=137438953472")); + + // ZERO. This is the suppression spelling and it must survive the parse as an + // ENGAGED optional: `nullopt` would fall through to the probe, which is the + // opposite of what the operator asked for. + const vllm::WeightResidencyConfig zero = + vllm::parse_weight_residency_extension_json( + R"({"vllm_cpp":{"device_fit":{"weight_budget_bytes":0}}})"); + REQUIRE(zero.device_weight_budget_bytes.has_value()); + CHECK(*zero.device_weight_budget_bytes == 0); + CHECK_FALSE(zero.empty()); + + // An absent `device_fit` object leaves it unset, and an empty one does too. + CHECK_FALSE(vllm::parse_weight_residency_extension_json( + R"({"vllm_cpp":{"mmap":{"enabled":true}}})") + .device_weight_budget_bytes.has_value()); + const vllm::WeightResidencyConfig empty_obj = + vllm::parse_weight_residency_extension_json( + R"({"vllm_cpp":{"device_fit":{}}})"); + CHECK_FALSE(empty_obj.device_weight_budget_bytes.has_value()); + CHECK(empty_obj.empty()); +} + +TEST_CASE("residency config: a misspelled device_fit key is REFUSED, never ignored") { + // The whole reason this parser enumerates: the mirrored parser ignores what it + // does not know, so a typo in a new level would start a server whose device-fit + // budget is the probe while the operator believes it is the number typed. On a + // box where the probe is smaller than the checkpoint, that typo is the refusal + // the operator was trying to suppress. + const std::string level = + RefusalMessage(R"({"vllm_cpp":{"device_fitt":{"weight_budget_bytes":1}}})"); + CHECK(Mentions(level, "unknown key \"vllm_cpp.device_fitt\"")); + CHECK(Mentions(level, "expected one of: mmap expert_stream device_fit")); + + // The hyphenated spelling of the new level, for the same reason the hyphenated + // `vllm-cpp` is pinned: every flag around it is hyphenated. + CHECK(Mentions( + RefusalMessage(R"({"vllm_cpp":{"device-fit":{"weight_budget_bytes":1}}})"), + "unknown key \"vllm_cpp.device-fit\"")); + + // And the field inside it. `weight_budget_byte` is the likeliest of these, + // because the singular reads correctly in English and the plural is the name. + const std::string field = + RefusalMessage(R"({"vllm_cpp":{"device_fit":{"weight_budget_byte":1}}})"); + CHECK(Mentions(field, "unknown key \"vllm_cpp.device_fit.weight_budget_byte\"")); + CHECK(Mentions(field, "expected one of: weight_budget_bytes")); + + // The name a reader might carry over from the environment variable. + CHECK(Mentions(RefusalMessage( + R"({"vllm_cpp":{"device_fit":{"device_weight_budget_bytes":1}}})"), + "unknown key \"vllm_cpp.device_fit.device_weight_budget_bytes\"")); +} + +TEST_CASE("residency config: a wrong-typed or NEGATIVE budget is REFUSED") { + CHECK(Mentions( + RefusalMessage(R"({"vllm_cpp":{"device_fit":{"weight_budget_bytes":"10"}}})"), + "\"vllm_cpp.device_fit.weight_budget_bytes\" must be an integer")); + CHECK(Mentions( + RefusalMessage(R"({"vllm_cpp":{"device_fit":{"weight_budget_bytes":1.5}}})"), + "\"vllm_cpp.device_fit.weight_budget_bytes\" must be an integer")); + CHECK(Mentions( + RefusalMessage(R"({"vllm_cpp":{"device_fit":{"weight_budget_bytes":true}}})"), + "\"vllm_cpp.device_fit.weight_budget_bytes\" must be an integer")); + + // NEGATIVE is refused and ZERO is not, which is the whole distinction between + // this field's helper and the one `slots` uses. A message that said "positive" + // here would be lying about a value the parser accepts. + const std::string neg = RefusalMessage( + R"({"vllm_cpp":{"device_fit":{"weight_budget_bytes":-1}}})"); + CHECK(Mentions(neg, "\"vllm_cpp.device_fit.weight_budget_bytes\"")); + CHECK(Mentions(neg, "must not be negative")); + CHECK(Mentions(neg, "-1")); + + // `device_fit` itself must be an object, and the message names the DOCUMENT + // path rather than a hardcoded prefix. + CHECK(Mentions(RefusalMessage(R"({"vllm_cpp":{"device_fit":5}})"), + "\"vllm_cpp.device_fit\" must be a JSON object")); +} + +TEST_CASE("residency config: the budget resolves env > config > probed total") { + ResidencyFixture fx; + constexpr size_t kProbed = 128ULL * 1024 * 1024 * 1024; + + // Neither input: the probe stands, byte-for-byte as before this key existed. + ::unsetenv("VT_DEVICE_WEIGHT_BUDGET_BYTES"); + CHECK(vllm::ResolveDeviceWeightBudgetBytes(kProbed) == kProbed); + + // Config only. + vllm::SetWeightResidencyConfig(vllm::parse_weight_residency_extension_json( + R"({"vllm_cpp":{"device_fit":{"weight_budget_bytes":4096}}})")); + CHECK(vllm::ResolveDeviceWeightBudgetBytes(kProbed) == 4096U); + + // A config ZERO reaches the resolver as zero, which `CheckDeviceWeightFit` + // reads as UNKNOWN and therefore as "do not refuse". The suppression spelling + // has to survive the resolver as well as the parser. + vllm::ResetWeightResidencyConfigForTesting(); + vllm::SetWeightResidencyConfig(vllm::parse_weight_residency_extension_json( + R"({"vllm_cpp":{"device_fit":{"weight_budget_bytes":0}}})")); + CHECK(vllm::ResolveDeviceWeightBudgetBytes(kProbed) == 0U); + + // Environment beats config, in BOTH directions. A benchmark arm switched by an + // exported variable is why this precedence exists, and an override that could + // not raise the budget back would not be one. + ::setenv("VT_DEVICE_WEIGHT_BUDGET_BYTES", "8192", 1); + CHECK(vllm::ResolveDeviceWeightBudgetBytes(kProbed) == 8192U); + ::setenv("VT_DEVICE_WEIGHT_BUDGET_BYTES", "0", 1); + CHECK(vllm::ResolveDeviceWeightBudgetBytes(kProbed) == 0U); + + // THE ENVIRONMENT GRAMMAR IS UNCHANGED: decimal digits only. A signed, spaced + // or garbage value is IGNORED, and what it now falls through to is the CONFIG + // rather than the probe. Reading "-1" as a budget would wrap to ULLONG_MAX and + // silently disable the guard, which is why the grammar is explicit. + vllm::ResetWeightResidencyConfigForTesting(); + vllm::SetWeightResidencyConfig(vllm::parse_weight_residency_extension_json( + R"({"vllm_cpp":{"device_fit":{"weight_budget_bytes":4096}}})")); + for (const char* bad : {"-1", " 64", "64x", "", "banana", "+64"}) { + ::setenv("VT_DEVICE_WEIGHT_BUDGET_BYTES", bad, 1); + CAPTURE(bad); + CHECK(vllm::ResolveDeviceWeightBudgetBytes(kProbed) == 4096U); + } + // ...and with no config either, the same bad values leave the probe standing. + vllm::ResetWeightResidencyConfigForTesting(); + for (const char* bad : {"-1", " 64", "64x", "", "banana"}) { + ::setenv("VT_DEVICE_WEIGHT_BUDGET_BYTES", bad, 1); + CAPTURE(bad); + CHECK(vllm::ResolveDeviceWeightBudgetBytes(kProbed) == kProbed); + } + ::unsetenv("VT_DEVICE_WEIGHT_BUDGET_BYTES"); +} + +TEST_CASE("residency config: the budget's override note asks whether the variable WINS") { + ResidencyFixture fx; + const vllm::WeightResidencyConfig c = + vllm::parse_weight_residency_extension_json( + R"({"vllm_cpp":{"device_fit":{"weight_budget_bytes":4096}}})"); + + ::unsetenv("VT_DEVICE_WEIGHT_BUDGET_BYTES"); + CHECK(c.DescribeEnvOverrides().empty()); + + ::setenv("VT_DEVICE_WEIGHT_BUDGET_BYTES", "8192", 1); + const std::string note = c.DescribeEnvOverrides(); + CHECK(Mentions(note, "VT_DEVICE_WEIGHT_BUDGET_BYTES")); + CHECK(Mentions(note, "device_weight_budget_bytes")); + + // `0` IS an override — it is the suppression value, not an absent one — so it + // must be announced. A predicate that tested for a positive number would drop + // exactly the value an operator is most surprised to have inherited. + ::setenv("VT_DEVICE_WEIGHT_BUDGET_BYTES", "0", 1); + CHECK(Mentions(c.DescribeEnvOverrides(), "VT_DEVICE_WEIGHT_BUDGET_BYTES")); + + // A value the resolver IGNORES is not an override, and announcing it would send + // the operator after a line that decides nothing (the #1122 L7 shape). + for (const char* bad : {"-1", " 64", "64x", "banana"}) { + ::setenv("VT_DEVICE_WEIGHT_BUDGET_BYTES", bad, 1); + CAPTURE(bad); + CHECK(c.DescribeEnvOverrides().empty()); + } + + // And a document that does not SET the budget is never reported for it, + // whatever the variable says. + ::setenv("VT_DEVICE_WEIGHT_BUDGET_BYTES", "8192", 1); + CHECK_FALSE(Mentions(vllm::parse_weight_residency_extension_json( + R"({"vllm_cpp":{"mmap":{"enabled":true}}})") + .DescribeEnvOverrides(), + "VT_DEVICE_WEIGHT_BUDGET_BYTES")); + ::unsetenv("VT_DEVICE_WEIGHT_BUDGET_BYTES"); +} + +TEST_CASE("residency config: ABSENT MEANS UNCHANGED for the budget, at both ends") { + ResidencyFixture fx; + ::unsetenv("VT_DEVICE_WEIGHT_BUDGET_BYTES"); + + // Two DIFFERENT partial documents, which is the shape that caught #1133 H2: a + // second document that restates the first's field cannot show a dropped field. + vllm::SetWeightResidencyConfig(vllm::parse_weight_residency_extension_json( + R"({"vllm_cpp":{"device_fit":{"weight_budget_bytes":4096},)" + R"("expert_stream":{"slots":8000}}})")); + vllm::SetWeightResidencyConfig(vllm::parse_weight_residency_extension_json( + R"({"vllm_cpp":{"mmap":{"enabled":true}}})")); + + const vllm::WeightResidencyConfig installed = + vllm::ActiveWeightResidencyConfig(); + REQUIRE(installed.device_weight_budget_bytes.has_value()); + CHECK(*installed.device_weight_budget_bytes == 4096); + REQUIRE(installed.expert_stream_slots.has_value()); + CHECK(*installed.expert_stream_slots == 8000); + REQUIRE(installed.mmap.has_value()); + CHECK(*installed.mmap == true); + CHECK(vllm::ResolveDeviceWeightBudgetBytes(999) == 4096U); + + // And a later document that DOES set it overwrites it, because that is what + // "set" means. `0` is a set value, so it must overwrite too rather than read as + // "unchanged" — the one place where the suppression spelling and the absent + // spelling would be confusable. + vllm::SetWeightResidencyConfig(vllm::parse_weight_residency_extension_json( + R"({"vllm_cpp":{"device_fit":{"weight_budget_bytes":0}}})")); + REQUIRE(vllm::ActiveWeightResidencyConfig() + .device_weight_budget_bytes.has_value()); + CHECK(*vllm::ActiveWeightResidencyConfig().device_weight_budget_bytes == 0); + CHECK(vllm::ResolveDeviceWeightBudgetBytes(999) == 0U); +} + +TEST_CASE("residency config: the budget LATCHES NOTHING, so a late install is accepted") { + ResidencyFixture fx; + ::unsetenv("VT_DEVICE_WEIGHT_BUDGET_BYTES"); + + // Read it, which is what a completed GGUF load does at its fit check. Nothing + // caches the answer, so nothing freezes. + CHECK(vllm::ResolveDeviceWeightBudgetBytes(4096) == 4096U); + CHECK_FALSE(vllm::WeightResidencyLatched()); + + // A second engine may therefore still set it, and the refusal must not fire. + // `expert_stream` and the slot geometry are the only two decisions that freeze, + // and this key is in neither. + CHECK_NOTHROW( + vllm::SetWeightResidencyConfig(vllm::parse_weight_residency_extension_json( + R"({"vllm_cpp":{"device_fit":{"weight_budget_bytes":8192}}})"))); + CHECK(vllm::ResolveDeviceWeightBudgetBytes(4096) == 8192U); + + // Even once the streaming decision HAS been taken, a budget-only document is + // not a change to it and is accepted. + (void)vllm::ResolveExpertStreamRequested(); + REQUIRE(vllm::WeightResidencyLatched(vllm::ResidencyLatch::kExpertStream)); + CHECK_NOTHROW( + vllm::SetWeightResidencyConfig(vllm::parse_weight_residency_extension_json( + R"({"vllm_cpp":{"device_fit":{"weight_budget_bytes":16384}}})"))); + CHECK(vllm::ResolveDeviceWeightBudgetBytes(4096) == 16384U); +} + +TEST_CASE("residency config: the SETTER refuses what the PARSER refuses") { + ResidencyFixture fx; + ::unsetenv("VT_DEVICE_WEIGHT_BUDGET_BYTES"); + + // THE OTHER DOOR. Every production caller reaches the process-global through + // `parse_weight_residency_extension_json`, which refuses a negative budget and a + // non-positive slot size. But `SetWeightResidencyConfig` is declared in a public + // header and takes the STRUCT, so a hand-built config is a legal way in that + // skips those rules entirely, and the resolver's own comment used to justify its + // cast by trusting a parser that had not necessarily run. + // + // The budget is the dangerous one rather than merely the wrong one: + // `ResolveDeviceWeightBudgetBytes` casts to `size_t`, so an installed `-1` + // resolves to SIZE_MAX — an effectively infinite budget that switches the + // load-time device-fit refusal OFF without a word. That is the exact + // "a budget the operator believes is set" failure the parser's own refusal text + // names. + vllm::WeightResidencyConfig negative_budget; + negative_budget.device_weight_budget_bytes = -1; + std::string message; + try { + vllm::SetWeightResidencyConfig(negative_budget); + message = "ACCEPTED (no throw)"; + } catch (const std::invalid_argument& e) { + message = e.what(); + } catch (const std::exception& e) { + message = std::string("WRONG EXCEPTION TYPE: ") + e.what(); + } + CAPTURE(message); + CHECK(Mentions(message, "device_weight_budget_bytes")); + CHECK(Mentions(message, "must not be negative")); + CHECK(Mentions(message, "-1")); + + // ...and nothing was installed, so the check the refusal protects still reads + // the probe. Without this the case would pass on an implementation that threw + // AFTER writing the value. + CHECK_FALSE( + vllm::ActiveWeightResidencyConfig().device_weight_budget_bytes.has_value()); + CHECK(vllm::ResolveDeviceWeightBudgetBytes(4096) == 4096U); + + // The two sizes get the parser's rule at this door too, so the two doors into + // one struct state one rule. + vllm::WeightResidencyConfig zero_slots; + zero_slots.expert_stream_slots = 0; + CHECK_THROWS_AS(vllm::SetWeightResidencyConfig(zero_slots), + std::invalid_argument); + vllm::WeightResidencyConfig negative_slot_bytes; + negative_slot_bytes.expert_stream_slot_bytes = -8; + CHECK_THROWS_AS(vllm::SetWeightResidencyConfig(negative_slot_bytes), + std::invalid_argument); + + // Everything the PARSER accepts, the setter still accepts — `0` for the budget + // above all, because it is this field's suppression spelling and a guard that + // refused it would delete the escape hatch the key exists to give. + vllm::WeightResidencyConfig zero_budget; + zero_budget.device_weight_budget_bytes = 0; + CHECK_NOTHROW(vllm::SetWeightResidencyConfig(zero_budget)); + CHECK(vllm::ResolveDeviceWeightBudgetBytes(4096) == 0U); +} diff --git a/tests/vllm/entrypoints/openai/test_serve_residency_config.cpp b/tests/vllm/entrypoints/openai/test_serve_residency_config.cpp index 103084a2c..c68ac58d1 100644 --- a/tests/vllm/entrypoints/openai/test_serve_residency_config.cpp +++ b/tests/vllm/entrypoints/openai/test_serve_residency_config.cpp @@ -33,6 +33,8 @@ #include #include #include +#include +#include #include #include #include @@ -53,7 +55,7 @@ constexpr const char* kPostParseBanner = "server: request logging"; // them. It does NOT name resolved values: the streaming answer is cached the first // time it is asked, so resolving it at install would move that decision ahead of the // weight load. That constraint binds `expert_stream` alone; the line reports the -// document for all five fields so it reports one kind of thing rather than a mixture. +// document for all six fields so it reports one kind of thing rather than a mixture. // The pair of lines is what lets a run whose document was overridden say so; see // CASE 5. constexpr const char* kInstallLine = "engine: weight residency"; @@ -215,12 +217,14 @@ TEST_CASE("serve: a mistyped vllm_cpp key aborts at startup, naming the key") { CHECK_FALSE(Contains(run.output, kInstallLine)); CHECK_FALSE(Contains(run.output, "SERVE_RC=0")); - // ...and it died BEFORE any model I/O. `server: loading model from` is NOT the - // marker for that — it is printed while parsing the configs, before the load - // starts, so asserting its absence here would be asserting the wrong thing - // (measured: the well-formed run below prints it too). The marker is the loader's - // OWN first complaint about the missing checkpoint, which the well-formed - // document reaches and the mistyped one never does. + // ...and it died BEFORE any model I/O. The marker is the loader's OWN first + // complaint about the missing checkpoint, which the well-formed document + // reaches and the mistyped one never does. Asserting the ABSENCE of + // `server: loading model from` would gate the parse's POSITION rather than its + // effect, and W2 moved that position: the parse now runs ahead of the + // architecture branch, so the abort happens before that line instead of after + // it (#1135). This assertion is stable across both positions, which is why it + // is the one this case makes. const std::string kLoaderReached = "model path is not a directory"; CHECK_FALSE(Contains(run.output, kLoaderReached)); @@ -292,3 +296,182 @@ TEST_CASE("serve: an exported VT_ override beats the document, and is reported") CHECK(Contains(clean.output, "expert_stream=on")); CHECK_FALSE(Contains(clean.output, "OVERRIDES")); } + +// ─── W2: the two SERVER entry points the document did not reach (#1135) ────── +// +// The cases above all drive the text-generation path. `--offload-config` was +// parsed inside that path, AFTER the architecture branch, so the pooling and +// transcription-only branches built their engine parameters without it and said +// nothing. That dropped both halves of the document — vLLM's mirrored +// `uva`/`prefetch` weight offload and the `vllm_cpp` residency extension — and had +// done so since before the extension existed. +// +// The parse now happens ONCE, ahead of the branch. These cases drive the two +// branches through the REAL `VllmServerMain`, which is the only thing that can +// tell "the pooling branch takes the document" from "the loader takes a document +// somebody handed it". +// +// A MODEL DIRECTORY IS NEEDED HERE, unlike above: the branch is chosen by reading +// `architectures` out of `config.json`, so a nonexistent path takes neither +// branch. The directory holds that one file and nothing else, so the load still +// fails immediately after the install — which is what keeps the install +// observable without a checkpoint. + +namespace { + +// A directory holding one `config.json` with the given `architectures`. Created +// in the PARENT and left in place for the re-exec'd child to read. The name has +// no spaces, because the child splits `VLLM_TEST_SERVE_ARGS` on them. +std::string MakeArchDir(const char* arch) { + std::string tmpl = std::filesystem::temp_directory_path().string() + + "/vllm-cpp-serve-arch-XXXXXX"; + std::vector buf(tmpl.begin(), tmpl.end()); + buf.push_back('\0'); + const char* made = ::mkdtemp(buf.data()); + REQUIRE(made != nullptr); + const std::string dir(made); + std::ofstream cfg(dir + "/config.json"); + REQUIRE(cfg.good()); + cfg << R"({"architectures":[")" << arch << R"("],"model_type":"llama"})"; + cfg.close(); + return dir; +} + +// The pooling architecture registered in this tree: `REGISTER_VLLM_MODEL` names +// it in `llama_embedding_registry.cpp`, whose `kLlamaEmbeddingInfo` sets +// `is_pooling_model = true`. +constexpr const char* kPoolingArch = "LlamaModel"; +// The transcription-only architecture: `REGISTER_VLLM_MODEL` names it in +// `parakeet_registry.cpp`, whose `kParakeetInfo` sets +// `supports_transcription_only = true`. +constexpr const char* kTranscriptionArch = "ParakeetForCTC"; + +constexpr const char* kPoolingBanner = "server: pooling (embedding) model"; +constexpr const char* kTranscriptionBanner = + "server: transcription-only model"; + +} // namespace + +TEST_CASE("serve: the POOLING path installs the residency document") { + const std::string dir = MakeArchDir(kPoolingArch); + const ChildRun run = RunServer( + "--model " + dir + + R"( --offload-config {"vllm_cpp":{"mmap":{"enabled":true,"prefault":false},"device_fit":{"weight_budget_bytes":4096}}})"); + INFO("child output:\n" << run.output); + + // It really is the pooling branch, and not the text path taking the document + // as it always did. Without this the case would pass on a build where the + // architecture peek failed and everything fell through to the text path. + CHECK(Contains(run.output, kPoolingBanner)); + + // ...and that branch's `EngineParams` carried the document to the loader. + CHECK(Contains(run.output, kInstallLine)); + CHECK(Contains(run.output, "mmap=on")); + CHECK(Contains(run.output, "prefault=off")); + CHECK(Contains(run.output, "device_weight_budget_bytes=4096")); + CHECK(run.status == 0); +} + +TEST_CASE("serve: the POOLING path takes the MIRRORED half too, and refuses a typo in it") { + // #1135 is not specific to the extension: `uva`/`prefetch` was dropped on this + // branch as well, and for longer. The proof that it arrived is the line + // `CreateWeightOffloader` prints when it CONSTRUCTS a backend + // (`model_loader.cpp`, the `choice.offloader->moves_weights()` arm). That line + // is printed only for a document that selects one, so it cannot appear for an + // `EngineParams` whose `offload_config` is unset — which is exactly what this + // branch used to build. It is not the totality REFUSAL + // (`RefuseUnsupportedWeightOffload`): that one runs after `LoadHfConfig`, which + // this one-key `config.json` does not get past. + const std::string kOffloaderInstalled = + "engine: offload_config installed UvaWeightOffloader"; + const std::string dir = MakeArchDir(kPoolingArch); + const ChildRun run = + RunServer("--model " + dir + + R"( --offload-config {"offload_backend":"uva","uva":{"cpu_offload_gb":10}})"); + INFO("child output:\n" << run.output); + CHECK(Contains(run.output, kPoolingBanner)); + CHECK(Contains(run.output, kOffloaderInstalled)); + // No `vllm_cpp` key, so no residency install: the two halves stay separate. + CHECK_FALSE(Contains(run.output, kInstallLine)); + CHECK(run.status == 0); + + // The negative control for the line above: without the flag the same branch on + // the same directory does not print it, so the assertion is the document being + // seen rather than a line that always prints. + const ChildRun bare = RunServer("--model " + dir); + INFO("child output:\n" << bare.output); + CHECK(Contains(bare.output, kPoolingBanner)); + CHECK_FALSE(Contains(bare.output, kOffloaderInstalled)); + + // A typo in the mirrored half aborts at startup on this branch too, because + // the parse is now ahead of the branch rather than inside one of them. It + // aborts BEFORE the branch is even chosen, which is why the pooling banner + // does not print. + const ChildRun typo = + RunServer("--model " + dir + + R"( --offload-config {"uva":{"cpu_offload_GB":10}})"); + INFO("child output:\n" << typo.output); + CHECK(Contains(typo.output, "unknown key \"uva.cpu_offload_GB\"")); + CHECK_FALSE(Contains(typo.output, kPoolingBanner)); +} + +TEST_CASE("serve: the POOLING path is unchanged without the flag") { + // The inertness control for the two cases above: the branch is chosen the same + // way and prints nothing new when no document is passed. + const std::string dir = MakeArchDir(kPoolingArch); + const ChildRun run = RunServer("--model " + dir); + INFO("child output:\n" << run.output); + CHECK(Contains(run.output, kPoolingBanner)); + CHECK_FALSE(Contains(run.output, kInstallLine)); + CHECK(run.status == 0); +} + +TEST_CASE("serve: the TRANSCRIPTION-only path REFUSES the document, naming what is missing") { + // The decided arm of #1135. `ParakeetTranscriber::FromDir` builds no + // `EngineParams` and calls no `LoadedEngine`, so no field of either half has a + // reader on this path. AGENTS.md: refuse an unimplemented arm with a message + // naming the missing part. Accepting and warning would leave a server running + // while it holds a placement instruction it does not follow, which is the + // failure this issue was filed about. + const std::string dir = MakeArchDir(kTranscriptionArch); + const ChildRun run = + RunServer("--model " + dir + + R"( --offload-config {"vllm_cpp":{"mmap":{"enabled":true}}})"); + INFO("child output:\n" << run.output); + + CHECK(Contains(run.output, "server: fatal:")); + CHECK(Contains(run.output, "--offload-config is not supported on a " + "transcription-only model")); + CHECK(Contains(run.output, "THE MISSING PART")); + CHECK(Contains(run.output, "ParakeetTranscriber")); + // It names the issue that owns the wiring, so a reader of the message can find + // the record rather than concluding the capability was forgotten. + CHECK(Contains(run.output, "#1195")); + // Nothing was installed, which is the whole point: the alternative shape + // accepts the document and ignores it. + CHECK_FALSE(Contains(run.output, kInstallLine)); + CHECK(run.status == 0); + + // The MIRRORED half is refused on the same terms. It was dropped here for + // longer than the extension has existed, so a fix that covered only the + // extension would leave the older half of the same bug in place. + const ChildRun mirrored = + RunServer("--model " + dir + + R"( --offload-config {"uva":{"cpu_offload_gb":10}})"); + INFO("child output:\n" << mirrored.output); + CHECK(Contains(mirrored.output, + "--offload-config is not supported on a transcription-only " + "model")); +} + +TEST_CASE("serve: the TRANSCRIPTION-only path is unchanged without the flag") { + // The control that keeps the refusal above a REFUSAL OF THE FLAG rather than a + // refusal of the path. Without a document the branch is entered exactly as + // before and fails later, on the checkpoint this directory does not have. + const std::string dir = MakeArchDir(kTranscriptionArch); + const ChildRun run = RunServer("--model " + dir); + INFO("child output:\n" << run.output); + CHECK(Contains(run.output, kTranscriptionBanner)); + CHECK_FALSE(Contains(run.output, "--offload-config is not supported")); + CHECK(run.status == 0); +} diff --git a/tests/vllm/entrypoints/test_cli_offload_config.cpp b/tests/vllm/entrypoints/test_cli_offload_config.cpp new file mode 100644 index 000000000..17268d63c --- /dev/null +++ b/tests/vllm/entrypoints/test_cli_offload_config.cpp @@ -0,0 +1,186 @@ +// `ENG-RESIDENCY-CONFIG` W2, issue #1135 — does `vllm-cli --offload-config` +// actually reach the loader? +// +// WHY THIS RUNS A BINARY. `tests/vllm/entrypoints/test_weight_residency_reach.cpp` +// already drives `vllm_engine_load` with the same JSON string and asserts the +// install, so the ABI is gated. What was not gated, and what #1135 is about, is +// the FLAG: `vllm-cli` had no `--offload-config` at all, and a test that called +// the ABI from here would have passed on the day the flag was missing. +// .agents/reachability.md's rule is the same one: enter through the production +// entry point, which for this row's third entry point is the command line of a +// built executable. +// +// THE MODEL DIRECTORY IS DELIBERATELY NONEXISTENT. `LoadedEngine::FromModelDir` +// installs the residency document in its first statement block, ahead of every +// path and weight operation, so a load that fails on a missing checkpoint still +// runs the install and prints its line. That is what makes the chain observable +// without a checkpoint. +// +// THE MUTATION this file exists for: delete the `mp.offload_config = ...` +// assignment in `examples/cli/main.cpp` and CASE 1 goes red while every other +// residency suite stays green. +#include + +#include + +#include +#include +#include +#include + +namespace { + +#ifndef VLLM_CLI_BINARY +#define VLLM_CLI_BINARY "" +#endif + +constexpr const char* kCliBinary = VLLM_CLI_BINARY; + +// `cmake -DVLLM_CPP_BUILD_TESTS=ON -DVLLM_CPP_BUILD_EXAMPLES=OFF` is a legal +// configuration, and there `VLLM_CLI_BINARY` is the empty string above: there is +// no `vllm-cli` to run and this gate cannot be run at all. +// +// It must not FAIL there, and it must not PASS there either. A `REQUIRE` on the +// define reported four failed cases on a build that simply did not include the +// binary — a verdict about the configuration wearing the shape of a verdict about +// the code — and returning early would have printed doctest's +// `assertions: 0 ... Status: SUCCESS!` banner, which is the same trap issue #463 +// files. So the process exits with CTest's SKIP_RETURN_CODE (77, registered on +// every test by `vllm_cpp_add_test` in tests/CMakeLists.txt) and `ctest` reports +// **Skipped**, which is the true result. Exiting rather than skipping one case is +// right because the define is a build-time constant: if it is empty, every case in +// this file is unrunnable, so the first one to ask ends the process for all of them. +[[noreturn]] void SkipGate() { + std::fprintf(stderr, + "\n*** GATE NOT RUN — SKIPPED (exit 77), this is NOT a pass ***\n" + "*** test_cli_offload_config: built without VLLM_CPP_BUILD_EXAMPLES," + " so there is no vllm-cli binary to run\n\n"); + std::fflush(stderr); + std::exit(77); +} + +// Called first in every case. Named rather than inlined so that a case added later +// cannot get the polarity wrong by writing a bare `REQUIRE`. +void RequireCliBinary() { + if (std::string(kCliBinary).empty()) SkipGate(); +} + +// Combined stdout and stderr, plus the child's exit status. `vllm-cli` writes +// its own progress to stderr and the engine writes the install line there too, +// so both streams are read together. +struct CliRun { + std::string output; + int status = -1; +}; + +CliRun RunCli(const std::string& args) { + const std::string cmd = std::string(kCliBinary) + " " + args + " 2>&1"; + CliRun run; + FILE* pipe = ::popen(cmd.c_str(), "r"); + REQUIRE(pipe != nullptr); + std::array buf{}; + while (std::fgets(buf.data(), static_cast(buf.size()), pipe) != nullptr) { + run.output += buf.data(); + } + const int closed = ::pclose(pipe); + REQUIRE(closed != -1); + run.status = WIFEXITED(closed) ? WEXITSTATUS(closed) : -1; + return run; +} + +bool Contains(const std::string& haystack, const std::string& needle) { + return haystack.find(needle) != std::string::npos; +} + +constexpr const char* kMissingModel = + "--model /nonexistent/vllm-cpp/cli-offload-config --prompt hi " + "--max-tokens 1"; + +// Printed by `LoadedEngine::FromModelDir` when it installs a residency document. +constexpr const char* kInstallLine = "engine: weight residency"; + +} // namespace + +TEST_CASE("vllm-cli: --offload-config's vllm_cpp half reaches the loader") { + RequireCliBinary(); + const CliRun run = RunCli( + std::string(kMissingModel) + + R"( --offload-config '{"vllm_cpp":{"mmap":{"enabled":true,"prefault":false},)" + R"("device_fit":{"weight_budget_bytes":4096}}}')"); + INFO("vllm-cli output:\n" << run.output); + + // The flag was accepted rather than reported as unknown. `vllm-cli` prints its + // usage and exits 2 on an unrecognised argument, so this is the assertion that + // separates "the flag exists" from "the flag was parsed as a model path". + CHECK_FALSE(Contains(run.output, "usage:")); + CHECK(Contains(run.output, "vllm-cli: loading model from")); + + // ...and the document reached the install, naming the fields the operator set. + CHECK(Contains(run.output, kInstallLine)); + CHECK(Contains(run.output, "mmap=on")); + CHECK(Contains(run.output, "prefault=off")); + CHECK(Contains(run.output, "device_weight_budget_bytes=4096")); + + // The load then failed on the deliberately missing checkpoint, which is what + // makes the install observable without one. + CHECK(Contains(run.output, "model load failed")); + CHECK(run.status == 1); +} + +TEST_CASE("vllm-cli: the MIRRORED half of the document reaches the loader too") { + // #1135 is not specific to the `vllm_cpp` extension: vLLM's own `uva` half had + // no way to reach this entry point either. The proof that it arrived is the + // line `CreateWeightOffloader` prints when it CONSTRUCTS a backend + // (`model_loader.cpp`, the `choice.offloader->moves_weights()` arm), which is + // printed only for a document that selects one. + RequireCliBinary(); + const std::string kOffloaderInstalled = + "engine: offload_config installed UvaWeightOffloader"; + const CliRun run = + RunCli(std::string("--model /nonexistent/vllm-cpp/cli-offload-uva ") + + "--prompt hi --max-tokens 1 " + + R"(--offload-config '{"offload_backend":"uva","uva":{"cpu_offload_gb":10}}')"); + INFO("vllm-cli output:\n" << run.output); + + CHECK_FALSE(Contains(run.output, "usage:")); + CHECK(Contains(run.output, kOffloaderInstalled)); + // A `vllm_cpp`-free document installs no residency, so the line must NOT + // print. Without this the case would also pass on an implementation that + // routed the mirrored half into the extension. + CHECK_FALSE(Contains(run.output, kInstallLine)); + CHECK(Contains(run.output, "model load failed")); + CHECK(run.status == 1); + + // The negative control: the same run without the flag prints neither line. + const CliRun bare = RunCli(kMissingModel); + INFO("vllm-cli output:\n" << bare.output); + CHECK_FALSE(Contains(bare.output, kOffloaderInstalled)); +} + +TEST_CASE("vllm-cli: a mistyped key in the document fails the load, naming the key") { + // The extension parser closes the whole document, and `vllm-cli` inherits that + // by parsing nothing itself: the refusal comes out of `vllm_engine_load`. A + // typo that quietly disabled this tier would be met as an out-of-memory kill + // minutes later instead of as a message. + RequireCliBinary(); + const CliRun run = + RunCli(std::string(kMissingModel) + + R"( --offload-config '{"vllm_cpp":{"device_fitt":{"weight_budget_bytes":1}}}')"); + INFO("vllm-cli output:\n" << run.output); + + CHECK(Contains(run.output, "unknown key \"vllm_cpp.device_fitt\"")); + CHECK(Contains(run.output, "expected one of: mmap expert_stream device_fit")); + CHECK_FALSE(Contains(run.output, kInstallLine)); + CHECK(run.status == 1); +} + +TEST_CASE("vllm-cli: no --offload-config prints no residency line at all") { + // The inertness case. Nearly every run passes no such flag, and those runs must + // print nothing new and install nothing. + RequireCliBinary(); + const CliRun run = RunCli(kMissingModel); + INFO("vllm-cli output:\n" << run.output); + CHECK(Contains(run.output, "vllm-cli: loading model from")); + CHECK_FALSE(Contains(run.output, kInstallLine)); + CHECK(run.status == 1); +} diff --git a/tests/vllm/entrypoints/test_gguf_device_fit_reach.cpp b/tests/vllm/entrypoints/test_gguf_device_fit_reach.cpp index 31489f8f2..bcb204181 100644 --- a/tests/vllm/entrypoints/test_gguf_device_fit_reach.cpp +++ b/tests/vllm/entrypoints/test_gguf_device_fit_reach.cpp @@ -28,6 +28,7 @@ #include "support/test_env.h" #include "vllm/config/device.h" +#include "vllm/config/weight_residency.h" #include "vllm/entrypoints/model_loader.h" #include "vllm/gguf_builder.h" #include "vllm/platforms/interface.h" @@ -316,3 +317,132 @@ TEST_CASE("device fit: an explicit CPU load is never refused, at any budget") { CHECK(message.find("cannot serve this GGUF") == std::string::npos); CHECK(message.find("tokenizer: GGUF missing kv") != std::string::npos); } + +// --- The budget as a CONFIG KEY (#1127), through the same loader -------------- +// +// The cases above move the budget with `VT_DEVICE_WEIGHT_BUDGET_BYTES`. The three +// below move it with `--offload-config`'s `vllm_cpp.device_fit.weight_budget_bytes`, +// arriving as `EngineParams::weight_residency`. The first two set NO variable at +// all; the third sets one DELIBERATELY, because its subject is the precedence +// between the two inputs rather than the config tier on its own. +// +// This is the reachability half of #1127 and not a second unit test of the +// resolver: `test_weight_residency_config` already builds the config by hand and +// calls `ResolveDeviceWeightBudgetBytes`, which proves the rule and says nothing +// about whether a document reaches it. The chain these cases traverse is +// `parse_weight_residency_extension_json` -> `EngineParams::weight_residency` -> +// `SetWeightResidencyConfig` (the install block at the top of `FromModelDir`) -> +// `DeviceWeightBudgetBytes` (the fit check, later in the SAME call). Measured: +// deleting the install call site in `LoadedEngine::FromModelDir`, and separately +// deleting the delegation inside `DeviceWeightBudgetBytes`, each turn this suite +// RED. +// +// The install and the fit check being in one function is what makes this +// observable without a checkpoint: the install runs before any path or weight +// operation, and the check runs before the tokenizer. + +TEST_CASE("device fit: the budget arrives as a CONFIG KEY, with no variable set") { + RegisterFakeStagingPlatform(); + TempFile f(BuildSyntheticMoeGguf()); + vllm::ResetWeightResidencyConfigForTesting(); + vllm_test::UnsetEnv("VT_DEVICE_WEIGHT_BUDGET_BYTES"); + + vllm::entrypoints::EngineParams params; + params.device = vllm::Device::kNamedPlatform; + // Parsed from the SAME string the `--offload-config` flag carries, rather than + // assembled field by field: a case that set the struct member directly would + // still pass if the parser never learned the key. + params.weight_residency = vllm::parse_weight_residency_extension_json( + R"({"vllm_cpp":{"device_fit":{"weight_budget_bytes":8191}}})"); + REQUIRE(params.weight_residency->device_weight_budget_bytes.has_value()); + + std::string message; + try { + (void)vllm::entrypoints::LoadedEngine::FromModelDir(f.path(), params); + } catch (const std::exception& e) { + message = e.what(); + } + vllm::ResetWeightResidencyConfigForTesting(); + + REQUIRE_FALSE(message.empty()); + CAPTURE(message); + // One byte under the footprint, exactly as the environment case is, so a + // comparison against the wrong quantity cannot pass by accident. + CHECK(message.find("cannot serve this GGUF") != std::string::npos); + CHECK(message.find(std::to_string(kStagedLowerBound)) != std::string::npos); + CHECK(message.find(std::to_string(kStagedLowerBound - 1)) != std::string::npos); + // The refusal names the config form as a way out, not only the variable, so an + // operator who set the budget with a document is told how to raise it with one. + CHECK(message.find("weight_budget_bytes") != std::string::npos); + CHECK(message.find("tokenizer") == std::string::npos); +} + +TEST_CASE("device fit: a config budget of ZERO suppresses the refusal") { + RegisterFakeStagingPlatform(); + TempFile f(BuildSyntheticMoeGguf()); + vllm::ResetWeightResidencyConfigForTesting(); + vllm_test::UnsetEnv("VT_DEVICE_WEIGHT_BUDGET_BYTES"); + + vllm::entrypoints::EngineParams params; + params.device = vllm::Device::kNamedPlatform; + // The platform's own probe is 0 in this fixture, so a budget of 0 is NOT what + // distinguishes this case from the default one. What it proves is that a + // configured zero reaches the check AS a zero rather than being dropped as a + // falsy value on the way, which is the one direction the merge and the resolver + // could each have got wrong. The positive control is the case above: the same + // file and platform refuse when the configured budget is 8191. + params.weight_residency = vllm::parse_weight_residency_extension_json( + R"({"vllm_cpp":{"device_fit":{"weight_budget_bytes":0}}})"); + + std::string message; + try { + (void)vllm::entrypoints::LoadedEngine::FromModelDir(f.path(), params); + } catch (const std::exception& e) { + message = e.what(); + } + // Read what the loader installed BEFORE clearing it: this is the assertion that + // the document reached the process-global rather than being carried past it. + const vllm::WeightResidencyConfig installed = + vllm::ActiveWeightResidencyConfig(); + vllm::ResetWeightResidencyConfigForTesting(); + + REQUIRE(installed.device_weight_budget_bytes.has_value()); + CHECK(*installed.device_weight_budget_bytes == 0); + + REQUIRE_FALSE(message.empty()); + CAPTURE(message); + CHECK(message.find("cannot serve this GGUF") == std::string::npos); + // The LATER error, asserted positively: without it, "no refusal" would also be + // true of a load that died earlier for an unrelated reason. + CHECK(message.find("tokenizer: GGUF missing kv") != std::string::npos); +} + +TEST_CASE("device fit: the VARIABLE beats the config key, through the loader") { + RegisterFakeStagingPlatform(); + TempFile f(BuildSyntheticMoeGguf()); + vllm::ResetWeightResidencyConfigForTesting(); + + vllm::entrypoints::EngineParams params; + params.device = vllm::Device::kNamedPlatform; + // The document suppresses the refusal; the variable puts a refusing budget + // back. The precedence exists so a benchmark arm is switchable without a + // restart, and this is that direction: the variable can turn a configured + // suppression back ON. + params.weight_residency = vllm::parse_weight_residency_extension_json( + R"({"vllm_cpp":{"device_fit":{"weight_budget_bytes":0}}})"); + vllm_test::SetEnv("VT_DEVICE_WEIGHT_BUDGET_BYTES", + std::to_string(kStagedLowerBound - 1)); + std::string message; + try { + (void)vllm::entrypoints::LoadedEngine::FromModelDir(f.path(), params); + } catch (const std::exception& e) { + message = e.what(); + } + vllm_test::UnsetEnv("VT_DEVICE_WEIGHT_BUDGET_BYTES"); + vllm::ResetWeightResidencyConfigForTesting(); + + REQUIRE_FALSE(message.empty()); + CAPTURE(message); + CHECK(message.find("cannot serve this GGUF") != std::string::npos); + CHECK(message.find(std::to_string(kStagedLowerBound - 1)) != std::string::npos); +}