diff --git a/.agents/issue-index.md b/.agents/issue-index.md index bdbc76def..99ee3b640 100644 --- a/.agents/issue-index.md +++ b/.agents/issue-index.md @@ -331,6 +331,13 @@ rather than merged. `scripts/check-agent-record.py` gates both. | [#1117](https://github.com/mudler/vllm.cpp/issues/1117) | `LTX25-A2VID-RECIPE` | `A2VidPipelineTwoStage` (`a2vid_two_stage.py:53` @ `fd4ded7f`) has no recipe row, so `pipeline_kind = a2vid_two_stage` gets the generic table refusal (`src/vllm/model_executor/models/ltx2_pipeline.cpp:1328-1332`) naming the pair rather than the missing machinery. [#922](https://github.com/mudler/vllm.cpp/issues/922) is CLOSED and closed the audio CONDITIONING, not the recipe: a supplied take rides `distilled_two_stage`, which [`ltx25-a2v-audio-input.md`](specs/ltx25-a2v-audio-input.md):438-446 already records as a different trajectory. Measured at `daeff67f2`: `git grep -n '"a2vid' -- src include tests docs examples` returns TWO hits, both upstream anchors inside `Fail`-message assertions (`tests/vllm/multimodal/test_ltx2_video.cpp:4363,:4427`), against a control of 4 for `"one_stage"` in `include/` alone. Four differences from the recipe it rides, each read at the pin: stage 1 is CFG/STG/modality-guided and caller-configured (`:230-240`, fed from `utils/args.py:947-1006`, `--a2v-guidance-scale` defaulting to `video_guider.modality_scale` = 3.0 at `utils/constants.py:54,:64`) where `distilled_two_stage` fixes `allow_guidance_override = false`; stage 1's schedule is scheduler-derived (`:225-227`) against our fixed `DistilledSigmas()`; stage 1 is plain Euler (`:229-258` passes no `stepper`, `utils/blocks.py:526-527`) against our `kEulerAncestral` on 2.5; and the AUDIO guider is the DEFAULT positive-only one (`:237-239`, `ltx-core components/guiders.py:200-210`) rather than the params table's cfg-7.0 row. Two non-schedule facts that must not be guessed: `--audio-path` is `required=True` (`:312-317`), and the distilled LoRA rides stage 2 ALONE (`:114` against `:107`) with `--distilled-lora` `required=True` (`utils/args.py:1140-1153`). Unblocked by `Ltx2GuidedDenoise` landing at `daeff67f2` (#1092/#1102), which [`ltx25-guided-video.md`](specs/ltx25-guided-video.md) `## Owed` names this arm against. Spec [`ltx25-a2vid-recipe.md`](specs/ltx25-a2vid-recipe.md) | enhancement | | [#1118](https://github.com/mudler/vllm.cpp/issues/1118) | `LTX25-A2VID-RECIPE` | LoRA adapters fuse ONCE, at load, into ONE weight set: `src/vllm/multimodal/ltx2_video.cpp:816-820` is the only `dit_options.loras.push_back` in the tree and it runs under `if (!lora_path.empty())`, so every phase of every recipe sees the same fused `im.dit.weights`. Three upstream pipelines build TWO `DiffusionStage`s from the same checkpoint with DIFFERENT adapter sets at `fd4ded7f`: `a2vid_two_stage.py:107` against `:114-119`, `ti2vid_two_stages.py:140` against `:151`, and `ti2vid_two_stages_hq.py:154,:165` at two separate strengths (`:92-101`). `distilled.py:131` builds ONE stage set, which is why `distilled_two_stage`, `dfr` and `retake` have never needed this and no gate has ever asked. Consequence for the arm landing with #1117: that row mirrors `--distilled-lora required=True` (`utils/args.py:1140-1153`) by refusing an `a2vid_two_stage` load with no `lora_path`, so the 3-step stage-2 schedule cannot run on a checkpoint with no distilled adapter — but it CANNOT mirror upstream leaving stage 1 WITHOUT that adapter, so stage 1's guided schedule runs against base + distilled LoRA where upstream runs it against the base alone. That divergence RENDERS, and the PIXELS it renders are not upstream's: it moves the trajectory, so the frames themselves differ, while the frame count, the shapes, the sample rate and the errors are all exactly what they were — nothing in the SHAPE of the result says anything is wrong, which is why it is filed rather than left in a comment. It is not undetectable, and saying so would be the more damaging error: the instrument that WOULD see it is a real-weights comparison against upstream's own render on the same checkpoint, take and seed, upstream's stage 1 on the base weights against ours on base + distilled. Two fix shapes, neither chosen: a second `Ltx2DitWeights` per adapter set (what upstream pays, two `from_checkpoint` calls at `a2vid_two_stage.py:103,:115`) or unfused runtime LoRA selectable per phase; the first doubles resident DiT weights and the second changes the GEMM path. Bounds #1093 and #921, which need the same seam. Listed under `## Owed` in [`ltx25-a2vid-recipe.md`](specs/ltx25-a2vid-recipe.md) | enhancement | | [#1134](https://github.com/mudler/vllm.cpp/issues/1134) | `BACKEND-ROCM` | The `VT_ATTN_DECODE_D128=1` ctest registration added by #767 cannot show the ROCm `d=128` decode arm REACHED the new kernel, and is empty off ROCm. `RegisteredDevices()` (`tests/vt/test_backend_cross_device.cpp:84-96`) enumerates `{kCUDA, kMETAL, kVULKAN, kXPU, kROCM}` and excludes `kCPU`, so on a CPU-only runner — which is what CI has — the new "Qwen3 geometry (bf16, GQA 2, head_dim 128)" case reports 1 test case, 0 assertions, exit 0, for BOTH registrations. On ROCm hardware the case's only backend assertion is `OpProviderStats::declines == 0`, and `OpProviderStats` counts at PROVIDER granularity, so it is identical with the flag set and unset; the NMSE bound passes on either kernel because the arm is correctness-complete. The two compose: there is no machine in this project on which the flag-ON registration distinguishes itself from the flag-OFF one. Disclosed in [`specs/rocm-decode-attn-d128.md`](specs/rocm-decode-attn-d128.md) §4, its `## Owed` section and its result banner, and §9 stop condition 2 is left OPEN rather than claimed discharged. Closing it needs a kernel-selection counter in `src/vt/rocm/rocm_paged_attn.hip` asserted to DIFFER between the two registrations; the CPU-runner half wants `kCPU` in `RegisteredDevices()` or a non-zero-assertion floor per [#463](https://github.com/mudler/vllm.cpp/issues/463). Not a duplicate of #463 (that is the unset-weights-env-var shape and does not describe the `declines` granularity half), #785 (a kernel that never LAUNCHES behind a dead `#if`, a code defect not a coverage one) or #900 (same family, LTX-2.5 subject) | bug | +| [#1123](https://github.com/mudler/vllm.cpp/issues/1123) | `ENG-EXPERT-STREAM` | `Qwen3.8-2.4T-A95B UD-Q1_0` (369.96 GiB) reached a serving state on `--device cuda` on a 119.631 GiB GB10 after 26 minutes and then died on the FIRST forward with `vt cuda: cudaMalloc: out of memory`, in the EngineCore busy loop rather than the loader, while the same checkpoint and binary serve on `--device cpu` (TTFT 667.0 s, 44.2 s/token, coherent). The log line could not name the allocation because `CudaBackend::Alloc` throws `"vt cuda: " + "cudaMalloc" + ": " + cudaGetErrorString(err)` and DISCARDS `bytes` (`src/vt/cuda/cuda_backend.cu:48-52,75-81`). Named by reading: `ResidentWeight` uploads the WHOLE stacked `[E*N,K]` keep-quant expert tower, `d.b.Alloc(w.bytes.size())` at `src/vllm/model_executor/models/qwen3_5.cpp:1011`, and BOTH switch positions reach it — the default grouped path (`MoeBlock:6615,6616,6620` -> `KqGrouped:5694`) and `VT_MOE_EXPERT_STREAM=1`, which disables grouping and then takes `KqExpertSlice:5595` -> `KqResidentSlice:5112` because the slot arm is guarded by `is_cpu()` (`:5578`). `BuildMoeMarlinResident` is NOT on this path (`MoeBlock:6555` needs `expert_*_fp4`, a GGUF populates `expert_*_kq`). Sized by re-censusing both GGUF tensor tables at revision `567d3e6ac26c5474b18311e619c04350fb9a5556` over all ten shards by HTTP range request, 1702 records parsed against 1702 declared in `split.tensors.count`: one IQ1_XXXS tower is **1,275,068,416 B (1.1875 GiB)**, the three Q2_K MTP-block towers are 2,818,572,288 B each, all 279 total **360,374,599,680 B = 335.62 GiB**, and `1,275,068,416 / 512 = 2,490,368` matches the W4 banner's `slot_bytes` exactly. Budget measured with the instrument that works where `nvidia-smi` answers `[N/A], [N/A], [N/A]`: `cudaMemGetInfo` on `dgx:gpu0` reports total `128,452,956,160` (119.631 GiB), free 113.677 GiB, `Integrated=1`. The load survives because a borrowed tower costs ZERO anonymous bytes; staging converts each into a real allocation, exhausting the pool after roughly 48 towers, partway through layer 16 of 93. FIXED here by a load-time refusal keyed on the measured condition (`needs_weight_staging` AND a known budget AND a LOWER-bound footprint above it), never on "CUDA + GGUF" and never on an architecture name, so a GGUF that fits still loads. The device-slot arm is NOT built and is [#1124](https://github.com/mudler/vllm.cpp/issues/1124) | bug | +| [#1124](https://github.com/mudler/vllm.cpp/issues/1124) | `ENG-EXPERT-STREAM` | `--device cuda` still cannot SERVE a larger-than-pool GGUF after [#1123](https://github.com/mudler/vllm.cpp/issues/1123); it refuses by name instead of dying mid-stream. The missing capability is a DEVICE expert slot store, and it is four pieces: `HostExpertSlotStore` is the only production `ExpertSlotStore` (`include/vllm/model_executor/host_expert_slot_store.h:28`, the only other subclass being a test double) while `include/vllm/model_executor/expert_streamer.h:8-9,30-31` claims "the production destination is a contiguous device-side slot array" and is FALSE today; the interface has no device-capable read, because `KqExpertSlice` reads back through `HostExpertSlotStore::Slot()`, the CONCRETE class (`qwen3_5.cpp:5258,5314`); the filler is `pread`-into-host, since `SlotForWrite` is handed straight to `::pread` (`expert_streamer.cpp:76-94`); and the consumer is device-gated by `is_cpu()` at `qwen3_5.cpp:5578`. Sized: 2790 slices per token at 2,490,368 B is 6.95 GB per token against a 119.631 GiB pool already holding the dense remainder. Deferred because W7 owns the pluggable backing store in the row's work breakdown, and the CPU arm's own decode bandwidth is still VOID (#912 F1 measured it with the step clock dead from token 3), so a device lane would be optimised against a number nobody has. Listed under `## Owed` in [`expert-streaming.md`](specs/expert-streaming.md) | gap | +| [#1126](https://github.com/mudler/vllm.cpp/issues/1126) | `ENG-EXPERT-STREAM` | `CudaBackend` never overrides `Backend::DeviceMemoryInfo`, and the seam's own comment says it does: `include/vt/backend.h:79-83` reads "ROCm/CUDA override with hipMemGetInfo/cudaMemGetInfo" while only `src/vt/rocm/rocm_backend.hip:338-345` does, and `cudaMemGetInfo` is called NOWHERE in the repository. The comment is corrected in prose by [#1123](https://github.com/mudler/vllm.cpp/issues/1123); the capability is this issue. The consequence is not only a comment: `Gemma4MoE` is the seam's only consumer, `FreeBytes` returns false on an absent probe (`src/vllm/model_executor/models/gemma4_moe.cpp:439-447`) and `MakeRoom` refuses on unknown by design (`:494-506`), so on EVERY CUDA device the device-expert LRU (`kMaxSlots = 24`, `kHeadroom = 1.5 GiB`) admits nothing and falls back to host H2D permanently, silently. That polarity is right for that call site; the defect is the missing probe. #1123 therefore probed `cudaMemGetInfo` in `CudaPlatform` (which already includes `` and already probes attributes at registration) and carried the total on `ResidencyPolicy`, touching nothing Gemma4 reads, because adding the override wakes another model's residency policy and that needs its own measurement. Measured on `dgx:gpu0`: total 128,452,956,160 (119.631 GiB), free 113.677 GiB, against `nvidia-smi` answering `[N/A], [N/A], [N/A]`. Listed under `## Owed` in [`expert-streaming.md`](specs/expert-streaming.md) | bug | +| [#1127](https://github.com/mudler/vllm.cpp/issues/1127) | `ENG-EXPERT-STREAM` | `VT_DEVICE_WEIGHT_BUDGET_BYTES`, added by [#1123](https://github.com/mudler/vllm.cpp/issues/1123) to override the probed device memory pool for the load-time fit refusal, should be a weight-residency CONFIG key rather than an environment variable, for the reason `ENG-RESIDENCY-CONFIG` gives for the five `VT_GGUF_*` / `VT_MOE_EXPERT_STREAM*` knobs it is converting. It was left as an environment variable ON PURPOSE: [#1110](https://github.com/mudler/vllm.cpp/issues/1110) / PR #1119 is in flight, adds exactly the `vllm_cpp` namespace inside `--offload-config` this key belongs in plus `include/vllm/config/weight_residency.h`, and touches the same `src/vllm/entrypoints/model_loader.cpp`, so landing a competing config surface first would create the conflict both changes then resolve. Closing it means a `device_weight_budget_bytes` key under that object, the loader reading `EngineParams::weight_residency` instead of `std::getenv`, and `docs/ENVIRONMENT.md` plus `docs/USAGE.md` following. Listed under `## Owed` in [`expert-streaming.md`](specs/expert-streaming.md) | gap | +| [#1136](https://github.com/mudler/vllm.cpp/issues/1136) | `ENG-EXPERT-STREAM` | The load-time GGUF fit bound of [#1123](https://github.com/mudler/vllm.cpp/issues/1123) can OVER-count, and therefore over-refuse. Per tensor it is `min(gguf_bytes, elems * model_dtype_bytes)`, a true lower bound on THAT tensor's staged size; the sum is not a lower bound on the load, because a tensor present in the file and never staged is a positive error. One such class exists on every default load: the MTP / `nextn` head is attached only under `params.speculative_config.has_value() && method == "mtp"` (`src/vllm/entrypoints/model_loader.cpp`, the `maybe_attach_mtp` GGUF arm), and the main model reads `block_count - nextn_predict_layers` blocks (`qwen3_5_gguf_weights.cpp:877-878`), so the head's blocks are outside its range. Measured on `unsloth/Qwen3.8-2.4T-A95B-GGUF` at `567d3e6ac26c5474b18311e619c04350fb9a5556`: block 92 is 20 tensors, **8,940,488,704 bytes (8.33 GiB) of 397,245,341,184**, so **2.2506 %**. A budget in `[what a default load stages, what the bound counts)` refuses a weight set that fits. The "under-count dominates" argument does NOT close this: the two errors are on different quantities and never cancel. NOT fixed, and the reason is the fix's own failure mode. Excluding those tensors means the bound taking a per-tensor staging POLICY as input, which is the caller's knowledge and not the file's, and an exclusion that is wrong under-counts toward zero — which restores exactly the 26-minute-load-then-`cudaMalloc`-OOM this row removed, on a device nobody on this fleet has to measure the change against. What IS done: `gguf_device_fit.h` states the direction, `tests/vllm/model_executor/test_gguf_device_fit.cpp` pins it executably with a counted-but-unstaged fixture and asserts both ends of the over-refusal window, `docs/USAGE.md` tells an operator that `VT_DEVICE_WEIGHT_BUDGET_BYTES` is the way out of it, and the spec's risk table records it. Also filed here rather than lost: the same review found `ResolveModelDeviceType` and `SelectQueueForModel` diverging on the AUTO arm (fixed in flow — the resolver now resolves through an attempted queue, pinned by two cases in `test_gguf_device_fit_reach.cpp`), the false "ROCm/CUDA override" comment in TWO places (both corrected in flow), and `CudaPlatform`'s policy assembly being reachable only in a CUDA build (extracted to `CudaResidencyPolicy` in `vllm/platforms/interface.h` and unit-tested on every host). Listed under `## Owed` in [`expert-streaming.md`](specs/expert-streaming.md) | bug | +| [#1139](https://github.com/mudler/vllm.cpp/issues/1139) | `KV-WARMUP-PROFILE` | All three upstream anchors on that row (`.agents/engine-matrix.md:112`) point at unrelated code at the current parity pin `555967922`, verified by reading the pinned tree: `vllm/v1/worker/gpu/model_runner.py:504` is inside a `DraftModelSpeculator.set_attn(...)` call, `:647` is a `torch.zeros(...)` argument in a `dummy_run=True` construction, and `vllm/v1/worker/gpu_worker.py:430` is a comment about `max_split_size_mb` inside `load_model`. The startup memory profile the row describes is `GPUWorker.determine_available_memory` (`gpu_worker.py:451-495`, `memory_profiling` at `:491-494` around `profile_run()` at `:495`) and `GPUModelRunner.profile_run` (`gpu/model_runner.py:682`); `model_memory_usage` is recorded AFTER the load at `gpu/model_runner.py:315`, which is why upstream never asks whether the weights will fit and why [#1123](https://github.com/mudler/vllm.cpp/issues/1123) has no upstream counterpart to mirror. Found while repairing [#1136](https://github.com/mudler/vllm.cpp/issues/1136): `gguf_device_fit.h` and `expert-streaming.md` had both COPIED the `:504,647` pair from this row, and both are corrected there, so this row is the surviving source. Filed and not fixed in flow because the fix is one cell in `.agents/engine-matrix.md`, which PR #1119 ([#1110](https://github.com/mudler/vllm.cpp/issues/1110)) is concurrently bumping alongside the hardcoded `ENGINE` count in `scripts/check-agent-record.py` — the record-lock hazard AGENTS.md names, and the reason the repairing session was told to leave both files alone. Most likely cause: correct at the previous `e24d1b24` pin and not reconciled when the pin advanced; whether other `INVENTORIED` rows citing `vllm/v1/worker/gpu/**` share the defect is a wider sweep than one cell | bug | +| [#1143](https://github.com/mudler/vllm.cpp/issues/1143) | `ENG-EXPERT-STREAM` | `src/vllm/entrypoints/model_loader.cpp` is cited by ABSOLUTE LINE NUMBER from **109 distinct sites across 45 files** — specs, matrices, docs and comments in other translation units — and the file is ~1640 lines that almost every engine and model row edits. Any edit near its top invalidates every citation below it, in files the editing change never opens. Measured on [#1136](https://github.com/mudler/vllm.cpp/issues/1136)'s repair, which inserts ~45 lines near line 100 (`AutoDeviceResolution` / `ResolveAutoDevice`, so the device resolver and the queue selector share one description): comparing the TEXT at every cited line between `e7d0a1f7c` and the repaired head gives **203 moved line references over 109 citing sites in 45 files, 10 unmoved**. What that number is NOT: a claim that 109 correct citations broke. Several were already stale — `.agents/model-matrix.md:197` cites `model_loader.cpp:184-223` as the "live loader" while line 184 at `e7d0a1f7c` is `static const bool once = [] {` inside the `VT_LOAD_STATS` helper. The finding is that the surface cannot survive an ordinary edit and that nobody can currently tell the two cases apart. NOT swept there, for two reasons: it is 109 sites in specs owned by other rows, and rewriting them all from the current tree would launder pre-existing debt into a clean-looking record. What #1136 DID fix is the two anchors it authored itself, plus adding them to its own anchor verifier so they could not go stale inside their own pull request — the defect the round before hit with `platforms/cuda.cpp:67`. Fix candidates, none chosen: cite SYMBOLS not lines across file boundaries (the only one that removes the class); an anchor gate holding each `file:line` against an expected substring (a checker change, and its expectation table is itself a shared-file lock); or split the 1640-line file, whose size is what makes the blast radius large. Listed under `## Owed` in [`expert-streaming.md`](specs/expert-streaming.md) until a row claims it | bug | | [#1146](https://github.com/mudler/vllm.cpp/issues/1146) | `ENV-LEASE-RUNTIME-STAGING` | #1129 is closed and its recorded cause is FALSIFIED by the worker image: it says a leased worker "cannot start Python" and lists `python3`, `pip`, `gcc`, `curl` and `git` as ABSENT, measured in one `dgx:gpu0` probe (job `ff28ada1-0cd3-4867-bf9b-f67050d0608b`). Measured 2026-08-17 on `thor:gpu0` through five `rc run` jobs (`6f4bdb03`, `9c0ebeac`, `8beba132`, `f60d945f`, `63c60a90`), that worker runs as `uid=0(root)` with `/usr/bin/gcc`, `/usr/bin/python3` and a working `apt-get`, and a relocated CUDA runtime staged on `/workspace` imports torch 2.13.0+cu130, reports `cuda available = True` on `NVIDIA Thor` capability (11,0), runs a bf16 matmul, and compiles and executes a Triton kernel (`TRITON_JIT_OK = 4096.0 PASS`, `PROBE5_RC=0`). So the image is provisionable per job, which is none of the three fleet-side fixes #1129 names as the only ways forward. Four walls stand between a staged runtime and a running one: the `pip --target` must run FROM the worker because the submitting host is `x86_64` and the workers are `aarch64`. `Python.h` is absent until `apt-get install python3-dev`. The NAS mount presents `file_mode=0664` so Triton cannot execute its own `ptxas-blackwell`, and `TRITON_PTXAS_PATH` does NOT fix it because it redirects only the plain `ptxas`. So `PYTHONPATH` is ORDERED, `/tmp/tp` before the NAS tree. SCOPE, and it is the point of the row: this is `thor:gpu0` at capability (11,0) ONLY, the GB10 is `sm_121a` and UNMEASURED, only `torch`, `triton` and `numpy` are staged so the pinned vLLM oracle is still NOT shown to run and #1129's consequence for the oracle-dependent rows is NARROWED rather than closed, the `+cu130` versus `release 12.8, V12.8.93` skew is recorded as observed and not adjudicated, and a prebuilt wheel does NOT shorten the route because an aarch64 vLLM wheel exists in general while our pin is neither among the wheels published for the one nightly commit nor on PyPI (the per-commit 404s prove nothing, because that URL scheme was never confirmed against a known-good case). Recipe, job IDs and staged-script sha256 values in [`lease-runtime-staging.md`](specs/lease-runtime-staging.md) | verification | | [#1033](https://github.com/mudler/vllm.cpp/issues/1033) | `GATE-ISSUE-INDEX-TABLE-SHAPE` | `check_table_shapes` (`scripts/check-agent-record.py:1292`) never ran on `.agents/issue-index.md`, so a malformed row there was invisible to every gate. The function already counted unescaped pipes per table line with exactly the right regex; its call site (`:1527-1530`) passed `roadmap_v1.md`, `coordination.md`, `*MATRIX_PATHS` and `*spec_paths`, and simply did not pass this path. Nothing else in the tree counts this file's cells, which made the index the ONLY markdown table in the record set with no shape gate. It is also the one record surface every change must write, with rows long enough to hide a stray pipe in a code span. Arming it reds exactly one row on `origin/main` at `100026481`: a pipe histogram over the index's 289 table lines reads `{5: 288, 9: 1}`, and the outlier is line 279, the [#1003](https://github.com/mudler/vllm.cpp/issues/1003) `ORACLE-LLAMACPP-REPIN-STOCK` row that arrived with `283c7e492` ([#1051](https://github.com/mudler/vllm.cpp/pull/1051)), carrying four unescaped pipes inside code spans at columns 2705, 3106, 3115 and 3338. FIXED IN FLOW: the path is added, the four pipes are escaped, and three cases in `tests/scripts/test_agent_record.py` hold it — one capturing the paths `main()` really hands the gate, one running it on the shipped file, one mutating a copy so the instrument is proven to fire. Repairing the row EDITS an append-only file, so `check-issue-index-append-only.py` is red on the branch and the exception is argued in the commit body, as `ff264cb82` ([#1025](https://github.com/mudler/vllm.cpp/pull/1025)) argued the same one: appending a corrected copy would leave the broken row in place and add a duplicate key, so the file only becomes well-formed by editing it where it sits. TWO PREMISES OF THE REPORT MEASURED FALSE and are recorded in the spec rather than quietly dropped: the checker does NOT stop at the first finding (one `errors` list, three findings in one run, exit 1 once), and the four pipes are NOT in a `git diff` piped into `grep` — that span does not exist in the row. Spec [`gate-issue-index-table-shape.md`](specs/gate-issue-index-table-shape.md) | bug | | [#1144](https://github.com/mudler/vllm.cpp/issues/1144) | — | `Res2sTwoStageRecipe` is this tree's port of `TI2VidTwoStagesHQPipeline` (`ti2vid_two_stages_hq.py:59` @ `fd4ded7f`), the one in-scope pipeline that runs the SAME adapter on BOTH stages at DIFFERENT strengths: it builds `distilled_lora_stage_1` and `distilled_lora_stage_2` from one path (`:92-101`) and hands one to each `DiffusionStage.from_checkpoint` (`:154`, `:165`), CLI-defaulted 0.25 and 0.5 (`utils/args.py:1174-1184`). Here neither phase sets `Ltx2PhaseRecipe::loras`, so both take the `kAllAdapters` default, and the engine carries ONE strength for the whole load (`lora_strength` absent is 1.0, `include/vllm/multimodal/ltx2_video.h:214-218`) — so both stages run at 1.0. Nothing refuses and nothing changes shape; only a render against upstream on the same checkpoint, take and seed sees it. A NEW FIELD ON `Ltx2PhaseRecipe` DOES NOT CLOSE IT: `Ltx2RebindDitLoras` early-returns on `currently_fused == fuse` where `currently_fused` is `checkpoint.lora_fused_tensors > 0`, a BOOLEAN, so it detects "already fused" and never "already fused AT THIS STRENGTH" — and HQ has both stages fused, so the no-op its own header advertises would swallow the transition and stage 2 would render at stage 1's strength. Closing it needs `bool fuse` to become a type carrying a strength AND `Ltx2DitCheckpoint` to record WHICH adapter state is applied, plus a `--distilled-lora-strength-stage-1`/`-stage-2` request pair, since one `lora_strength` extra cannot spell two values. The re-materialize-and-write-back mechanism #1118 landed is untouched by that change. Filed because [#921](https://github.com/mudler/vllm.cpp/issues/921), which owned the per-phase strength in [`ltx25-phase-lora.md`](specs/ltx25-phase-lora.md), was CLOSED as completed on 2026-08-17 by `LTX25-RES2S-LOOP` (`4d7748646`, PR [#1125](https://github.com/mudler/vllm.cpp/pull/1125)); that row named the distilled LoRA per stage as out of scope and correct to leave, but did not list it under its own `## Owed`, so the debt outlived its issue with no open owner. Found during the review repair of PR [#1140](https://github.com/mudler/vllm.cpp/pull/1140). Listed under `## Owed` in [`ltx25-phase-lora.md`](specs/ltx25-phase-lora.md) | bug | diff --git a/.agents/specs/expert-streaming.md b/.agents/specs/expert-streaming.md index 807cc7544..f3d74317b 100644 --- a/.agents/specs/expert-streaming.md +++ b/.agents/specs/expert-streaming.md @@ -1360,6 +1360,232 @@ boundary). The Windows repair is NOT mutation-proven: no MSVC is reachable from this host, and the checker that would have caught it statically is the subject of #1107. +## `--device cuda` loads for 26 minutes and then dies: the allocation, named (#1123) + +17 August 2026. The same checkpoint that serves on `--device cpu` reaches a +serving state on `--device cuda` on the same box and the same binary, and then +dies on the first request with `vt cuda: cudaMalloc: out of memory` inside the +EngineCore busy loop. The load succeeded, so the failure is the forward. + +The log line cannot say which allocation failed or how big it was: +`CudaBackend::Alloc` is `Check(cudaMalloc(&p, bytes), "cudaMalloc")` +(`src/vt/cuda/cuda_backend.cu:77-81`) and `Check` composes +`"vt cuda: " + what + ": " + cudaGetErrorString(err)` +(`cuda_backend.cu:48-52`), where `what` is the literal `"cudaMalloc"`. `bytes` +is in scope and discarded. So the size was established from the code and the +checkpoint rather than from the message. + +### The allocation + +`ResidentWeight` (`src/vllm/model_executor/models/qwen3_5.cpp:963-1025`) aliases +the host bytes when `GetPlatform(...).is_cpu()` and otherwise **uploads the whole +tensor**: `const size_t nb = w.bytes.size(); void* p = d.b.Alloc(nb);` +(`qwen3_5.cpp:1010-1011`). For a routed-expert weight, `w` is the STACKED +`[E*N,K]` keep-quant tower, so `nb` is one tower — every expert of one matrix of +one layer, in one contiguous `cudaMalloc`. + +**Both switch positions of the keep-quant MoE path reach that same line**, which +is why no knob avoids it. Every `f:N` below is a CALL SITE — the line inside `f` +that invokes the next hop — never a definition line, so the chain can be walked +one `sed` at a time: + +| Configuration | Path | Reaches | +|---|---|---| +| default (`VT_QWEN35_GROUPED_MOE` unset ⇒ on) | `MoeBlock:6615,6616,6620` → `KqGrouped:5694` | `ResidentWeight(d, w_kq)` | +| `VT_MOE_EXPERT_STREAM=1` (which DISABLES grouping, `:5670-5676`) | `ExpertMlpKq:5651,5652` → `MatmulF32Slice:5611` → `KqExpertSlice:5595` | `KqResidentSlice:5114` → `ResidentWeight` | + +`KqExpertSlice`'s slot arm is guarded by `is_cpu()` (`qwen3_5.cpp:5578`), so on a +device platform it falls through before the store is even constructed. That is +the `## Owed` line "Streaming serves the CPU-resident borrowed tower only" +observed from the other end: the lane that makes the model fit has no device arm, +and the device path therefore asks `cudaMalloc` for every expert byte. + +Ruled out by reading, not by assumption. `src/vt/cuda/cuda_moe.cu` and +`cuda_glue.cu` contain no allocation at all. `BuildMoeMarlinResident` +(`qwen3_5.cpp:6010-6215`; its `E ×` per-expert allocations are `:6049-6064`, plus +two repack temporaries at `:6094-6095`) does allocate unpooled, and it +is NOT on this path: `MoeBlock` takes the fp4/Marlin arm at `:6555`, guarded by +`const bool fp4 = !w.expert_gate_fp4.empty()` at `:6548`, and a GGUF keep-quant +load populates `expert_*_kq`, not `expert_*_fp4`. + +### The size, measured from the checkpoint + +Both GGUF tensor tables of `unsloth/Qwen3.8-2.4T-A95B-GGUF` at revision +`567d3e6ac26c5474b18311e619c04350fb9a5556` were re-censused independently of the +earlier census in this spec, by HTTP range request over all ten shards, with no +tensor data downloaded: **1702 tensor records parsed against the 1702 declared in +`split.tensors.count`**, and the two numbers agreeing is the coverage claim. +Shard 1 carries the 58 metadata keys and ZERO tensors; shards 2-10 carry the +table. + +| Encoding | ggml | Tensors | Bytes | % bytes | +|---|---|---|---|---| +| IQ1_XXXS (routed experts) | 66 | 276 | 351,918,882,816 | 88.59 | +| Q5_K | 13 | 420 | 23,391,633,408 | 5.89 | +| Q6_K | 14 | 162 | 8,876,851,200 | 2.23 | +| Q2_K (the `nextn` MTP block's 3 towers) | 10 | 3 | 8,455,716,864 | 2.13 | +| Q4_K | 12 | 2 | 2,288,517,120 | 0.58 | +| F32 | 0 | 838 | 2,171,133,440 | 0.55 | +| Q8_0 | 8 | 1 | 142,606,336 | 0.04 | +| total | | 1702 | 397,245,341,184 | 369.96 GiB | + +`expert_count = 512`, `embedding_length = 8192`, +`expert_feed_forward_length = 2048`, `expert_used_count = 10`, +`block_count = 93`. So one tower is: + +| Tower | Bytes | | Count | +|---|---|---|---| +| IQ1_XXXS `ffn_{gate,up,down}_exps`, dims `[8192,2048,512]` / `[2048,8192,512]` | **1,275,068,416** | 1.1875 GiB | 276 | +| Q2_K `ffn_{gate,up,down}_exps` (block 92, the MTP layer) | **2,818,572,288** | 2.6250 GiB | 3 | +| all `*_exps` | **360,374,599,680** | **335.62 GiB** | 279 | + +`1,275,068,416 / 512 = 2,490,368` bytes per expert slice, which is exactly the +`slot_bytes=2490368` the W4 banner printed, so the arithmetic here and the +running lane agree on the same weight. + +**So the answer to "which allocation and how big" is: `qwen3_5.cpp:1011`, +1,275,068,416 bytes at a time (2,818,572,288 for three of them), 279 times, +335.62 GiB in total.** + +### Why it fails, and why the loader does not + +The budget is measurable, and `nvidia-smi` is the wrong instrument for it: on +this box it answers `[N/A], [N/A], [N/A]` for `memory.total,memory.free, +memory.used`, and the `rc` fleet label likewise records `vram=[N/A]M`. +`cudaMemGetInfo` answers honestly. Measured on `dgx:gpu0` under an `rc` hold, +through `libcudart.so.13` in the `vllmcpp-build:gb10` image: + +``` +cudaGetDeviceCount rc = 0 count = 1 +cudaMemGetInfo rc = 0 +free = 122059919360 (113.677 GiB) +total = 128452956160 (119.631 GiB) +attr Integrated rc=0 value=1 +attr UnifiedAddressing rc=0 value=1 +``` + +`total` is EXACTLY `/proc/meminfo MemTotal` (125442340 kB) times 1024, and equals +the fleet's own `mem_total_bytes=128452956160`. **One unified pool, correctly +reported by the CUDA runtime and not by `nvidia-smi`.** + +Against that pool, 335.62 GiB of tower staging is 2.8x the whole machine, and +the total device-resident weight demand (towers plus the dense remainder, which +this spec measured at 62 GiB resident on the CPU arm) is over 3x. The load +survives because a borrowed tower costs ZERO anonymous bytes — that is the +finding in "Why a 370 GiB model fits in 119 GiB" above. Staging converts each +borrow into a real allocation, so the pool is exhausted after roughly +`(119.6 - 62) / 1.1875 ≈ 48` towers, i.e. partway through layer 16 of 93, on the +first forward. Hence: 26 minutes to READY, then death mid-stream. + +### The refusal: what it keys on, and what it deliberately does not + +Loading for 26 minutes and dying mid-stream is the worst of the three available +behaviours, and AGENTS.md already says which one is right: refuse an +unimplemented arm at load with a message that names the missing part. This +change lands that refusal and nothing else. It does NOT build the device-slot +arm; that stays owed below. + +The predicate is keyed on the MEASURED condition, never on "CUDA + GGUF" and +never on an architecture name, because a GGUF that genuinely fits the pool must +still load: + +``` +refuse ⇔ needs_weight_staging() ∧ budget_known ∧ staged_lower_bound > budget +``` + +Three properties are deliberate. + +**A PER-TENSOR lower bound, and it is wrong in both directions.** Per tensor the +bound is `min(gguf_bytes, elems × model_dtype_bytes)`. A weight the loader keeps +quantized is staged verbatim, which is `gguf_bytes`; a weight it expands is staged +at the model dtype, which is `elems × 2` for bf16. Taking the minimum makes each +term a true lower bound on THAT tensor's staged size. + +It does not follow that the sum is a lower bound on the load, and the first +version of this section claimed it did — "so the refusal can never over-refuse". +The review that caught it supplied the counter-example, which is present on every +default load: + +- **Over-count.** A tensor counted and never staged is a positive error. The MTP / + `nextn` block is attached only when a speculator is configured, so on a default + load block 92 of the target checkpoint — 20 tensors, 8,940,488,704 bytes, 8.33 + GiB, **2.2506 %** — is counted and not staged. A budget in + `[what a default load stages, what this counts)` refuses a weight set that fits. +- **Under-count.** The bound omits KV cache, activations, scratch pools and the + CUDA context, so a checkpoint at 0.95x of the pool still passes here and still + dies later. + +The two errors are on DIFFERENT quantities and do not cancel, so "the under-count +dominates" is not an argument that the refusal is safe — it is an argument about a +number the refusal never compares. `gguf_device_fit.h` said this correctly from the +day it landed; the spec and the commit body did not, which is why the wording here +is now the header's. `test_gguf_device_fit` carries a counted-but-unstaged fixture +so the over-count direction is executable rather than described, and both remainders +are owed: the over-count to [#1136](https://github.com/mudler/vllm.cpp/issues/1136), +because closing it means teaching the bound which tensors THIS load will stage, +which is load policy and not a property of the file; the under-count to the startup +memory profile `KV-WARMUP-PROFILE` owns. A headroom fraction invented here would be +the guess this bound exists to avoid, in either direction. + +**`total`, not `free`.** `free` at load time carries the page cache and whatever +else the box is doing, so it makes the refusal a function of contention. `total` +is a device property. + +**Unknown is not a verdict.** `ResidencyPolicy::device_memory_total_bytes` is 0 +on every platform that does not probe one, and 0 means UNKNOWN. A caller that +cannot learn the budget declines to decide, so no non-CUDA device and no CUDA +build without the probe changes behaviour. This is the polarity +`gemma4_moe.cpp:506` chose for the opposite reason (it refuses the device +allocation on unknown, because a hung `hipMalloc` is worse than a host +fallback); here the risk runs the other way, since refusing a load on an unknown +budget would break every device whose budget nothing reports. + +**Which platforms this actually covers**, because the first version of this +section and `docs/USAGE.md` both got it wrong in the same way. The two predicates +coincide on exactly one platform: `needs_weight_staging()` is true only on +`CudaPlatform` (`src/vllm/platforms/cuda.cpp:71`) and a budget is probed only +there. So **every** NVIDIA GPU this build runs on gets both the probe and the +refusal — a discrete card is `CudaPlatform` too, not a separate case. ROCm, +Vulkan and Metal answer `needs_weight_staging() == false` (ROCm says so +explicitly, `src/vllm/platforms/rocm.cpp:74`), which means they read the mapping +where it lies and have no staging allocation to fail: the refusal is not "owed" +to them, it is inapplicable. What IS owed on ROCm is the separate +`Backend::DeviceMemoryInfo` capability (#1126). + +The probe is added to `CudaPlatform`, which already includes `` +and already probes device attributes at registration, and NOT to +`Backend::DeviceMemoryInfo`. That seam's comment claimed "ROCm/CUDA override with +hipMemGetInfo/cudaMemGetInfo" and only ROCm does +(`src/vt/rocm/rocm_backend.hip:338-345`). The comment is **corrected in this +change**, in the two places that carried it: `include/vt/backend.h:78-93` on the +seam, and `gemma4_moe.cpp:440-448` on the only call site — the second copy was +found by this round's audit and is why the first correction alone would have left +the claim in the tree. Overriding the seam would also silently wake `Gemma4MoE`'s +device-expert LRU, whose `MakeRoom` refuses on CUDA today precisely because the +query is absent (`gemma4_moe.cpp:506`). Waking another model's residency policy is +a behaviour change with its own measurement, so it is filed rather than done. + +### Tests, and how the device branch is reached on a CPU-only host + +`needs_weight_staging()` is true on exactly one platform in this tree +(`src/vllm/platforms/cuda.cpp:71`), so the branch is unreachable from the real +loader on a host with no CUDA device — the untestable-device-branch shape this +row has hit repeatedly. It is reached here by registering a FAKE staging +platform in the CUDA lookup slot, which is the instrument +`tests/vllm/entrypoints/test_device_selection.cpp` already established for +exactly this reason, in its own executable so the global registry cannot leak +into other suites. + +| Case | Instrument | +|---|---| +| the arithmetic, both directions and the boundary | `GgufStagedWeightFootprint` / `CheckDeviceWeightFit` over a table: `>` refuses, `==` and `<` do not, unknown budget does not, a non-staging platform does not, an F32 tensor is counted at bf16 and a quantized one at its GGUF size | +| the bound's OVER-count direction | a fixture carrying a `blk.N.nextn.*` tensor a default load never stages: the footprint counts it, and both ends of the resulting over-refusal window are asserted, so the direction cannot be claimed away again | +| the AUTO arm names the device the load will RUN on | a fake staging platform whose backend's `CreateQueue()` can be made to throw, driven through `FromModelDir` twice at the same budget: it refuses when the queue can be created, and refuses NOTHING when it cannot, because that load runs on CPU | +| the CUDA residency policy assembles the probed budget | `CudaResidencyPolicy` in `platforms/interface.h`, unit-tested on every host, so the assignment is no longer reachable only in a CUDA build | +| the refusal is REACHED from the loader | `LoadedEngine::FromModelDir` on a synthetic `qwen35moe` GGUF with the fake staging platform registered and a small `VT_DEVICE_WEIGHT_BUDGET_BYTES`: the thrown message is the fit refusal | +| a fitting GGUF still loads | the SAME call with a generous budget: the throw is a LATER, different one (the synthetic file has no tokenizer), which is what proves the check let it through rather than that it never ran | +| the CPU arm is untouched | the same file with `device=cpu` never refuses, whatever the budget | + ## Owed Carried debt for this row. Each item names why it is not closed here. @@ -1377,6 +1603,12 @@ Carried debt for this row. Each item names why it is not closed here. | **The CUDA arms of [#1029](https://github.com/mudler/vllm.cpp/issues/1029)'s grouped gate have not run on a device.** | Recorded in that issue, which stays open for it. Unchanged by this repair. | | **Streaming serves the CPU-resident borrowed tower only.** A staged device weight takes `KqResidentSlice`, so there is no device-slot arm. | Deliberate for phase 1 (copying a device-resident weight through host slots moves MORE bytes); W7 owns the pluggable backing store. | | **The grouped keep-quant MoE path and streaming are mutually exclusive.** `VT_MOE_EXPERT_STREAM=1` disables grouping and says so once on stderr. | Grouping stages the whole tower, which is what streaming exists to avoid. Making them compose needs a slot-aware grouped GEMM, which is its own row. | +| **`--device cuda` still cannot SERVE a larger-than-pool GGUF; it only refuses by name now.** The device-slot arm is the missing capability: a `DeviceExpertSlotStore` behind `ExpertSlotStore`, a read accessor on that interface (`KqExpertSlice` reads `HostExpertSlotStore::Slot()`, the CONCRETE class, so the seam cannot be swapped today), a device filler that is not `pread`-into-host (`ExpertSlotStore::SlotForWrite` is handed straight to `::pread`, `expert_streamer.cpp:76-94`), and lifting the `is_cpu()` guard at `qwen3_5.cpp:5578`. Sized by the measurement above: 2790 slices per token at 2,490,368 bytes is 6.95 GB per token against a 119.631 GiB pool that already holds the dense remainder. Tracked as [#1124](https://github.com/mudler/vllm.cpp/issues/1124). | It is a campaign, not a fix: W7 (the pluggable backing store) is its declared owner in the work breakdown, and the CPU arm's own I/O rate is still unmeasured on a live cache two rows above. Building a device lane on top of a host lane whose bandwidth number is void would be optimising against a number nobody has. | +| **The fit bound omits everything that is not a weight.** KV cache, activations, the scratch pools and the CUDA context are not counted, so a checkpoint at 0.95x of the pool passes the refusal and still dies on the first forward. | A headroom fraction invented here would be exactly the guess the per-tensor bound exists to avoid. The number wants the startup memory profile that `KV-WARMUP-PROFILE` owns (`INVENTORIED`; upstream's is `GPUWorker.determine_available_memory`, `vllm/v1/worker/gpu_worker.py:451-495`, around `profile_run`, `vllm/v1/worker/gpu/model_runner.py:682`), which is a different row. Those two anchors are stated here from the pinned tree because that row's own three anchors are stale at the current pin, and `gguf_device_fit.h` had copied two of them — filed as [#1139](https://github.com/mudler/vllm.cpp/issues/1139), owned by `KV-WARMUP-PROFILE`, blocked here only by the `engine-matrix.md` record lock #1119 holds. | +| **The fit bound also counts too MUCH, and that direction can refuse a load that fits.** A tensor present in the file and not staged by THIS load is a positive over-count. On a default load that is the MTP / `nextn` block: 8,940,488,704 bytes, 8.33 GiB, 2.2506 % of the target checkpoint. A budget in that window refuses a weight set that would have fitted. | Not closed here. Closing it means the bound taking a per-tensor staging POLICY as input, which is the caller's knowledge and not the file's, and the exclusion's own failure mode is an under-count to nothing — which restores the 26-minute-then-OOM this row exists to remove, on a device nobody here has to measure it on. So the direction is stated in `gguf_device_fit.h`, pinned executably by `test_gguf_device_fit`, exposed to operators in `docs/USAGE.md`, and tracked as [#1136](https://github.com/mudler/vllm.cpp/issues/1136). `VT_DEVICE_WEIGHT_BUDGET_BYTES` is the way out of the window in the meantime. | +| **`Backend::DeviceMemoryInfo` has no CUDA override.** Only ROCm implements it (`src/vt/rocm/rocm_backend.hip:338-345`), so `Gemma4MoE`'s device-expert LRU refuses on every CUDA device and falls back to host H2D silently. | The CAPABILITY is [#1126](https://github.com/mudler/vllm.cpp/issues/1126) and is not built here, because adding the override wakes a landed residency policy that needs its own measurement. What IS done here is the false COMMENT, corrected in both places that carried it: `include/vt/backend.h:78-93` and `gemma4_moe.cpp:440-448`. The second copy was found by this round's own audit; correcting only the seam would have left the claim in the tree, which is the shape this row keeps hitting. | +| **`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). | ## Risks/decisions diff --git a/CMakeLists.txt b/CMakeLists.txt index 9798db800..0cefa2ddc 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -717,6 +717,7 @@ add_library(vllm STATIC src/vllm/model_executor/model_loader/gguf_reader.cpp src/vllm/model_executor/model_loader/gguf_dequant.cpp src/vllm/model_executor/model_loader/gguf_keep_quant.cpp + src/vllm/model_executor/model_loader/gguf_device_fit.cpp src/vllm/model_executor/model_loader/nvfp4_dequant.cpp src/vllm/model_executor/model_loader/awq_gptq_dequant.cpp src/vllm/model_executor/model_loader/mxfp4_dequant.cpp diff --git a/docs/ENVIRONMENT.md b/docs/ENVIRONMENT.md index a9a4ffa43..34b8238c1 100644 --- a/docs/ENVIRONMENT.md +++ b/docs/ENVIRONMENT.md @@ -81,6 +81,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 keep `64`. 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` | | `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 keep the default. 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 | | `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 | +| `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 | ## Rollback and bisect switches diff --git a/docs/USAGE.md b/docs/USAGE.md index 0dad6fff0..0c7c4035c 100644 --- a/docs/USAGE.md +++ b/docs/USAGE.md @@ -3629,6 +3629,69 @@ A run whose `steps` is 0, or whose `exhausted` is large, is not a measurement of streaming, whatever the startup line said. See [`docs/ENVIRONMENT.md`](ENVIRONMENT.md) for every knob and its parsing rules. +### `--device cuda` refuses a checkpoint it cannot hold + +Streaming is a **host** capability. The GGUF mapping is borrowed in place on the +CPU path, so a routed-expert tower costs no resident bytes, which is the whole +reason a 369.96 GiB checkpoint serves on a 119.631 GiB box. A weight-staging +device has no such lane: it copies every tower into device memory, one +`cudaMalloc` per stacked `[E*N,K]` tower. + +For `Qwen3.8-2.4T-A95B UD-Q1_0` that is 276 towers of 1,275,068,416 bytes plus +three of 2,818,572,288, so 335.62 GiB in total, against a pool `cudaMemGetInfo` +reports as +128,452,956,160 bytes (119.631 GiB). Until that lane exists +([#1124](https://github.com/mudler/vllm.cpp/issues/1124)), the engine **refuses +at load** and names what is missing: + +```text +device 'cuda' cannot serve this GGUF: staging its weights needs at least N bytes +(X GiB) of device memory across T tensors, the largest single allocation being M +bytes (Y GiB, ''), and this device's memory pool is B bytes (Z GiB). +THE MISSING PART: ... there is no device-side expert slot store and no device +streaming lane ... Use device=cpu, which serves this checkpoint today, or a +checkpoint that fits the pool. +``` + +It used to load for 26 minutes, report ready, and then die on the first request +with `vt cuda: cudaMalloc: out of memory` from inside the engine's busy loop +([#1123](https://github.com/mudler/vllm.cpp/issues/1123)). + +The refusal is keyed on the measured condition and not on the device or the file +format, so **a GGUF that fits the pool still loads on `--device cuda`**. Three +things it deliberately does not do: + +- it never fires on a platform that does not stage weights, so every + `--device cpu` load is unchanged; +- it never fires when no budget is known. Today exactly one platform stages + weights (CUDA) and exactly one probes a budget (CUDA, with `cudaMemGetInfo`), + so **every NVIDIA GPU this build runs on — discrete or GB10 — gets both the + probe and the refusal**, while ROCm, Vulkan and Metal answer + `needs_weight_staging() == false`: they read the GGUF mapping where it already + lies, so there is no staging allocation to fail and nothing for this check to + decide. What is owed there is the `Backend::DeviceMemoryInfo` probe CUDA does + not implement ([#1126](https://github.com/mudler/vllm.cpp/issues/1126)), which + is a different capability; +- it counts **weights only**. The KV cache, activations, scratch pools and the + driver context are not in the bound, so a checkpoint just under the pool + passes this check 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 + `VT_DEVICE_WEIGHT_BUDGET_BYTES` 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. + +**The instrument matters here.** `nvidia-smi +--query-gpu=memory.total,memory.free,memory.used` answers `[N/A], [N/A], [N/A]` +on a GB10, because host and device share one pool. `cudaMemGetInfo` answers +honestly, and its `total` is EXACTLY `/proc/meminfo MemTotal` +(125442340 kB) times 1024. Do not size this from `nvidia-smi`. + ## SSE keepalives on long prefill Async chat/completion streams can emit SSE **comment** frames (`:\n\n`) while diff --git a/include/vllm/entrypoints/model_loader.h b/include/vllm/entrypoints/model_loader.h index cc5eadaaa..3a979bea9 100644 --- a/include/vllm/entrypoints/model_loader.h +++ b/include/vllm/entrypoints/model_loader.h @@ -199,6 +199,31 @@ struct EngineParams { vt::Queue SelectQueueForModel(std::string_view architecture, vllm::Device device); +// The device type `SelectQueueForModel` will pick. +// +// It exists because the load-time GGUF device-fit refusal (issue #1123) has to +// know the target device BEFORE any weight I/O, and the load's own queue is not +// created until after the weights are loaded. Throws for an explicitly named +// device that this build/process cannot serve, exactly as the queue selector +// does; the auto arm falls back to `kCPU` instead of throwing, also exactly as it +// does. +// +// The two agree because both arms run one implementation, and on the AUTO arm +// that implementation CREATES A QUEUE and destroys it. #1136 measured why the +// cheaper version was wrong: `SelectQueueForModel`'s auto arm falls back to CPU +// when `CreateQueue()` throws, so a resolver that only asked `CurrentPlatform()` +// answered `'cuda'` on a box where the load would run on CPU, and the fit refusal +// then rejected a checkpoint by naming a device nothing was going to run on. +// +// The cost of agreeing is one extra stream created and destroyed, and it is bounded +// by where this function is called: the load-time GGUF fit check is the only caller +// outside `SelectQueueForModel` itself, so a safetensors load pays nothing, an +// explicitly named device pays nothing (that arm creates no queue here), and an +// auto-arm GGUF load pays one `CreateQueue`/`DestroyQueue` pair. That is not free, +// and it is smaller than removing a working load. +vt::DeviceType ResolveModelDeviceType(std::string_view architecture, + vllm::Device device); + // Owns the full V1 engine stack (config + weights + tokenizer + Scheduler + // runner -> Executor -> EngineCore; Input/OutputProcessor -> LLMEngine) for a // registered model. The concrete weights/forward are held behind LoadedModel; diff --git a/include/vllm/model_executor/model_loader/gguf_device_fit.h b/include/vllm/model_executor/model_loader/gguf_device_fit.h new file mode 100644 index 000000000..c0bff5b16 --- /dev/null +++ b/include/vllm/model_executor/model_loader/gguf_device_fit.h @@ -0,0 +1,141 @@ +// vllm.cpp original. Pinned vLLM (555967922) has no GGUF load format — the whole +// tree carries two incidental mentions of the word and no loader — so there is no +// upstream counterpart to mirror. +// +// The closest upstream idea is the startup memory profile, and it answers a +// different question. `GPUWorker.determine_available_memory` +// (`vllm/v1/worker/gpu_worker.py:451-495`) runs `memory_profiling` around +// `profile_run` (`vllm/v1/worker/gpu/model_runner.py:682`) and passes the weight +// bytes IN as a known quantity, `weights_memory=model_memory_usage` +// (`gpu_worker.py:493`). That quantity is recorded AFTER the load has finished +// (`gpu/model_runner.py:315`). So upstream measures what is left once the weights +// are resident; it never asks whether they will fit, because by then it has paid +// for them. This file asks the question upstream does not. +// +// ENG-EXPERT-STREAM, issue #1123: a load-time answer to "can this device +// actually hold this checkpoint's weights?". +// +// The failure this exists to remove: `Qwen3.8-2.4T-A95B UD-Q1_0` (369.96 GiB) +// reached a serving state on `--device cuda` on a 119.631 GiB GB10 after 26 +// minutes and then died on the FIRST forward with `vt cuda: cudaMalloc: out of +// memory`. The load succeeds because a keep-quant expert tower is BORROWED from +// the GGUF mapping and costs zero anonymous bytes; the forward dies because a +// weight-staging platform copies each borrowed tower into device memory +// (`ResidentWeight`, `qwen3_5.cpp:1011`, `d.b.Alloc(w.bytes.size())` on a +// STACKED `[E*N,K]` tower: 276 towers of 1,275,068,416 bytes plus 3 of +// 2,818,572,288, so 360,374,599,680 bytes = 335.62 GiB in total). +// +// See `.agents/specs/expert-streaming.md`, section "`--device cuda` loads for 26 +// minutes and then dies", for the measurement this is keyed on. +#pragma once + +#include +#include +#include + +#include "vllm/model_executor/model_loader/gguf_reader.h" + +namespace vllm { + +// A lower bound on the device-resident bytes a weight-staging platform must +// allocate to stage EVERY TENSOR IN THIS FILE. +// +// Read that scope literally, because it is where the one over-count comes from: it +// is NOT a lower bound on what a particular load stages, since a load may stage a +// subset of the file. The per-tensor term is exact-or-low; the SET is exact-or-high. +// +// Per tensor the bound is `min(gguf_bytes, elems * model_dtype_bytes)`: a weight +// the loader keeps quantized is staged verbatim (`gguf_bytes`), a weight it +// expands is staged at the model dtype (`elems * 2` for bf16), and which of the +// two happens is a per-tensor loader policy this header deliberately does not +// try to predict. Taking the minimum makes the per-tensor term a true lower +// bound on that tensor's staged size. +// +// Both directions of error are named here rather than claimed away, because a +// bound whose error direction is unstated is not a bound. +// +// * It can OVER-count, by including a tensor the loader never stages. The one +// such class in this tree is the MTP / `nextn` block, which is attached only +// when a speculator is configured +// (`model_loader.cpp`, the `speculative_config->method == "mtp"` guard). +// Measured on the target checkpoint, that is block 92: 20 tensors, +// 8,940,488,704 of 397,245,341,184 bytes, so 2.2506 %. +// * It UNDER-counts by everything that is NOT a weight: KV cache, +// activations, the device scratch pool and the driver context. That term is +// far larger than 2.2506 %. +// +// The two errors are on DIFFERENT quantities and do not cancel, so neither +// rescues the other. In particular the refusal CAN over-refuse: a budget in +// [what a default load stages, what this counts) rejects a weight set that fits. +// On the target checkpoint that window is 8.33 GiB wide on a 369.96 GiB +// checkpoint, `VT_DEVICE_WEIGHT_BUDGET_BYTES` is the operator's way out of it, +// and `test_gguf_device_fit` pins the direction rather than leaving it described +// (issue #1136). Closing it means teaching the bound which tensors THIS load will +// stage, which is load policy and not a property of the file, so it is owed and +// not invented here. The under-count is owed to the startup memory profile +// (`KV-WARMUP-PROFILE`) for the same reason: an invented headroom fraction here +// would be the guess this bound exists to avoid. +struct GgufStagedFootprint { + // The bound, in bytes. The name is accurate for what it measures — the sum of + // per-tensor lower bounds over the file's whole tensor table — and it is NOT a + // lower bound on one load's staging, for the reason above. + size_t lower_bound_bytes = 0; + // How many tensor records went into it. A caller that reports a footprint + // without this cannot say how many things it examined. + size_t tensor_count = 0; + // The largest single tensor, which is the largest single contiguous + // allocation the staging path will ask the driver for. An aggregate that fits + // is not the same as a contiguous block that fits. + size_t largest_tensor_bytes = 0; + std::string largest_tensor_name; +}; + +// `model_dtype_bytes` is the resolved model dtype's size (2 for bf16, which is +// what every GGUF path here loads at). vLLM resolves ONE model dtype and every +// layer inherits it, so one value is the faithful shape. +GgufStagedFootprint GgufStagedWeightFootprint(const GgufFile& gguf, + size_t model_dtype_bytes = 2); + +// The budget to compare a footprint against, in bytes, or 0 for UNKNOWN. +// +// `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. +// +// 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 +// of contention. `total` is a device property. +size_t DeviceWeightBudgetBytes(size_t device_memory_total_bytes); + +// The verdict. `refuse == false` with a zero budget means "not decided", +// which is NOT the same as "it fits" — see the comment on the budget above. +struct DeviceWeightFit { + bool refuse = false; + size_t needed_bytes = 0; + size_t budget_bytes = 0; + // Empty unless `refuse`. Names the device, both numbers, the missing + // capability and what to do instead. + std::string message; +}; + +// The predicate, in one place so one description exists: +// +// refuse <=> needs_weight_staging AND budget != 0 AND needed > budget +// +// Keyed on the MEASURED condition, never on "CUDA + GGUF" and never on an +// architecture name, so a GGUF that genuinely fits the pool still loads. +// Strictly greater than: a checkpoint whose footprint exactly equals the budget is +// not refused here. The footprint is approximate in both directions, so equality is +// not evidence of anything, and the tie goes to attempting the load. +DeviceWeightFit CheckDeviceWeightFit(const GgufFile& gguf, + std::string_view device_name, + bool needs_weight_staging, + size_t budget_bytes, + size_t model_dtype_bytes = 2); + +} // namespace vllm diff --git a/include/vllm/platforms/interface.h b/include/vllm/platforms/interface.h index ba21a058a..ebfb15439 100644 --- a/include/vllm/platforms/interface.h +++ b/include/vllm/platforms/interface.h @@ -53,6 +53,23 @@ struct ResidencyPolicy { // Optional soft cap (bytes) on pooled device scratch; 0 == uncapped (today). // Consumed by the DevicePool: a discrete GPU sets a bound; GB10 leaves it 0. size_t device_pool_cap_bytes = 0; + // Total device memory this platform can draw weight allocations from (bytes). + // 0 == UNKNOWN, and unknown must never be read as "unlimited" or as "nothing + // fits": a caller that cannot learn the budget declines to decide. Consumed by + // the load-time GGUF fit refusal (`gguf_device_fit.h`, issue #1123). + // + // TOTAL rather than FREE, because `free` at load time carries the page cache + // and whatever else the box is doing, which would make a load-time verdict a + // function of contention. + // + // This is NOT `vt::Backend::DeviceMemoryInfo`, which is a live free/total + // probe whose only consumer is `Gemma4MoE`'s device-expert LRU and which + // `CudaBackend` does not override at all -- so wiring CUDA into that seam + // would wake a landed residency policy that is currently dead on CUDA, which + // is a behaviour change with its own measurement to make (issue #1126). + // Probed once at platform registration; on a GB10 `cudaMemGetInfo` reports + // 128452956160 (119.631 GiB) where `nvidia-smi` reports `[N/A]`. + size_t device_memory_total_bytes = 0; }; // Residency DECISIONS derived from a ResidencyPolicy — the single, testable place @@ -75,6 +92,31 @@ inline bool ShouldInterleaveLoadStream(const ResidencyPolicy& policy, return marlin_committed && policy.release_host_weights_after_upload; } +// The CUDA platform's residency policy, assembled from the values its registrar +// probes. `src/vllm/platforms/cuda.cpp` calls this and supplies +// `cudaMemGetInfo`'s `total`; it holds no policy of its own. +// +// A free function HERE, rather than a body inside that file, because that +// translation unit compiles ONLY in a CUDA build. While the assembly lived there, +// nothing on a host without a CUDA toolkit could reach it — no test, and no +// mutation, which is why "delete the `device_memory_total_bytes` assignment" was +// recorded as an owed mutation by #1123 instead of being proven. This header +// compiles everywhere, so tests/vllm/platforms/test_platform.cpp pins the +// assembly on every host and `cuda.cpp` is left holding only the probe it alone +// can make (#1136). The probe itself is still not mutation-proven; a CUDA build +// is the only thing that reaches the `cudaMemGetInfo` call. +// +// `device_memory_total_bytes` is 0 when the probe failed, which the load-time +// GGUF fit refusal reads as UNKNOWN and never as "nothing fits". +inline ResidencyPolicy CudaResidencyPolicy(size_t device_memory_total_bytes) { + ResidencyPolicy p; + p.release_host_weights_after_upload = true; // freed after the Marlin build + p.uses_device_memory_pool = true; // qwen3_5.cpp DevicePool + p.device_pool_cap_bytes = 0; // uncapped + p.device_memory_total_bytes = device_memory_total_bytes; + return p; +} + // The selection inputs of vllm/platforms/cuda.py::_get_backend_priorities @ pin // e24d1b24 (`use_mla`, `device_capability`, `num_heads`, `kv_cache_dtype`) plus // the sparse flag that `AttentionBackend.is_sparse()` / diff --git a/include/vt/backend.h b/include/vt/backend.h index 61d47bc6c..462a5d5bf 100644 --- a/include/vt/backend.h +++ b/include/vt/backend.h @@ -76,8 +76,21 @@ class Backend { virtual bool DeviceMemoryIsHostAddressable() const { return false; } // Optional device free/total VRAM probe (bytes). Default false = unknown. - // ROCm/CUDA override with hipMemGetInfo/cudaMemGetInfo so model code can - // size LRU caches without including vendor headers (device-leakage). + // ROCm overrides it with hipMemGetInfo (src/vt/rocm/rocm_backend.hip) so model + // code can size LRU caches without including vendor headers (device-leakage). + // + // CUDA does NOT override it. This comment claimed "ROCm/CUDA" until #1123 + // measured what that costs: `Gemma4MoE`'s device-expert LRU is the seam's only + // consumer, its `FreeBytes` returns false on an absent probe and `MakeRoom` + // then refuses the device upload (both in gemma4_moe.cpp), so on EVERY CUDA + // device that cache admits nothing and falls back to host H2D, silently. + // Adding the override therefore WAKES a landed residency policy and needs its own + // measurement; that is issue #1126, and this line says what is true until then. + // + // The load-time GGUF fit refusal deliberately does not read this seam: it is a + // live free/total probe, and a load-time budget must not be a function of + // contention. It carries its own `total` on + // `vllm::platforms::ResidencyPolicy::device_memory_total_bytes` instead. virtual bool DeviceMemoryInfo(size_t* /*free_bytes*/, size_t* /*total_bytes*/) const { return false; } diff --git a/src/vllm/entrypoints/model_loader.cpp b/src/vllm/entrypoints/model_loader.cpp index 30b7577b8..28348080b 100644 --- a/src/vllm/entrypoints/model_loader.cpp +++ b/src/vllm/entrypoints/model_loader.cpp @@ -14,6 +14,7 @@ #include #include #include +#include #include #include #include @@ -22,6 +23,7 @@ #include #include "vllm/model_executor/weight_offloader.h" +#include "vllm/model_executor/model_loader/gguf_device_fit.h" #include "vllm/model_executor/model_loader/gguf_reader.h" #include "vllm/model_executor/model_loader/safetensors_reader.h" #include "vllm/model_executor/models/deepseek_v4.h" // deepseek4 GGUF dispatch arm @@ -57,15 +59,97 @@ namespace fs = std::filesystem; // a failure to serve the named device PROPAGATES instead of falling back to // CPU (mirror of vLLM never substituting an explicitly named device, // vllm/config/device.py:61-66). -vt::Queue SelectQueueForModel(std::string_view architecture, - vllm::Device device) { +namespace { + +// The auto arm of the resolution below, WITHOUT creating a queue. Extracted so +// the queue selector and the load-time device-fit refusal (issue #1123) read one +// description of "which device will this model run on" rather than two that can +// drift. May throw, exactly as `CurrentPlatform()` can, and every caller keeps +// the try/catch the original code had around it. +vt::DeviceType AutoAcceleratorDeviceType(std::string_view architecture) { + const vllm::platforms::Platform& plat = vllm::platforms::CurrentPlatform(); + const vt::DeviceType dev = plat.device_type(); + // A PARTIAL backend (Metal today: 15 of 75 ops) must be able to decline a + // model whose kernels it has not registered. The default answer is `true`, + // so CUDA and CPU selection is byte-unchanged. + if (dev != vt::DeviceType::kCPU && + (architecture.empty() || plat.supports_model_architecture(architecture))) { + return dev; + } + return vt::DeviceType::kCPU; +} + +// The AUTO arm, resolved by ATTEMPTING the queue. One implementation, so +// `ResolveModelDeviceType` and `SelectQueueForModel` cannot answer differently. +// +// Asking `CurrentPlatform()` alone is not enough, and #1136 measured why. This +// arm has always fallen back to CPU when `CreateQueue()` throws — "a platform can +// be registered while CreateQueue still fails, and CPU must remain reachable" — +// so on such a box a platform query answers `kCUDA` while the load runs on the +// CPU queue. The load-time device-fit refusal reads the query, and it therefore +// refused a checkpoint by naming a device nothing was going to run on, removing a +// load that previously served on CPU. Whether `CreateQueue()` fails is knowable +// only by calling it, so it is called here, once, and the queue goes to whichever +// caller wants one. +struct AutoDeviceResolution { + vt::DeviceType device = vt::DeviceType::kCPU; + // Set exactly when `device != kCPU`: the queue whose creation PROVED it. + std::optional queue; +}; + +AutoDeviceResolution ResolveAutoDevice(std::string_view architecture) { + AutoDeviceResolution out; + try { + const vt::DeviceType dev = AutoAcceleratorDeviceType(architecture); + if (dev != vt::DeviceType::kCPU) { + // Order matters: `device` is set only AFTER the queue exists, so a throw + // leaves the CPU answer rather than a device nothing can serve. + vt::Queue q = vt::GetBackend(dev).CreateQueue(); + out.queue = q; + out.device = dev; + } + } catch (const std::exception&) { + // No usable accelerator; CPU, which is what this arm has always returned. + } + return out; +} + +} // namespace + +vt::DeviceType ResolveModelDeviceType(std::string_view architecture, + vllm::Device device) { if (device != vllm::Device::kAuto) { const vllm::platforms::Platform* named_platform = vllm::platforms::FindPlatformByName(vllm::DeviceName(device)); - const vt::DeviceType resolved = LoadedEngine::ResolveExplicitDeviceType( + // Propagates for an explicitly named absent device, which is the refusal + // vllm/config/device.py:61-66 mirrors and must not be swallowed here. + return LoadedEngine::ResolveExplicitDeviceType( device, named_platform == nullptr ? std::nullopt : std::optional{named_platform->device_type()}); + } + AutoDeviceResolution resolved = ResolveAutoDevice(architecture); + // The queue was created only to learn whether it CAN be created. `vt::Queue` is + // a NON-OWNING handle (a raw `cudaStream_t`) with no destructor, so dropping the + // value would leak the stream. + // + // Through the FREE `vt::DestroyQueue`, not `Backend::DestroyQueue`: that is what + // this file's only other queue teardown does (`load_queue`, below), it is what + // `vt/backend.h` asks of new code so device index and queue cleanup are never + // ambient, and it adds the `Synchronize` and the handle/id clearing the method + // does not. The CREATE side deliberately stays `GetBackend(...).CreateQueue()`, + // because that is the call this arm has always made and switching it would move + // the production queue-selection path onto the drop-in resource ABI — a + // behaviour change, which this repair is not. + if (resolved.queue.has_value()) vt::DestroyQueue(*resolved.queue); + return resolved.device; +} + +vt::Queue SelectQueueForModel(std::string_view architecture, + vllm::Device device) { + if (device != vllm::Device::kAuto) { + const vt::DeviceType resolved = + ResolveModelDeviceType(architecture, device); if (resolved == vt::DeviceType::kCPU) { return vt::Queue{vt::Device{vt::DeviceType::kCPU, 0}, nullptr}; } @@ -83,24 +167,14 @@ vt::Queue SelectQueueForModel(std::string_view architecture, // the single line that stood between the Metal backend and running a model. // It now asks the PLATFORM seam, which is the tree's own answer to "which // device is this process running on": CurrentPlatform() walks - // {kCUDA, kXPU, kVULKAN, kMETAL, kCPU} and returns the first whose backend - // actually probed a device (src/vllm/platforms/platform.cpp:38-40), so on a - // CUDA box this selects EXACTLY the queue the old code did, byte for byte, - // and on the M4 it selects Metal. The try/catch stays: a platform can be + // {kCUDA, kROCM, kXPU, kVULKAN, kMETAL, kTENSTORRENT, kCPU} and returns the + // first whose backend actually probed a device + // (src/vllm/platforms/platform.cpp:91-98), so on a CUDA box this selects + // EXACTLY the queue the old code did, byte for byte, and on the M4 it selects + // Metal. The try/catch stays, now inside `ResolveAutoDevice`: a platform can be // registered while CreateQueue still fails, and CPU must remain reachable. - try { - const vllm::platforms::Platform& plat = vllm::platforms::CurrentPlatform(); - const vt::DeviceType dev = plat.device_type(); - // A PARTIAL backend (Metal today: 15 of 75 ops) must be able to decline a - // model whose kernels it has not registered. The default answer is `true`, - // so CUDA and CPU selection is byte-unchanged. - if (dev != vt::DeviceType::kCPU && - (architecture.empty() || plat.supports_model_architecture(architecture))) { - return vt::GetBackend(dev).CreateQueue(); - } - } catch (const std::exception&) { - // No usable accelerator; fall through to CPU. - } + AutoDeviceResolution resolved = ResolveAutoDevice(architecture); + if (resolved.queue.has_value()) return *resolved.queue; return vt::Queue{vt::Device{vt::DeviceType::kCPU, 0}, nullptr}; } @@ -1313,7 +1387,36 @@ std::unique_ptr LoadedEngine::FromModelDir( // Resolve before tokenizer/weight work so unsupported architecture errors // are deterministic and match registry.py rather than being masked by a // later source-specific missing-tensor/tokenizer error. - (void)ModelRegistry::Resolve(config); + const ModelRegistration& gguf_arch = ModelRegistry::Resolve(config); + // Issue #1123: refuse a GGUF whose weights cannot be STAGED onto the target + // device, here, before any weight I/O and before the tokenizer. + // + // `Qwen3.8-2.4T-A95B UD-Q1_0` (369.96 GiB) reached a serving state on + // `--device cuda` on a 119.631 GiB GB10 after 26 minutes and then died on + // the FIRST forward with `vt cuda: cudaMalloc: out of memory`. The load + // succeeds because a keep-quant expert tower is BORROWED from this mapping + // and costs zero anonymous bytes; the forward dies because a + // weight-staging device copies each tower into device memory + // (`ResidentWeight`, qwen3_5.cpp:1011 -- 276 towers of 1,275,068,416 + // bytes plus 3 of 2,818,572,288, so 335.62 GiB). Loading for 26 minutes and dying + // mid-stream is the worst of the available behaviours. + // + // Placed AFTER Resolve so an unsupported-architecture error keeps its + // priority and the error ordering this branch documents is unchanged, and + // BEFORE the tokenizer and the weights because everything after this point + // is the cost the refusal exists to avoid paying. The predicate lives in + // `gguf_device_fit.h`; it decides nothing on a platform that does not stage + // weights (every CPU load) and nothing when no budget is known. + { + const platforms::Platform& target = platforms::GetPlatform( + ResolveModelDeviceType(gguf_arch.architecture, params.device)); + const DeviceWeightFit fit = CheckDeviceWeightFit( + gguf, vt::DeviceTypeName(target.device_type()), + target.needs_weight_staging(), + DeviceWeightBudgetBytes( + target.residency_policy().device_memory_total_bytes)); + if (fit.refuse) throw std::runtime_error(fit.message); + } tok::Tokenizer tokenizer = tok::Tokenizer::FromGguf(gguf); // Dense-vs-MoE GGUF dispatch now happens through the registry: the bench // branch's inline `IsDenseArch` split is superseded by diff --git a/src/vllm/model_executor/model_loader/gguf_device_fit.cpp b/src/vllm/model_executor/model_loader/gguf_device_fit.cpp new file mode 100644 index 000000000..56f152667 --- /dev/null +++ b/src/vllm/model_executor/model_loader/gguf_device_fit.cpp @@ -0,0 +1,117 @@ +// 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 + +namespace vllm { +namespace { + +// Bytes -> "N.NN GiB", so a refusal reads as a size rather than as 19 digits. +// Both the raw byte count and the GiB appear in the message: the first is what a +// reader can grep for in the code, the second is what an operator compares +// against the box. +std::string Gib(size_t bytes) { + const double gib = static_cast(bytes) / (1024.0 * 1024.0 * 1024.0); + std::string s = std::to_string(gib); + const size_t dot = s.find('.'); + if (dot != std::string::npos && s.size() > dot + 3) s.resize(dot + 3); + return s + " GiB"; +} + +} // namespace + +GgufStagedFootprint GgufStagedWeightFootprint(const GgufFile& gguf, + size_t model_dtype_bytes) { + GgufStagedFootprint out; + for (const GgufTensorInfo& t : gguf.Tensors()) { + size_t elems = 1; + for (const int64_t d : t.shape) { + if (d <= 0) { // A malformed dim cannot be reasoned about; contribute the + elems = 0; // on-disk size alone rather than a bogus expanded size. + break; + } + elems *= static_cast(d); + } + // The expanded size, when it is knowable. `elems == 0` means the shape was + // unusable, and then the on-disk size is the only defensible term. + const size_t expanded = elems == 0 ? t.nbytes : elems * model_dtype_bytes; + const size_t staged = expanded < t.nbytes ? expanded : t.nbytes; + out.lower_bound_bytes += staged; + ++out.tensor_count; + if (staged > out.largest_tensor_bytes) { + out.largest_tensor_bytes = staged; + out.largest_tensor_name = t.name; + } + } + return out; +} + +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; +} + +DeviceWeightFit CheckDeviceWeightFit(const GgufFile& gguf, + std::string_view device_name, + bool needs_weight_staging, + size_t budget_bytes, + size_t model_dtype_bytes) { + DeviceWeightFit fit; + fit.budget_bytes = budget_bytes; + // A platform that does not stage weights reads them where they already are, so + // a borrowed tower costs it nothing and there is nothing to compare. This is + // the branch every CPU load takes, and it must be free of any behaviour + // change: no footprint is even computed. + if (!needs_weight_staging) return fit; + // 0 == UNKNOWN, and unknown is not a verdict. Refusing a load because nothing + // reported a budget would break every device whose budget nothing probes. + if (budget_bytes == 0) return fit; + + const GgufStagedFootprint fp = + GgufStagedWeightFootprint(gguf, model_dtype_bytes); + fit.needed_bytes = fp.lower_bound_bytes; + if (fp.lower_bound_bytes <= budget_bytes) return fit; + + fit.refuse = true; + fit.message = + "device '" + std::string(device_name) + + "' cannot serve this GGUF: staging its weights needs at least " + + std::to_string(fp.lower_bound_bytes) + " bytes (" + + Gib(fp.lower_bound_bytes) + ") of device memory across " + + std::to_string(fp.tensor_count) + " tensors, the largest single " + "allocation being " + std::to_string(fp.largest_tensor_bytes) + " bytes (" + + Gib(fp.largest_tensor_bytes) + ", '" + fp.largest_tensor_name + + "'), and this device's memory pool is " + std::to_string(budget_bytes) + + " bytes (" + Gib(budget_bytes) + "). THE MISSING PART: the " + "larger-than-memory lane that makes a checkpoint like this fit is " + "HOST-ONLY. The GGUF mapping is borrowed in place on the CPU path and " + "costs no resident bytes, while a weight-staging device copies every " + "expert tower into device memory; there is no device-side expert slot " + "store and no device streaming lane (ENG-EXPERT-STREAM, issues #1123 and " + "#1124). Use device=cpu, which serves this checkpoint today, or a " + "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 " + "and restores that late failure; it does not make the model fit."; + return fit; +} + +} // namespace vllm diff --git a/src/vllm/model_executor/models/gemma4_moe.cpp b/src/vllm/model_executor/models/gemma4_moe.cpp index 238d09669..ede647835 100644 --- a/src/vllm/model_executor/models/gemma4_moe.cpp +++ b/src/vllm/model_executor/models/gemma4_moe.cpp @@ -437,7 +437,15 @@ struct DevExpertLru { bool Enabled() { return BudgetBytes() > 0; } - // Free VRAM via Backend::DeviceMemoryInfo (ROCm/CUDA). No HIP in this TU. + // Free VRAM via Backend::DeviceMemoryInfo. No HIP in this TU. + // + // ROCm ONLY: `CudaBackend` does not override that seam, so this returns false on + // every CUDA device and `MakeRoom` below then refuses the device upload, which + // makes this whole cache dead on CUDA today. That is issue #1126, not an + // accident of this call site — the refuse-on-unknown polarity here is correct, + // because an Alloc without headroom has hung hipMalloc. This comment said + // "(ROCm/CUDA)" until #1123 measured it (the same false claim as the one on + // `vt::Backend::DeviceMemoryInfo` itself). static bool FreeBytes(Dev d, size_t* free_out) { *free_out = 0; size_t free_b = 0, tot_b = 0; diff --git a/src/vllm/platforms/cuda.cpp b/src/vllm/platforms/cuda.cpp index 9eb934f11..d1436be9a 100644 --- a/src/vllm/platforms/cuda.cpp +++ b/src/vllm/platforms/cuda.cpp @@ -5,6 +5,7 @@ // usable GPU). Compiled only in CUDA builds (CMake target_sources gate). #include +#include #include #include "vllm/platforms/cuda_attn_priority.h" @@ -17,8 +18,11 @@ namespace { class CudaPlatform final : public Platform { public: - CudaPlatform(int cc_major, int cc_minor, bool integrated) - : cap_{cc_major, cc_minor}, integrated_{integrated} {} + CudaPlatform(int cc_major, int cc_minor, bool integrated, + size_t device_memory_total_bytes) + : cap_{cc_major, cc_minor}, + integrated_{integrated}, + device_memory_total_bytes_{device_memory_total_bytes} {} DeviceType device_type() const override { return DeviceType::kCUDA; } Backend& backend() const override { return vt::GetBackend(DeviceType::kCUDA); } @@ -89,14 +93,21 @@ class CudaPlatform final : public Platform { // still overrides (house A/B convention). // * uses_device_memory_pool = true + device_pool_cap_bytes = 0: the DevicePool // scratch reuse, uncapped, exactly as today. + // * device_memory_total_bytes = cudaMemGetInfo's `total`, probed at + // registration (issue #1123). NEW data, consumed only by the load-time + // GGUF fit refusal; nothing that read this struct before sees a change. // A discrete GPU sets different values (e.g. a pool cap) and NO model code is // touched — that is the item-2 additive win. + // + // The four assignments themselves live in `CudaResidencyPolicy` + // (`vllm/platforms/interface.h`), not here, because this translation unit + // compiles only in a CUDA build: while they were inline, nothing on a host + // without a CUDA toolkit could reach them, which is why #1123 had to record + // "delete the device_memory_total_bytes assignment" as an unproven mutation. + // test_platform.cpp now pins the assembly on every host (#1136). What stays + // CUDA-only here is the `cudaMemGetInfo` probe below and the value it threads. ResidencyPolicy residency_policy() const override { - ResidencyPolicy p; - p.release_host_weights_after_upload = true; // freed after Marlin build (today) - p.uses_device_memory_pool = true; // qwen3_5.cpp DevicePool - p.device_pool_cap_bytes = 0; // uncapped - return p; + return CudaResidencyPolicy(device_memory_total_bytes_); } // Capability-ordered attention-backend priority — a faithful port of @@ -125,6 +136,8 @@ class CudaPlatform final : public Platform { private: DeviceCapability cap_; bool integrated_ = false; + // cudaMemGetInfo's `total`, probed once at registration; 0 == UNKNOWN. + size_t device_memory_total_bytes_ = 0; }; // Registers kCUDA during static init (registration must complete before main() @@ -151,13 +164,31 @@ struct Registrar { if (cudaDeviceGetAttribute(&integrated, cudaDevAttrIntegrated, 0) != cudaSuccess) { integrated = 0; } + // ResidencyPolicy::device_memory_total_bytes (issue #1123) — probe once here, + // beside the other device probes. `nvidia-smi` is the WRONG instrument for + // this on a GB10: `--query-gpu=memory.total,memory.free,memory.used` answers + // `[N/A], [N/A], [N/A]` because host and device share one pool, and the `rc` + // fleet label records `vram=[N/A]M` for the same reason. `cudaMemGetInfo` + // answers honestly. Measured on dgx:gpu0 through libcudart.so.13: + // total = 128452956160 (119.631 GiB), free = 122059919360 (113.677 GiB), + // and `total` is EXACTLY `/proc/meminfo MemTotal` (125442340 kB) times 1024. + // + // A query failure leaves 0 = UNKNOWN, which the consumer treats as "do not + // decide" rather than as "nothing fits". `free_bytes` is read and discarded: + // this is a load-time budget, and `free` makes it a function of contention. + size_t total_bytes = 0; + size_t free_bytes = 0; + if (cudaMemGetInfo(&free_bytes, &total_bytes) != cudaSuccess) { + total_bytes = 0; + } // GCC 13 false-positive: -Wdangling-pointer mis-flags a static local with a // vtable constructed from automatic ints, though CudaPlatform copies both // into cap_ by value (no pointer/reference to major/minor is retained). The // static outlives the registrar as RegisterPlatform requires. #pragma GCC diagnostic push #pragma GCC diagnostic ignored "-Wdangling-pointer" - static CudaPlatform platform(major, minor, integrated != 0); // device 0 only + static CudaPlatform platform(major, minor, integrated != 0, + total_bytes); // device 0 only #pragma GCC diagnostic pop RegisterPlatform(DeviceType::kCUDA, &platform); } diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt index b85c0509b..21214ac89 100644 --- a/tests/CMakeLists.txt +++ b/tests/CMakeLists.txt @@ -384,6 +384,16 @@ target_include_directories(test_model_registry PRIVATE ${CMAKE_SOURCE_DIR}/src) vllm_cpp_add_test(test_platform vllm/platforms/test_platform.cpp) vllm_cpp_add_test(test_device_selection vllm/entrypoints/test_device_selection.cpp) +# ENG-EXPERT-STREAM #1123: the load-time GGUF device-fit refusal, in two targets +# for two different questions. The arithmetic runs everywhere; the REACHABILITY +# gate needs a fake WEIGHT-STAGING platform in the CUDA lookup slot, because +# `needs_weight_staging()` is true on exactly one platform in this tree and the +# branch is otherwise unreachable on a host with no CUDA device. Its own +# executable for the same reason test_device_selection is one. +vllm_cpp_add_test(test_gguf_device_fit + vllm/model_executor/test_gguf_device_fit.cpp) +vllm_cpp_add_test(test_gguf_device_fit_reach + vllm/entrypoints/test_gguf_device_fit_reach.cpp) vllm_cpp_add_test(test_linear_method vllm/model_executor/layers/test_linear_method.cpp) target_include_directories(test_linear_method PRIVATE ${CMAKE_SOURCE_DIR}/src) vllm_cpp_add_test(test_decode_graph_sizes vllm/models/test_decode_graph_sizes.cpp) diff --git a/tests/vllm/entrypoints/test_gguf_device_fit_reach.cpp b/tests/vllm/entrypoints/test_gguf_device_fit_reach.cpp new file mode 100644 index 000000000..fff17d21f --- /dev/null +++ b/tests/vllm/entrypoints/test_gguf_device_fit_reach.cpp @@ -0,0 +1,318 @@ +// ENG-EXPERT-STREAM, issue #1123 — the REACHABILITY gate for the load-time +// device-fit refusal. The arithmetic is gated in test_gguf_device_fit; this file +// answers the different question that tree has carried green before: does the +// production loader actually ask? +// +// A test that constructs the predicate by hand proves the predicate works and +// never proves anything reaches it. So this drives +// `LoadedEngine::FromModelDir`, the loader entry point every consumer uses, and +// asserts the thrown MESSAGE. Deleting the call site in `model_loader.cpp` makes +// the refusing case throw the LATER tokenizer error instead, which is red here. +// +// Why a fake platform. `needs_weight_staging()` is true on exactly one platform +// in this tree (`src/vllm/platforms/cuda.cpp:71`), so on a host with no CUDA device +// the branch is unreachable from the real loader — the untestable-device-branch +// shape this row has hit repeatedly. A fake staging platform registered in the +// CUDA lookup slot reaches it, which is the instrument +// `tests/vllm/entrypoints/test_device_selection.cpp` established for exactly +// this reason. It is a SEPARATE executable for the same reason that one is: +// registering into the global platform/backend registries must not leak into +// other suites. +#include + +#include +#include +#include +#include +#include + +#include "support/test_env.h" +#include "vllm/config/device.h" +#include "vllm/entrypoints/model_loader.h" +#include "vllm/gguf_builder.h" +#include "vllm/platforms/interface.h" +#include "vt/backend.h" + +namespace { + +using gguf_test::GgufModelBuilder; +using gguf_test::StrKv; +using gguf_test::TempFile; +using gguf_test::U32Kv; + +// A backend that allocates on the host. Nothing in this file runs a forward; it +// exists so the fake platform has a `Backend&` to return, and so +// `SelectQueueForModel` has something to create a queue from if it is ever +// asked (this file never gets that far — the refusal fires first, and on the +// permitting arm the tokenizer throws first). +class HostBackend final : public vt::Backend { + public: + void* Alloc(size_t bytes) override { + return std::malloc(bytes == 0 ? 1 : bytes); + } + void Free(void* p) override { std::free(p); } + void Memset(vt::Queue&, void* p, int value, size_t bytes) override { + std::memset(p, value, bytes); + } + void Copy(vt::Queue&, void* dst, const void* src, size_t bytes) override { + std::memcpy(dst, src, bytes); + } + // "A platform can be registered while CreateQueue still fails" is the reason + // `SelectQueueForModel`'s AUTO arm wraps this call in a try/catch and falls + // back to CPU (`ResolveAutoDevice`, `model_loader.cpp:100-115`). This flag + // reproduces that box on a + // host with no CUDA device, so the resolver the fit refusal reads can be + // checked against the queue the load will actually run on. A flag rather than + // a second backend, because the registry is global and process-wide: a second + // registration would leak into the other cases in this file. + bool create_queue_throws = false; + vt::Queue CreateQueue() override { + if (create_queue_throws) { + throw std::runtime_error("fake backend: no usable device"); + } + ++queues_created; + return vt::Queue{vt::Device{vt::DeviceType::kCUDA, 0}, nullptr}; + } + // Counted so a case can prove the resolver ATTEMPTED a queue rather than + // merely returning the same answer for another reason, and that it handed the + // probe queue back. `vt::Queue` is a non-owning handle with no destructor, so a + // resolver that dropped the value would leak the stream and nothing would say + // so; these two counters are what makes that observable on a fake backend. + int queues_created = 0; + int queues_destroyed = 0; + void DestroyQueue(vt::Queue&) override { ++queues_destroyed; } + bool UnifiedMemory() const override { return true; } +}; + +// The one property under test: a platform that STAGES weights, carrying a +// budget on its residency policy exactly as `CudaPlatform` now does. +class StagingPlatform final : public vllm::platforms::Platform { + public: + StagingPlatform(HostBackend& backend, size_t budget) + : backend_(backend), budget_(budget) {} + + vt::DeviceType device_type() const override { return vt::DeviceType::kCUDA; } + vt::Backend& backend() const override { return backend_; } + vllm::platforms::DeviceCapability get_device_capability() const override { + return {12, 1}; + } + std::vector supported_dtypes() const override { + return {vt::DType::kBF16}; + } + bool needs_weight_staging() const override { return true; } + vllm::platforms::ResidencyPolicy residency_policy() const override { + vllm::platforms::ResidencyPolicy p; + p.device_memory_total_bytes = budget_; + return p; + } + + private: + HostBackend& backend_; + size_t budget_ = 0; +}; + +HostBackend& Backend() { + static HostBackend backend; + return backend; +} + +// The budget is deliberately supplied through the POLICY here (0), and moved by +// `VT_DEVICE_WEIGHT_BUDGET_BYTES` in each case, so both halves of +// `DeviceWeightBudgetBytes` are exercised through the production path: the +// unknown-policy arm and the override arm. +StagingPlatform& Platform() { + static StagingPlatform platform(Backend(), /*budget=*/0); + return platform; +} + +void RegisterFakeStagingPlatform() { + vt::RegisterBackend(vt::DeviceType::kCUDA, &Backend()); + vllm::platforms::RegisterPlatform(vt::DeviceType::kCUDA, &Platform()); +} + +// A synthetic `qwen35moe` GGUF: enough hparams for `HfConfigFromGguf` and +// `ModelRegistry::Resolve` to succeed, so the fit check is reached at its real +// position in the ladder (AFTER architecture resolution) and BEFORE any weight +// I/O. It carries no tokenizer, so the arm that is ALLOWED through fails LATER +// and DIFFERENTLY -- measured, not assumed: +// +// tokenizer: GGUF missing kv "tokenizer.ggml.model" +// +// That is the NEXT step after the check, and it is what makes the permitting +// case meaningful: the load got past the check. +// +// Its total staged footprint is small and asserted below rather than assumed. +std::string BuildSyntheticMoeGguf() { + GgufModelBuilder b; + b.AddKv(StrKv("general.architecture", "qwen35moe")); + b.AddKv(U32Kv("qwen35moe.embedding_length", 64)); + b.AddKv(U32Kv("qwen35moe.block_count", 2)); + b.AddKv(U32Kv("qwen35moe.attention.head_count", 4)); + b.AddKv(U32Kv("qwen35moe.attention.head_count_kv", 2)); + b.AddKv(U32Kv("qwen35moe.attention.key_length", 16)); + b.AddKv(U32Kv("qwen35moe.expert_count", 4)); + b.AddKv(U32Kv("qwen35moe.expert_used_count", 2)); + b.AddKv(U32Kv("qwen35moe.expert_feed_forward_length", 32)); + b.AddKv(U32Kv("qwen35moe.expert_shared_feed_forward_length", 32)); + b.AddKv(U32Kv("qwen35moe.ssm.group_count", 2)); + b.AddKv(U32Kv("qwen35moe.ssm.time_step_rank", 4)); + b.AddKv(U32Kv("qwen35moe.ssm.state_size", 8)); + b.AddKv(U32Kv("qwen35moe.ssm.conv_kernel", 4)); + b.AddKv(U32Kv("qwen35moe.full_attention_interval", 4)); + b.AddKv(U32Kv("qwen35moe.context_length", 256)); + // Both of these are REQUIRED by HfConfigFromGguf (`ReqFloat`, + // qwen3_5_gguf_weights.cpp:843-847), which runs BEFORE the fit check. Omitting + // them made the load throw "missing metadata key" during the config parse and + // the refusing case never reached the check at all — caught because the case + // asserted the MESSAGE rather than merely that something threw. + b.AddKv(gguf_test::F32Kv("qwen35moe.rope.freq_base", 1000000.0F)); + b.AddKv(gguf_test::F32Kv("qwen35moe.attention.layer_norm_rms_epsilon", 1e-6F)); + // One F32 tensor of 4096 elements: 16384 bytes on disk, 8192 expanded to + // bf16, so the staged lower bound is 8192. + b.AddTensor("token_embd.weight", {64, 64}, /*ggml_type=*/0, + std::string(4096 * 4, '\0')); + return b.Build(); +} + +constexpr size_t kStagedLowerBound = 8192; + +std::string ThrownMessage(const std::string& gguf_path, vllm::Device device) { + vllm::entrypoints::EngineParams params; + params.device = device; + try { + (void)vllm::entrypoints::LoadedEngine::FromModelDir(gguf_path, params); + } catch (const std::exception& e) { + return e.what(); + } + return std::string(); +} + +} // namespace + +TEST_CASE("device fit: the loader REFUSES a GGUF that exceeds the staging budget") { + RegisterFakeStagingPlatform(); + TempFile f(BuildSyntheticMoeGguf()); + + // One byte under the footprint. Chosen at the boundary so the case cannot pass + // by accident on an implementation that compares the wrong quantity. + vllm_test::SetEnv("VT_DEVICE_WEIGHT_BUDGET_BYTES", + std::to_string(kStagedLowerBound - 1)); + const std::string message = ThrownMessage(f.path(), vllm::Device::kNamedPlatform); + vllm_test::UnsetEnv("VT_DEVICE_WEIGHT_BUDGET_BYTES"); + + REQUIRE_FALSE(message.empty()); + CAPTURE(message); + // The refusal, not some other failure: the message names the device, the + // measured need, the budget it exceeded, and the missing capability. + 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); + CHECK(message.find("HOST-ONLY") != std::string::npos); + CHECK(message.find("device=cpu") != std::string::npos); + // And it fires BEFORE the tokenizer and therefore before any weight I/O, which + // is the whole point of refusing at load: everything after this point is the + // 26 minutes the refusal exists to avoid paying. + CHECK(message.find("tokenizer") == std::string::npos); +} + +TEST_CASE("device fit: a GGUF that FITS the budget is let through to the next stage") { + RegisterFakeStagingPlatform(); + TempFile f(BuildSyntheticMoeGguf()); + + // Exactly the footprint: the boundary on the permitting side, so a mutation + // that turns `>` into `>=` is red here rather than merely unnoticed. + vllm_test::SetEnv("VT_DEVICE_WEIGHT_BUDGET_BYTES", + std::to_string(kStagedLowerBound)); + const std::string message = ThrownMessage(f.path(), vllm::Device::kNamedPlatform); + vllm_test::UnsetEnv("VT_DEVICE_WEIGHT_BUDGET_BYTES"); + + // It still throws — this synthetic file carries no tokenizer — and that is what + // makes the case meaningful: the throw is a DIFFERENT one, from the step AFTER + // the check, which proves the check let it through rather than that it never + // ran. Asserting the later message positively is the point: a case that only + // asserted the absence of the refusal would also pass if the loader had died + // earlier for an unrelated reason, which is exactly how the first draft of + // this file passed while the config parse was throwing. + REQUIRE_FALSE(message.empty()); + CAPTURE(message); + CHECK(message.find("cannot serve this GGUF") == std::string::npos); + CHECK(message.find("HOST-ONLY") == std::string::npos); + CHECK(message.find("tokenizer: GGUF missing kv") != std::string::npos); +} + +// --- The AUTO arm: the refusal must name the device the load will RUN on ------ +// +// `SelectQueueForModel`'s auto arm falls back to CPU when `CreateQueue()` throws, +// and its own comment says why: "a platform can be registered while CreateQueue +// still fails, and CPU must remain reachable". A resolver that only asked +// `CurrentPlatform()` answered `kCUDA` on such a box, so the fit refusal REFUSED +// a checkpoint by naming a device nothing was going to run on — a load that +// previously served on CPU. These two cases are the pair: the same file, the same +// budget, the same platform, differing only in whether the queue can be created. +TEST_CASE("device fit: the AUTO arm refuses when the accelerator queue CAN be created") { + RegisterFakeStagingPlatform(); + TempFile f(BuildSyntheticMoeGguf()); + Backend().create_queue_throws = false; + const int created_before = Backend().queues_created; + const int destroyed_before = Backend().queues_destroyed; + + vllm_test::SetEnv("VT_DEVICE_WEIGHT_BUDGET_BYTES", + std::to_string(kStagedLowerBound - 1)); + const std::string message = ThrownMessage(f.path(), vllm::Device::kAuto); + vllm_test::UnsetEnv("VT_DEVICE_WEIGHT_BUDGET_BYTES"); + + // The POSITIVE control for the case below. Without it, "no refusal" there could + // mean the auto arm never selects the fake platform at all, and the pair would + // prove nothing. + REQUIRE_FALSE(message.empty()); + CAPTURE(message); + CHECK(message.find("cannot serve this GGUF") != std::string::npos); + CHECK(message.find("HOST-ONLY") != std::string::npos); + // The resolution went through an ATTEMPTED queue rather than a bare platform + // query, which is the only way it can agree with the queue selector. + CHECK(Backend().queues_created == created_before + 1); + // And it gave the probe queue back. The refusal throws before + // `SelectQueueForModel` runs, so this load creates exactly one queue and + // destroys exactly one: a resolver that leaked it reads 1 created, 0 destroyed. + CHECK(Backend().queues_destroyed == destroyed_before + 1); +} + +TEST_CASE("device fit: the AUTO arm refuses NOTHING when the accelerator queue cannot be created") { + RegisterFakeStagingPlatform(); + TempFile f(BuildSyntheticMoeGguf()); + Backend().create_queue_throws = true; + + vllm_test::SetEnv("VT_DEVICE_WEIGHT_BUDGET_BYTES", + std::to_string(kStagedLowerBound - 1)); + const std::string message = ThrownMessage(f.path(), vllm::Device::kAuto); + vllm_test::UnsetEnv("VT_DEVICE_WEIGHT_BUDGET_BYTES"); + Backend().create_queue_throws = false; + + REQUIRE_FALSE(message.empty()); + CAPTURE(message); + // This load runs on CPU, and the CPU platform does not stage weights, so there + // is nothing to refuse. Asserting the LATER tokenizer error positively rather + // than merely the absence of the refusal: a case that only checked the absence + // would also pass if the loader had died earlier for an unrelated reason. + CHECK(message.find("cannot serve this GGUF") == std::string::npos); + CHECK(message.find("HOST-ONLY") == std::string::npos); + CHECK(message.find("tokenizer: GGUF missing kv") != std::string::npos); +} + +TEST_CASE("device fit: an explicit CPU load is never refused, at any budget") { + RegisterFakeStagingPlatform(); + TempFile f(BuildSyntheticMoeGguf()); + + // A budget of one byte, which every checkpoint exceeds. The CPU platform does + // not stage weights, so the predicate must not even look — this is the arm + // that keeps `--device cpu` byte-identical, and it is the arm the measured + // 370 GiB checkpoint actually serves on. + vllm_test::SetEnv("VT_DEVICE_WEIGHT_BUDGET_BYTES", "1"); + const std::string message = ThrownMessage(f.path(), vllm::Device::kCPU); + vllm_test::UnsetEnv("VT_DEVICE_WEIGHT_BUDGET_BYTES"); + + REQUIRE_FALSE(message.empty()); // no tokenizer, as above + CAPTURE(message); + CHECK(message.find("cannot serve this GGUF") == std::string::npos); + CHECK(message.find("tokenizer: GGUF missing kv") != std::string::npos); +} diff --git a/tests/vllm/model_executor/test_gguf_device_fit.cpp b/tests/vllm/model_executor/test_gguf_device_fit.cpp new file mode 100644 index 000000000..c12ea473c --- /dev/null +++ b/tests/vllm/model_executor/test_gguf_device_fit.cpp @@ -0,0 +1,252 @@ +// ENG-EXPERT-STREAM, issue #1123. The arithmetic and the predicate behind the +// load-time refusal of a GGUF whose weights cannot be staged onto the target +// device. The REACHABILITY half — that the loader actually asks — is a separate +// binary, test_gguf_device_fit_reach, because it has to register a fake staging +// platform in a global registry. +// +// Why the numbers here are the ones they are: a Q8_0 block is 34 bytes per 32 +// elements, so `elems * 2` (bf16) is 64 and the on-disk size is the smaller +// term; an F32 tensor is 4 bytes per element, so `elems * 2` is the smaller +// term. One file with both therefore pins BOTH arms of +// `min(gguf_bytes, elems * model_dtype_bytes)` in a single sum, and a mutation +// that drops either arm changes it. +#include + +#include +#include +#include +#include + +#include "support/test_env.h" +#include "vllm/gguf_builder.h" +#include "vllm/model_executor/model_loader/gguf_device_fit.h" +#include "vllm/model_executor/model_loader/gguf_reader.h" + +namespace { + +using gguf_test::GgufModelBuilder; +using gguf_test::StrKv; +using gguf_test::TempFile; + +// One Q8_0 block: f16 scale + 32 int8 quants = 34 bytes for 32 elements. +std::string Q8Block() { + std::string b(2, '\0'); + b[0] = '\0'; + b[1] = '\x3c'; // f16 1.0, little-endian + for (int i = 0; i < 32; ++i) b.push_back(static_cast(i)); + return b; +} + +// A GGUF with exactly two tensors: +// "t_q8" Q8_0, 32 elements -> 34 bytes on disk, 64 bytes expanded to bf16 +// "t_f32" F32, 8 elements -> 32 bytes on disk, 16 bytes expanded to bf16 +// So the staged lower bound is min(34,64) + min(32,16) = 34 + 16 = 50. +constexpr size_t kExpectedLowerBound = 50; +constexpr size_t kExpectedTensors = 2; +constexpr size_t kExpectedLargest = 34; // "t_q8" + +std::string BuildTwoTensorGguf() { + GgufModelBuilder b; + b.AddKv(StrKv("general.architecture", "llama")); + b.AddTensor("t_q8", {32}, /*ggml_type=*/8, Q8Block()); + b.AddTensor("t_f32", {4, 2}, /*ggml_type=*/0, std::string(32, '\1')); + return b.Build(); +} + +// The same file plus one tensor of the MTP / `nextn` block, which a DEFAULT load +// never stages: the head is attached only under +// `params.speculative_config.has_value() && method == "mtp"` +// (`src/vllm/entrypoints/model_loader.cpp:1452-1453`), and the main model reads +// `block_count - nextn_predict_layers` blocks +// (`qwen3_5_gguf_weights.cpp:877-878`), so the head's blocks are outside its +// range. The footprint counts it anyway, because it takes the whole tensor +// table. A second Q8_0 block, so its staged term is min(34, 64) = 34. +constexpr size_t kNextnStaged = 34; + +std::string BuildGgufWithUnstagedNextnBlock() { + GgufModelBuilder b; + b.AddKv(StrKv("general.architecture", "llama")); + b.AddTensor("t_q8", {32}, /*ggml_type=*/8, Q8Block()); + b.AddTensor("t_f32", {4, 2}, /*ggml_type=*/0, std::string(32, '\1')); + b.AddTensor("blk.1.nextn.eh_proj.weight", {32}, /*ggml_type=*/8, Q8Block()); + return b.Build(); +} + +} // namespace + +TEST_CASE("gguf_device_fit: the footprint takes min(on-disk, expanded) per tensor") { + TempFile f(BuildTwoTensorGguf()); + const vllm::GgufFile gguf = vllm::GgufFile::Open(f.path()); + REQUIRE(gguf.Tensors().size() == kExpectedTensors); + + const vllm::GgufStagedFootprint fp = vllm::GgufStagedWeightFootprint(gguf); + // The count is asserted, not assumed: a bound that cannot say how many + // tensors it examined has not reported one. + CHECK(fp.tensor_count == kExpectedTensors); + CHECK(fp.lower_bound_bytes == kExpectedLowerBound); + CHECK(fp.largest_tensor_bytes == kExpectedLargest); + CHECK(fp.largest_tensor_name == "t_q8"); + + // Summing is enough BECAUSE the two arms disagree: an implementation that + // always took the on-disk size would give 34 + 32 = 66, one that always + // expanded would give 64 + 16 = 80, and both differ from 50. That is why the + // fixture carries one tensor of each kind rather than two of one kind. +} + +TEST_CASE("gguf_device_fit: a wider model dtype cannot raise the on-disk term") { + TempFile f(BuildTwoTensorGguf()); + const vllm::GgufFile gguf = vllm::GgufFile::Open(f.path()); + // f32 model dtype: the Q8_0 term stays 34 (min against 128) and the F32 term + // becomes min(32, 32) = 32. So the sum moves to 66 and NOT to 34 + 128. + const vllm::GgufStagedFootprint fp = + vllm::GgufStagedWeightFootprint(gguf, /*model_dtype_bytes=*/4); + CHECK(fp.tensor_count == kExpectedTensors); + CHECK(fp.lower_bound_bytes == 66); +} + +// The bound's ONE over-count direction, made executable rather than only +// described. Every other case in this file runs on a fixture whose tensors are +// all staged, so the footprint there happens to EQUAL the true staged size and +// the boundary cases cannot tell an exact quantity from an over-counted one. +// +// This case separates them. It exists because the spec and the commit body for +// this change asserted the bound was a lower bound "so the refusal can never +// over-refuse", which is false: a tensor counted and never staged is a positive +// over-count, and one is present on every default load. On the measured +// checkpoint that is the `nextn` block, 8,940,488,704 of 397,245,341,184 bytes +// (2.2506 %). The two error directions are on DIFFERENT quantities and do not +// cancel, so "the under-count dominates" does not rescue the claim. Recorded and +// owned by issue #1136; the header states the direction, and this pins it. +// +// NOTE for whoever closes #1136 by teaching the bound which tensors this load will +// stage: this case is SUPPOSED to go red then, and it is not an obstacle. It +// characterises today's contract, so changing the contract means changing it here +// too — deliberately, in the same commit, rather than discovering later that the +// bound quietly stopped counting something. +TEST_CASE("gguf_device_fit: a tensor the loader never stages is COUNTED, so the bound can over-refuse") { + TempFile f(BuildGgufWithUnstagedNextnBlock()); + const vllm::GgufFile gguf = vllm::GgufFile::Open(f.path()); + REQUIRE(gguf.Tensors().size() == kExpectedTensors + 1); + + const vllm::GgufStagedFootprint fp = vllm::GgufStagedWeightFootprint(gguf); + CHECK(fp.tensor_count == kExpectedTensors + 1); + CHECK(fp.lower_bound_bytes == kExpectedLowerBound + kNextnStaged); + + // The consequence at the boundary: a default load stages + // `kExpectedLowerBound` bytes, the predicate compares + // `kExpectedLowerBound + kNextnStaged`, and every budget in between refuses a + // weight set that fits. Both ends of that window are asserted, so a change + // that narrowed or widened the over-count moves this case. + for (const size_t budget : + {kExpectedLowerBound, kExpectedLowerBound + kNextnStaged - 1}) { + CAPTURE(budget); + const vllm::DeviceWeightFit fit = + vllm::CheckDeviceWeightFit(gguf, "cuda", true, budget); + CHECK(fit.refuse); + CHECK(fit.needed_bytes == kExpectedLowerBound + kNextnStaged); + } + // At the counted total it does not refuse, which pins the over-count to + // exactly this tensor and nothing more. + const vllm::DeviceWeightFit ok = vllm::CheckDeviceWeightFit( + gguf, "cuda", true, kExpectedLowerBound + kNextnStaged); + CHECK_FALSE(ok.refuse); +} + +TEST_CASE("gguf_device_fit: a non-staging platform is never refused, at any budget") { + TempFile f(BuildTwoTensorGguf()); + const vllm::GgufFile gguf = vllm::GgufFile::Open(f.path()); + for (const size_t budget : {size_t{0}, size_t{1}, kExpectedLowerBound - 1, + kExpectedLowerBound, size_t{1} << 40}) { + CAPTURE(budget); + const vllm::DeviceWeightFit fit = vllm::CheckDeviceWeightFit( + gguf, "cpu", /*needs_weight_staging=*/false, budget); + CHECK_FALSE(fit.refuse); + CHECK(fit.message.empty()); + // Nothing is even computed on this arm, which is what makes every CPU load + // byte-identical to before. + CHECK(fit.needed_bytes == 0); + } +} + +TEST_CASE("gguf_device_fit: an unknown budget is not a verdict") { + TempFile f(BuildTwoTensorGguf()); + const vllm::GgufFile gguf = vllm::GgufFile::Open(f.path()); + const vllm::DeviceWeightFit fit = vllm::CheckDeviceWeightFit( + gguf, "cuda", /*needs_weight_staging=*/true, /*budget_bytes=*/0); + CHECK_FALSE(fit.refuse); + CHECK(fit.message.empty()); + CHECK(fit.budget_bytes == 0); +} + +TEST_CASE("gguf_device_fit: refuses strictly above the budget, and not at or below it") { + TempFile f(BuildTwoTensorGguf()); + const vllm::GgufFile gguf = vllm::GgufFile::Open(f.path()); + + SUBCASE("one byte under the footprint refuses") { + const vllm::DeviceWeightFit fit = vllm::CheckDeviceWeightFit( + gguf, "cuda", true, kExpectedLowerBound - 1); + CHECK(fit.refuse); + CHECK(fit.needed_bytes == kExpectedLowerBound); + CHECK(fit.budget_bytes == kExpectedLowerBound - 1); + } + SUBCASE("exactly the footprint does NOT refuse") { + const vllm::DeviceWeightFit fit = + vllm::CheckDeviceWeightFit(gguf, "cuda", true, kExpectedLowerBound); + CHECK_FALSE(fit.refuse); + CHECK(fit.needed_bytes == kExpectedLowerBound); + } + SUBCASE("a generous budget does NOT refuse") { + const vllm::DeviceWeightFit fit = + vllm::CheckDeviceWeightFit(gguf, "cuda", true, size_t{1} << 40); + CHECK_FALSE(fit.refuse); + CHECK(fit.message.empty()); + } +} + +TEST_CASE("gguf_device_fit: the refusal names the device, both numbers, the missing part and the remedy") { + TempFile f(BuildTwoTensorGguf()); + const vllm::GgufFile gguf = vllm::GgufFile::Open(f.path()); + const vllm::DeviceWeightFit fit = + vllm::CheckDeviceWeightFit(gguf, "cuda", true, /*budget_bytes=*/8); + REQUIRE(fit.refuse); + const std::string& m = fit.message; + // A refusal that does not name what is missing is the behaviour this change + // exists to replace, so each half is asserted rather than the message length. + CHECK(m.find("device 'cuda'") != std::string::npos); + CHECK(m.find(std::to_string(kExpectedLowerBound)) != std::string::npos); + CHECK(m.find("8 bytes") != std::string::npos); + CHECK(m.find("t_q8") != std::string::npos); + CHECK(m.find("HOST-ONLY") != std::string::npos); + CHECK(m.find("device=cpu") != std::string::npos); + CHECK(m.find("#1123") != std::string::npos); + CHECK(m.find("VT_DEVICE_WEIGHT_BUDGET_BYTES") != std::string::npos); +} + +TEST_CASE("gguf_device_fit: VT_DEVICE_WEIGHT_BUDGET_BYTES overrides, and a typo does NOT disable the guard") { + SUBCASE("unset: the platform's probe is the budget") { + vllm_test::UnsetEnv("VT_DEVICE_WEIGHT_BUDGET_BYTES"); + CHECK(vllm::DeviceWeightBudgetBytes(4096) == 4096); + CHECK(vllm::DeviceWeightBudgetBytes(0) == 0); + } + SUBCASE("set: the override wins in both directions") { + vllm_test::SetEnv("VT_DEVICE_WEIGHT_BUDGET_BYTES", "64"); + CHECK(vllm::DeviceWeightBudgetBytes(4096) == 64); + CHECK(vllm::DeviceWeightBudgetBytes(0) == 64); + vllm_test::UnsetEnv("VT_DEVICE_WEIGHT_BUDGET_BYTES"); + } + SUBCASE("an explicit 0 disables the check, which is a documented escape") { + vllm_test::SetEnv("VT_DEVICE_WEIGHT_BUDGET_BYTES", "0"); + CHECK(vllm::DeviceWeightBudgetBytes(4096) == 0); + vllm_test::UnsetEnv("VT_DEVICE_WEIGHT_BUDGET_BYTES"); + } + SUBCASE("trailing garbage is IGNORED, not read as 0") { + // Reading "12x" as 0 would silently disable the refusal on a typo, which is + // the invisible-fallback shape this tree refuses. The probe must survive. + for (const char* bad : {"12x", "x", "-1", " 64", "64 ", "1e9", ""}) { + CAPTURE(bad); + vllm_test::SetEnv("VT_DEVICE_WEIGHT_BUDGET_BYTES", bad); + CHECK(vllm::DeviceWeightBudgetBytes(4096) == 4096); + } + vllm_test::UnsetEnv("VT_DEVICE_WEIGHT_BUDGET_BYTES"); + } +} diff --git a/tests/vllm/platforms/test_platform.cpp b/tests/vllm/platforms/test_platform.cpp index d3f7c3540..05cf0ddfb 100644 --- a/tests/vllm/platforms/test_platform.cpp +++ b/tests/vllm/platforms/test_platform.cpp @@ -384,3 +384,42 @@ TEST_CASE("residency_policy carries per-platform values the model consumes") { CHECK_FALSE(ShouldReleaseHostWeights(cpu, /*marlin=*/true, /*env=*/true)); CHECK_FALSE(ShouldInterleaveLoadStream(cpu, /*marlin=*/true)); } + +// The CUDA platform's POLICY ASSEMBLY, on every host. +// +// `CudaPlatform::residency_policy()` lives in `src/vllm/platforms/cuda.cpp`, which +// compiles only in a CUDA build, so while the four assignments lived there nothing +// on a CPU-only host could reach them: #1123 recorded "delete the +// `device_memory_total_bytes` assignment" as an OWED mutation for exactly that +// reason. The assembly is now a free function in the platform header, `cuda.cpp` +// calls it with its probe, and this case pins every field it sets (#1136). +// +// What is still NOT pinned here, and is not claimed to be: the `cudaMemGetInfo` +// call and the constructor threading in `cuda.cpp`. Those need a CUDA build. +TEST_CASE("CudaResidencyPolicy assembles the CUDA policy, budget included") { + using vllm::platforms::CudaResidencyPolicy; + + // GB10 as measured: cudaMemGetInfo total = 128452956160 (119.631 GiB). + const size_t kGb10Total = 128452956160U; + const ResidencyPolicy probed = CudaResidencyPolicy(kGb10Total); + // The three fields that predate #1123, unchanged: the CUDA path frees the host + // mirror after the Marlin build, pools device scratch, and leaves it uncapped. + CHECK(probed.release_host_weights_after_upload); + CHECK(probed.uses_device_memory_pool); + CHECK(probed.device_pool_cap_bytes == 0); + // The field #1123 added. This assertion is the one the owed mutation wanted: + // deleting the assignment leaves 0, which the fit refusal reads as UNKNOWN, so + // a checkpoint that cannot fit would load and die on the first forward again. + CHECK(probed.device_memory_total_bytes == kGb10Total); + // And it must be the ARGUMENT, not a constant: a second value moves it. + CHECK(CudaResidencyPolicy(4096).device_memory_total_bytes == 4096); + + // A failed probe is 0 = UNKNOWN, and 0 must survive as 0 rather than being + // substituted. The derived decisions are unaffected by the budget either way, + // which is what makes this field additive. + const ResidencyPolicy unknown = CudaResidencyPolicy(0); + CHECK(unknown.device_memory_total_bytes == 0); + CHECK(unknown.release_host_weights_after_upload); + CHECK(ShouldReleaseHostWeights(unknown, /*marlin=*/true, /*env=*/true)); + CHECK(ShouldInterleaveLoadStream(unknown, /*marlin=*/true)); +}