feat(ENG-RESIDENCY-CONFIG): the disk-residency tier is a config key, an absent field means unchanged at BOTH ends of the install, and a typo anywhere in the document is refused (#1110, #1109, #1122, #1133) - #1119
Conversation
…ce, and why the mirror stays untouched
vllm.cpp offloads weights at two tiers and only one was reachable from a config
surface. `--offload-config '{"uva":{"cpu_offload_gb":N,"cpu_offload_params":
["experts"]}}'` is wired from the server flag through `EngineParams` to
`LoadedEngine::FromModelDir`. The tier that actually makes `Qwen3.8-2.4T-A95B
UD-Q1_0` (370 GiB) serve on a 119 GB box — mmap residency, load-time prefault,
and the expert-stream lane with its slot count and slot size — was environment
only, and `Qwen35ExpertStreamRequested` (`qwen3_5.cpp:5146-5154`) read one
`getenv` and had no other input. That asymmetry is the more awkward for the first
tier's config being expert-aware already, so a user reasonably reaches for
`--offload-config` to control expert placement and finds nothing for the tier the
big-model case needs.
This spec takes option 2 of #1110: a namespaced `vllm_cpp` key inside the
existing flag. One user-facing flag for one user-facing concept, and
`include/vllm/config/offload.h` stays a byte-faithful transcription of
`vllm/config/offload.py` — upstream has no disk tier, so the extension is
vllm.cpp-original by construction rather than by preference.
Three findings shape it, and each is recorded with what it costs. An unknown
top-level key is already accepted by `parse_offload_config_json`, which is what
makes the namespaced key possible and also why the extension refuses an unknown
key of its own: a misspelled `vllm_cpp` would otherwise silently disable the tier
holding the model in memory, discovered as an out-of-memory kill rather than an
error. Precedence is environment variable over JSON config over built-in default,
because those variables exist so a benchmark arm is switchable without restarting
the server and an A/B in flight depends on it. And every knob latches in a
function-local static, so a config installed after the first read would be
silently ignored — the install therefore throws on a late non-empty install, and
the loader installs in `FromModelDir`'s first statement block, ahead of all
weight I/O.
`VT_MOE_EXPERT_STREAM_STATS_EVERY` stays environment-only by decision, argued in
the Port map: it changes a diagnostic cadence rather than what the process
reserves, so it is the instrument and not the configuration.
Issue index rows for #1110 and for #1109, whose `VT_GGUF_PREFAULT` default the
implementation corrects in flow rather than paper over while editing the same
table.
FOLLOWING_AGENTS_PROTOCOL
Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: ClaudeCode:claude-opus-5 [claude-code]
…ow, and three mutations found it was watching the wrong things (#1110, #1109) The five knobs that keep a 370 GiB checkpoint inside 119 GB are reachable from `--offload-config` under a `vllm_cpp` key: `mmap.enabled`, `mmap.prefault`, `expert_stream.enabled`, `expert_stream.slots` and `expert_stream.slot_bytes`. The chain is the server flag or the C ABI's `offload_config` string -> `EngineParams::weight_residency` -> the install in `LoadedEngine::FromModelDir`'s first statement block, ahead of the offloader and every path and weight operation, because each knob is read through a static that latches on first use. Closes #1110. `include/vllm/config/offload.h` is untouched. It transcribes `vllm/config/offload.py` line for line and upstream has no disk tier at the pin, so the extension is vllm.cpp-original by construction. Two parsers read one string, each seeing only its own half, which keeps one flag for one concept and needs no new ABI field. The extension refuses a key it does not know, because the mirrored parser ignores unknown keys and a silently dropped `{"vllm_cpp":{"mmapp":...}}` starts a server that does not borrow its weights and is met as an out-of-memory kill rather than an error. Precedence is env var, then config, then built-in default, and the install prints one line naming what it installed plus any variable shadowing it. Those variables exist so a benchmark arm is switchable without a restart, so `VT_X=0` has to beat a config `true`. Each knob keeps its own resolver, its own env name and its own historical polarity: `VT_MOE_EXPERT_STREAM` examines only the first character, so `VT_MOE_EXPERT_STREAM=false` is ON, and normalising that would have changed what a value means in a row about where values come from. `VT_MOE_EXPERT_STREAM_STATS_EVERY` stays environment-only, argued in the spec. Also fixes #1109: `docs/ENVIRONMENT.md:50` gave `VT_GGUF_PREFAULT`'s default as off while the code has always defaulted it on. This change writes that default into a resolver and a config key, so leaving the table contradicting it would put the contradiction inside one commit. Three mutations were findings rather than confirmations, and each is repaired here. The server-level suite stayed green with the install call site deleted, because the log line was built from `params` and not from the installed config, so the log and the install were independent statements; it reads the global back now. With the prefault site mutated to never consult its resolver, three GGUF suites stayed green (39/39, 6/6, 11/11), because byte-transparency is equally true of a prefault that never ran; a span counter is the observable. And `test_expert_stream_mixed_slot` set the slot count to its default 64, so the value could not be told from the site ignoring the variable; it sets 96 now and asserts the geometry the store was built with. Both coverage gaps predate this row. Sixteen mutations in total, each applied alone with the sha256 before and after, the build's exit status beside every result, a non-zero case count required, and the tree restored by byte copy and verified. Full gate: 502 of 502, exit 0, clean build. Row `ENG-RESIDENCY-CONFIG` is ACTIVE and not DONE. Nothing has driven the 370 GiB checkpoint through the JSON form on a box that can hold it: dgx.casa was unreachable at the SSH layer for this work, everything here is CPU-local, and the measurement is recorded under the spec's `## Owed`. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: ClaudeCode:claude-opus-5 [claude-code]
… env table resolved by key `docs/ENVIRONMENT.md` conflicted on the four `VT_MOE_EXPERT_STREAM*` rows. #1100 rewrote the `STATS_EVERY` row to document the new end-of-process final line, and this branch had appended a "also settable as `--offload-config`'s `vllm_cpp....`" clause to each of the four. Resolved per AGENTS.md `## Records`: origin/main's complete rows were taken and this branch's scoped clause re-applied to each, so the incoming prose survives verbatim and no unrelated key moved. The other files auto-merged, including `qwen3_5.cpp`, whose expert-stream constructor and `Qwen35ExpertStreamRequested` were verified by reading rather than trusted to the three-way result. The full gate is re-run on the merged tree; the pre-merge run does not carry. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: ClaudeCode:claude-opus-5 [claude-code]
… the per-expert-slice decode path (#1110) `ResolveExpertStreamRequested` marked the process-global latch by taking the `weight_residency` mutex, and it did so on every call including the cached ones. That function is reached once per expert slice: `KqExpertSlice` -> `Qwen35ExpertStream::Get` -> `Qwen35ExpertStreamRequested`. So the row put a process-wide lock in the decode loop of the lane it exists to make configurable, while its own spec, matrix row and pull request all state that it changes no kernel, no allocation and no performance axis. Whichever of those is wrong, they cannot both stand. The flag is now a `std::atomic<bool>` and every mark is a relaxed store, so the hot path costs an atomic write instead of a lock acquisition. Relaxed is sufficient: the flag only has to be observed as true by a LATER `SetWeightResidencyConfig`, and that call takes the mutex, so it synchronises with nothing weaker than it already needed. `WeightResidencyLatched` no longer locks either. Found by reading rather than by a mutation, which is why it is recorded in the spec's Evidence beside the three mutations that were also findings: no mutation of a correctness guarantee could have surfaced it, because the behaviour was already correct. M6 (remove the late-install throw) and the new M17 (never mark the latch) were both re-run against the atomic form and both stay RED, each with 11 cases and 2 failing, so the guarantee the flag carries is still gated. Focused suites unchanged and green: `test_weight_residency_config` 11 cases / 138 assertions, `test_weight_residency_reach` 5 / 39, `test_serve_residency_config` 5 / 55, `test_gguf_keep_quant` 39 / 6093, `test_expert_stream_mixed_slot` 1 / 181. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: ClaudeCode:claude-opus-5 [claude-code]
…ecode, no conflict `a6df72777` (#1113) landed the NemotronH paged forward. It touches no file this branch edits, so the three-way merge is clean rather than resolved, and the ratchets it shares with this branch were re-checked rather than assumed: `check-public-doc-tables` is green, so neither its `docs/STATUS.md` and `docs/BENCHMARKS.md` rows nor this branch's pushed the long-paragraph or oversized-cell counts over their pins. The gate re-runs on the merged tree; a clean merge is not a built one. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: ClaudeCode:claude-opus-5 [claude-code]
…ATURES parser count, vt::Conv1d, no conflict FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5-1m [claude-code]
…tream knob, and narrow the late-install throw (#1122) The fresh review of #1119 returned FAIL: 2 high, 3 medium, 8 low, and 14 of 69 sentences about the code wrong. This repairs all thirteen findings and audits the prose around them. H1. `RejectUnknownKeys` enumerated only INSIDE `vllm_cpp`, so a misspelled TOP-LEVEL key parsed to an empty config while the header claimed a typo "is an error". `{"vllm-cpp":{...}}` is the likeliest spelling of all, because every flag around it is hyphenated, and it started a server running this tier at its defaults -- prefault ON, streaming OFF, the two settings the 370 GiB case exists to change -- so the operator met the typo as an out-of-memory kill. The claim is made true rather than deleted, and the precondition the finding demanded was checked first. There is no `--offload-config` string anywhere in the vLLM tree at the pin, so the whole JSON document is vllm.cpp's own and no upstream-legal document can be broken by refusing; and vLLM builds its config dataclasses with `@config`, whose body sets `ConfigDict(extra="forbid")` under the comment "Extra fields are forbidden by default" (`vllm/config/utils.py:68-69` @ 555967922). Refusing is therefore the mirror-faithful polarity and the tolerance was the deviation. The enumeration lives in the EXTENSION parser, which reads the same string at both production entry points, so `parse_offload_config_json` stays a byte-faithful transcription of the mirror. The four legal top-level keys are `offload_backend`, `uva`, `prefetch` and `vllm_cpp`. M2. Nothing gated the headline knob reaching its decision: rewiring `ResolveExpertStreamRequested` to read `.mmap` instead of `.expert_stream` -- the adjacent field of the same type -- left all four suites green. That answer is cached in a function-local static, so a process can observe what it resolved exactly once, and every other suite has to spend its one observation elsewhere. The observation now has a binary of its own. M1 is a decision. The late-install throw fired on a legitimate two-model process -- load A with no residency config, load B carrying `vllm_cpp`, and B could not load -- and its stated reason was false for two of the five knobs: `mmap` resolves per load and this row had already deleted prefault's static. It is NARROWED rather than kept-and-documented, because a hard failure on a legal second load is worse than the thing it prevents, and because the reason can then be true. Reading a knob is not taking a decision; what cannot be retaken is the streaming answer and the slot store's geometry. So the latch is per-decision (`ResidencyLatch::kExpertStream`, `kExpertStreamGeometry`), the shared resolvers mark nothing, and the refusal fires only on a field a taken decision has frozen. An equal re-install, an empty install, and a document touching only `mmap`/`prefault` are all accepted. H2. `include/vllm.h` still told a C client that `offload_config` was "ACCEPTED BUT NOT YET ACTED ON ... no weight moves yet". The same string now carries `vllm_cpp`, which does move weights, and the row added `VLLM_ERR_INVALID_ARGUMENT` refusals that fire through `vllm_engine_load`. The inert sentence is scoped to the three mirrored keys and the live half is documented beside it, refusals included. M3. The install line prints the installed DOCUMENT, not resolved values, and it cannot print resolved values without moving the streaming decision ahead of the weight load. Reading the global back is still the right mechanism -- it is what makes the line evidence that the install ran -- so the mechanism is unchanged and the three sentences that claimed otherwise are corrected, along with the test name and the two test comments that stated opposite positions. L1 the message path is built from the document rather than a hardcoded `vllm_cpp.` prefix, and the tests now assert the message and not only the throw. L2 the relaxed atomic keeps its ordering and loses its untrue reason: a mutex acquire does not synchronise-with a relaxed store made outside it. What carries it is that the flags publish one monotonic bit and no config state. The window ordering cannot close is recorded rather than denied. L3 `ActiveWeightResidencyConfig` returns by value, copied under the lock, instead of a reference read after it was released. L4 and L5 are recorded under `## Owed` and in `docs/USAGE.md`: `--offload-config` reaches neither `vllm-cli` nor the server's pooling and transcription paths, which build their engine parameters without the document at all -- the mirrored half included, and since before this key existed. L6 two `docs/ENVIRONMENT.md` rows enumerated a resolution that is false when a config is present. L7 `DescribeEnvOverrides` reported presence, so `VT_MOE_EXPERT_STREAM_SLOTS=banana` was announced as an override the resolver ignores; it now asks the same predicate the resolver uses. L8 the `EnvOn` anchor is `61-66`. The #1110 and #1109 index rows are left where they are. Relocating a row to the end is a delete plus an add, which is what the append-only rule forbids and what makes two branches' relocations merge into a duplicate under `merge=union`; the new row is appended properly. Eight mutations, each applied alone, with `git diff --stat` and the file's sha256 before and after, the build's exit status and an ENOSPC count beside every result, the LAST `test cases:` match read, and the tree restored by byte copy with the sha256 compared. All red. N4, the coarse latch restored inside `ResolveResidencyBool`, first left the reach suite GREEN while the config suite went red, because a per-field comparison is insensitive to where the flag is marked unless the document changes the frozen field; the reach case was widened to the shape the reviewer actually measured and goes red too. That first result is recorded in the spec, because a green mutation is a statement about the test. L3 and L2 have no mutation and say so: a single-threaded suite cannot distinguish a copy from a reference, and both were sound code with a wrong sentence. Gate on the merged tree: build rc 0 with zero warnings and zero ENOSPC, `ctest -j 6` 505/505 rc 0 (504 before, plus the new binary), focused suites `test_weight_residency_config` 14 cases / 188 assertions, `test_expert_stream_latch` 1 / 9, `test_weight_residency_reach` 6 / 54, `test_serve_residency_config` 6 / 64, all rc 0. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5-1m [claude-code]
…deo recipe, MUSIC3 DiT staging, ROCm GetBlas TLS Bring the branch onto d1e5e9b so the trailer and commit-style gates, which scope themselves to origin/main..HEAD, actually examine this tree instead of skipping. No file in the three incoming commits overlaps this row. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5-1m [claude-code]
…e install, and the nested typo the ABI already promised to refuse (#1133) The schema says an absent field means unchanged. The code meant two other things by it, one at each end of the install, and each produced its own user-visible failure on a legal two-model process. Fixing them as one thing is the point of this change. `FrozenFields` treated absent as A CHANGE. It compared `in.expert_stream` against the stored optional, and `nullopt != engaged` is true, so once a process had latched the streaming decision, `{"vllm_cpp":{"mmap":{"enabled":true}}}` on a second engine threw `std::logic_error` out of `LoadedEngine::FromModelDir` and came back as `VLLM_ERR_MODEL_LOAD`. #1122's M1 was narrowed, not removed: the same hard failure on a legal load, now reached by a document that merely omits the latched field. The message made it worse by asserting that "accepting this would record a configuration the engine is not running" — for a document asking for exactly what the process had resolved, the engine was running it. The install treated absent as A CLEAR. `g.config = config` replaced wholesale, so a partial second document turned `expert_stream=on` with 8000 slots into OFF with 64, with no diagnostic, and the slot store reads those values lazily — on the first slice taken, which can be long after the second engine loaded. The wholesale replace predates the per-field narrowing; the narrowing widened its reach, because before it any differing document after a decision threw, so the drop could not happen once anything had been decided. Both are now what the schema says. A field is scored frozen only when the document SETS it. The comparison is against THE DECISION ACTUALLY TAKEN — the cached streaming answer, the geometry the store was built with — resolved through the same function the production resolver calls, so a document that agrees with the running engine installs, and so does one the environment overrides anyway, while one that would make a resolver return something else still throws. The install merges field by field, which makes the empty document a no-op by construction instead of by a special case. The refusal message now quotes the decision rather than the stored document, because after the merge the stored document is not what is in force either. The streaming flag carries its own answer to make that comparison possible: one tri-state atomic (nothing decided / decided off / decided on) rather than a bool plus a value, so a reader takes one relaxed load and there is no ordering question between a fact and a value in two places. The geometry's numbers move from a file-static into `Global` under `mu`, beside the flag's writer. Both stores are idempotent and both are monotonic, which is what carries `relaxed` here — not a synchronises-with edge, which a relaxed store outside `mu` does not give. Third defect, the same shape one level over. `include/vllm.h` claimed an unknown key ANYWHERE in the document is refused, and the enumeration closed the top level and the inside of `vllm_cpp` and stopped. Measured: `{"uva":{"cpu_offload_gbb":1}}`, `{"uva":{"cpu_offload_GB":10}}` and `{"prefetch":{"offload_groupsize":8}}` were all ACCEPTED, each giving a 0 GiB budget or a group size of 0 under a document the operator believes configures offloading. The claim is made TRUE rather than scoped down, because refusing is mirror-faithful here too: `UVAOffloadConfig` (`vllm/config/offload.py:15-16` @ 555967922) and `PrefetchOffloadConfig` (`:47-48`) each carry `@config`, whose body sets `ConfigDict(extra="forbid")` (`vllm/config/utils.py:68-69`), so upstream refuses a nested typo and the tolerance was the deviation. The enumeration stays in the EXTENSION parser and lists NAMES only, so `parse_offload_config_json` is untouched and keeps sole ownership of those fields' types, defaults and bounds. WHY EIGHT MUTATIONS AND A FRESH REVIEW MISSED BOTH BEHAVIOUR DEFECTS, which is the part worth a case shape rather than another sentence. Every latch case installed a COPY of the first document — `mmap_too = cfg`, `mmap_only = sizes` — or the empty one, so the second install always restated the frozen field at the value it already had, and that is the one shape in which neither defect can arise. Two DIFFERENT partial documents in one process is the shape that distinguishes them, and it is now a case at the unit level and again through `FromModelDir` and `vllm_engine_load`. The reach case has to read the global rather than the return code: a refused install and a missing checkpoint both leave `vllm_engine_load` as `VLLM_ERR_MODEL_LOAD`, so `mmap` asked for FALSE against an installed TRUE is the discriminator. Neither new case assumes the streaming answer's value: it is a per-process static, so the cases read it and assert relative to it, or a `-tc` run of one case would pass or fail on test order rather than on the code. Prose repairs, all from the same review. The spec's surviving "the install line prints the RESOLVED values"; `qwen3_5.cpp` stating the pre-narrowing contract; the function-local static attributed to `Qwen35ExpertStreamRequested` when it lives in `ResolveExpertStreamRequested` and that function is a pure delegation; "refuses a late install" for "refuses a late change"; the reach case's "exactly as every GGUF load in the tree does", when load A stops on the missing checkpoint and the case calls `FromEnv()` itself two lines later; "every knob resolves lazily", when the constraint binds `expert_stream` alone and the other four could be resolved at install; "one line" for one line plus a conditional second; and the claim that an absent `vllm_cpp` key is "byte-identical to the engine before this row existed", which a `{"typo":1}` abort falsifies for the set of documents accepted even though it holds for the engine a legal document produces. Eight stale `path:line` anchors in this row's own spec, found by checking anchors last against the final tree. Seven were already stale at `c7fa7084b`, verified by reading the same line numbers out of `git show HEAD~1:<path>`: the Port map's four `VT_MOE_EXPERT_STREAM*` rows pointed at a comment and at three lines of the statistics printer rather than at the constructor beside them, the "refuses by name" citation pointed at magic-static commentary, `GgufLoadPolicy::FromEnv` was cited 85 lines past itself, the `EngineParams::offload_config` field was off by one, and `include/vllm.h:436` fell mid-sentence. Three stale anchors belonging to OTHER rows are named in the spec and left alone rather than edited, because correcting them would widen this diff into three unrelated subsystems. The previous commit's body reported `ctest -j 6` as 505/505 where the tree measured 506/506. That body is history and cannot be corrected in place; the count for this tree is below, and the pull request body — which is what the squash lands — carries it. `## Owed`'s unreached-entry-point gap named #1122 as its owner, the very issue this pull request closes, so on landing it would have had no open issue. Filed as #1135 and pointed there, in the spec and in `docs/USAGE.md`. Gate on the merged tree (`origin/main` at `d1e5e9bc0`, merged in so the trailer and commit-style gates examine this branch instead of skipping). CLEAN rebuild from an empty build directory, 974 translation units: exit 0, zero compiler warnings, zero ENOSPC lines. `ctest --test-dir build -j 3`: 507/507, exit 0, two pre-existing environment skips (`test_modelopt_mixed_precision_checkpoint`, `test_voxtral_e2e`). Earlier runs at `-j 4` and `-j 6` on a box at load average 47-62 starved `test_engine_core_proc` and `test_serve_low_tools`; each passes on a serial re-run, which is what `.agents/verification.md` requires before calling a starved test a regression. `scripts/agent-preflight.sh --staged` exit 0, all gates green, no gate skipped. Focused, counts from the last `test cases:` match and every rc captured directly: `test_weight_residency_config` 17 cases / 245 assertions, `test_weight_residency_reach` 7 / 76, `test_expert_stream_latch` 1 / 9, `test_serve_residency_config` 6 / 64 (1 skipped, the re-exec'd child). FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5-1m [claude-code]
…e install, and the nested typo the ABI already promised to refuse (#1133) The schema says an absent field means unchanged. The code meant two other things by it, one at each end of the install, and each produced its own user-visible failure on a legal two-model process. Fixing them as one thing is the point of this change. `FrozenFields` treated absent as A CHANGE. It compared `in.expert_stream` against the stored optional, and `nullopt != engaged` is true, so once a process had latched the streaming decision, `{"vllm_cpp":{"mmap":{"enabled":true}}}` on a second engine threw `std::logic_error` out of `LoadedEngine::FromModelDir` and came back as `VLLM_ERR_MODEL_LOAD`. #1122's M1 was narrowed, not removed: the same hard failure on a legal load, now reached by a document that merely omits the latched field. The message made it worse by asserting that "accepting this would record a configuration the engine is not running" — for a document asking for exactly what the process had resolved, the engine was running it. The install treated absent as A CLEAR. `g.config = config` replaced wholesale, so a partial second document turned `expert_stream=on` with 8000 slots into OFF with 64, with no diagnostic, and the slot store reads those values lazily — on the first slice taken, which can be long after the second engine loaded. The wholesale replace predates the per-field narrowing; the narrowing widened its reach, because before it any differing document after a decision threw, so the drop could not happen once anything had been decided. Both are now what the schema says. A field is scored frozen only when the document SETS it. The comparison is against THE DECISION ACTUALLY TAKEN — the cached streaming answer, the geometry the store was built with — resolved through the same function the production resolver calls, so a document that agrees with the running engine installs, and so does one the environment overrides anyway, while one that would make a resolver return something else still throws. The install merges field by field, which makes the empty document a no-op by construction instead of by a special case. The refusal message now quotes the decision rather than the stored document, because after the merge the stored document is not what is in force either. The streaming flag carries its own answer to make that comparison possible: one tri-state atomic (nothing decided / decided off / decided on) rather than a bool plus a value, so a reader takes one relaxed load and there is no ordering question between a fact and a value in two places. The geometry's numbers move from a file-static into `Global` under `mu`, beside the flag's writer. Both stores are idempotent and both are monotonic, which is what carries `relaxed` here — not a synchronises-with edge, which a relaxed store outside `mu` does not give. Third defect, the same shape one level over. `include/vllm.h` claimed an unknown key ANYWHERE in the document is refused, and the enumeration closed the top level and the inside of `vllm_cpp` and stopped. Measured: `{"uva":{"cpu_offload_gbb":1}}`, `{"uva":{"cpu_offload_GB":10}}` and `{"prefetch":{"offload_groupsize":8}}` were all ACCEPTED, each giving a 0 GiB budget or a group size of 0 under a document the operator believes configures offloading. The claim is made TRUE rather than scoped down, because refusing is mirror-faithful here too: `UVAOffloadConfig` (`vllm/config/offload.py:15-16` @ 555967922) and `PrefetchOffloadConfig` (`:47-48`) each carry `@config`, whose body sets `ConfigDict(extra="forbid")` (`vllm/config/utils.py:68-69`), so upstream refuses a nested typo and the tolerance was the deviation. The enumeration stays in the EXTENSION parser and lists NAMES only, so `parse_offload_config_json` is untouched and keeps sole ownership of those fields' types, defaults and bounds. WHY EIGHT MUTATIONS AND A FRESH REVIEW MISSED BOTH BEHAVIOUR DEFECTS, which is the part worth a case shape rather than another sentence. Every latch case installed a COPY of the first document — `mmap_too = cfg`, `mmap_only = sizes` — or the empty one, so the second install always restated the frozen field at the value it already had, and that is the one shape in which neither defect can arise. Two DIFFERENT partial documents in one process is the shape that distinguishes them, and it is now a case at the unit level and again through `FromModelDir` and `vllm_engine_load`. The reach case has to read the global rather than the return code: a refused install and a missing checkpoint both leave `vllm_engine_load` as `VLLM_ERR_MODEL_LOAD`, so `mmap` asked for FALSE against an installed TRUE is the discriminator. Neither new case assumes the streaming answer's value: it is a per-process static, so the cases read it and assert relative to it, or a `-tc` run of one case would pass or fail on test order rather than on the code. Prose repairs, all from the same review. The spec's surviving "the install line prints the RESOLVED values"; `qwen3_5.cpp` stating the pre-narrowing contract; the function-local static attributed to `Qwen35ExpertStreamRequested` when it lives in `ResolveExpertStreamRequested` and that function is a pure delegation; "refuses a late install" for "refuses a late change"; the reach case's "exactly as every GGUF load in the tree does", when load A stops on the missing checkpoint and the case calls `FromEnv()` itself two lines later; "every knob resolves lazily", when the constraint binds `expert_stream` alone and the other four could be resolved at install; "one line" for one line plus a conditional second; and the claim that an absent `vllm_cpp` key is "byte-identical to the engine before this row existed", which a `{"typo":1}` abort falsifies for the set of documents accepted even though it holds for the engine a legal document produces. TEN stale or imprecise `path:line` anchors in this row's own spec, found by checking anchors last against the final tree. All ten were already wrong at `c7fa7084b`: for `qwen3_5.cpp` and `include/vllm.h` that was verified by reading the same line numbers out of `git show HEAD~1:<path>`, and for the other three files it follows from this change not touching them. The Port map's four `VT_MOE_EXPERT_STREAM*` rows pointed at a comment and at three lines of the statistics printer rather than at the constructor beside them, the "refuses by name" citation pointed at magic-static commentary, `GgufLoadPolicy::FromEnv` was cited 85 lines past itself, the `EngineParams::offload_config` field was off by one, `include/vllm.h:436` fell mid-sentence, and the `FromModelDir` range started one line late and ended 30 lines early. The count was first written as eight and corrected by auditing it against the actual corrections. Three stale anchors belonging to OTHER rows are named in the spec and left alone rather than edited, because correcting them would widen this diff into three unrelated subsystems. The previous commit's body reported `ctest -j 6` as 505/505 where the tree measured 506/506. That body is history and cannot be corrected in place; the count for this tree is below, and the pull request body — which is what the squash lands — carries it. `## Owed`'s unreached-entry-point gap named #1122 as its owner, the very issue this pull request closes, so on landing it would have had no open issue. Filed as #1135 and pointed there, in the spec and in `docs/USAGE.md`. Gate on the merged tree (`origin/main` at `d1e5e9bc0`, merged in so the trailer and commit-style gates examine this branch instead of skipping). CLEAN rebuild from an empty build directory, 974 translation units: exit 0, zero compiler warnings, zero ENOSPC lines. The disk ran between 89% and 95% full throughout, and a full disk produces a link error that reads exactly like broken code, so an ENOSPC count is printed beside every build and mutation result. `ctest --test-dir build -j 3`: 507/507, exit 0, two pre-existing environment skips (`test_modelopt_mixed_precision_checkpoint`, `test_voxtral_e2e`). Earlier runs at `-j 4` and `-j 6` on a box at load average 47-62 starved `test_engine_core_proc` and `test_serve_low_tools`; each passes on a serial re-run, which is what `.agents/verification.md` requires before calling a starved test a regression, and both are green in the `-j 3` run above. Not claimed: that they were red before this change, because no `ctest` was run on the pre-change tree. `scripts/agent-preflight.sh --staged` exit 0, all gates green, no gate skipped. Focused, counts from the last `test cases:` match and every rc captured directly: `test_weight_residency_config` 17 cases / 250 assertions, `test_weight_residency_reach` 7 / 76, `test_expert_stream_latch` 1 / 9, `test_serve_residency_config` 6 / 64 (1 skipped, the re-exec'd child). FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5-1m [claude-code]
… ROCm head_dim=128 decode arm Bring the branch onto a7583ac so the trailer and commit-style gates, which scope themselves to origin/main..HEAD, keep examining this tree. Neither incoming commit touches a file this row changes. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5-1m [claude-code]
…n merge doubled `merge=union` on `.agents/issue-index.md` concatenates both sides of a changed region without deduplicating, so merging `a7583ac75` — which appends #1129, #1130 and #1134 after the same rows this branch had already taken from `d1e5e9bc0` — emitted #837, #1117 and #1118 twice, byte-identical each time. `check-agent-record.py` refuses that by name, and it is right to: under union a duplicate is exactly what two branches appending the same issue look like, so it cannot distinguish this from a real double-append. Reconciled the way AGENTS.md prescribes for a keyed record rather than by hand- deleting three lines: take the complete target-branch version of the file, then re-apply this branch's scoped edit. The result is asserted to be `a7583ac75`'s file as a BYTE-EXACT PREFIX, with exactly this branch's five rows appended after it (#1110, #1109, #1122, #1133, #1135) and every one of main's 315 rows textually unchanged. `check-agent-record.py` goes from three ERRORs to OK. This removes two duplicate rows and adds none, so the append-only rule is satisfied in the sense that matters: no issue loses its row, and no row's text is edited. What is deleted is a merge artifact that was never authored. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5-1m [claude-code]
…ushing over it `19e6d5b1a` was pushed, then amended twice — once to correct claims this session's own audit found wrong, once to add a case and a mutation — which made the local tip a sibling of the remote one rather than a descendant, so the push was rejected as a non-fast-forward. It is rejected correctly: a force-push is not available here for any branch, and the reviewed head every review anchors to is reachable only through history that is never rewritten. So the pushed commit is merged back in as a parent instead. The resulting TREE is byte-identical to `55cd0fe43` — `git diff 55cd0fe` is empty — because `19e6d5b1a`'s content is the same work minus the later corrections. The one conflict, in the row's spec, is where the corrected paragraph meets the uncorrected one, and it resolves to the corrected side. `.agents/issue-index.md` was re-checked after the union driver ran again: no row is duplicated and `check-agent-record.py` is OK. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5-1m [claude-code]
|
Third repair round, at Disposition
The H3 decisionEnumerate, in the extension parser, names only. Three reasons, in order of weight. It is mirror-faithful rather than local strictness, and that was checked before it was written: It keeps It adds no refusal a user could previously rely on. The mirrored parser runs first at both entry points and already threw the same sentence for a non-object The alternative the round allowed — keep the enumeration shallow and scope the four over-claiming sentences down — was rejected because the sentences describe the behaviour a user needs, and The case shape, and its mutationsThe hole was that every latch case installed a COPY of the first document (
Neither new case assumes the streaming answer's value. It is a per-process static, so each case reads it and asserts relative to it; hardcoding
R12 is the one worth reading. The first message assertion was R15 is GREEN and stays recorded as GREEN. Clearing the built geometry on reset is coherence with Claims audit58 claims examined, 5 wrong, all 5 corrected before the push. The three rounds before this were 14/69, 41/46 and 21/55.
Every guarantee stated in the code, the spec and the two bodies now names the case or mutation that makes it true, or says which it does not have. The three that do not have one, said so explicitly: L3's stale count (history, not correctable in place), the reset's geometry clear (R15 green, coherence not behaviour), and the eight prose repairs (a reader is their gate). Deferred
|
… executable, which pr-size has been refusing since #1110 landed `pr-size` is a required check and has been RED on this pull request since `26a97b75e`, including on the reviewed head `c7fa7084b`. Neither earlier repair round mentioned it. The reason is exact and correct: `scripts/check-pr-size.py` classes any change to `scripts/check-*.py` as a `governance_checker` and demands a non-empty edit in the paired test file, and this row changed two of them with no test edit at all. What it changed in each was a MARK, not a rule: `ENGINE_ROWS` 157 -> 158 in `check-agent-record.py`, and one entry in `RUNNABLE_BASELINE` in `check-gate-commands.py`. That is why it slipped past three rounds of review — there is no branch to invert. It is also exactly why the contract is right to fire: the constant was the only artifact, so "158 is the real row count" and "this row's Gates section names commands that can fail" were plausible rather than checkable, which is the state the evidence contract exists to refuse. Three cases, in the idiom the neighbouring credits already use for #117, #606, #633, #632, #81 and SERVE-RECIPE-ARGS. `test_residency_config_row_is_inside_the_engine_ratchet` names the row in the engine matrix and requires the count to agree, so the 158 and the row are one semantic change. It also fixes where the row BELONGS: it sits between two offload rows that could each plausibly have absorbed it, so "genuinely new" fails mechanically if it is ever folded into a neighbour without the count following. `test_residency_config_is_credited_for_real_commands` pins three things rather than one: the row is in the exact pin, its Gates section really does yield a command that can fail, and the spec says why no GPU, oracle or throughput leg is implicated. A CPU-only credit whose record is silent about the missing arms reads as coverage. `test_dropping_residency_config_from_the_pin_breaks_it` is the mutation in the direction the re-pin actually moved, which is what separates a row pinned because it entered the population from a row pinned to quiet a gate. Mutation-proven, and one of the two mutations took three attempts to become valid, which is the part worth recording. Dropping the entry from `RUNNABLE_BASELINE` turns both new gate-commands cases RED with all 39 cases running. Removing the row from the engine matrix alone aborted `AgentRecordMutationTests.setUpClass` on a dirty baseline, so only 51 of 78 cases ran and the new case NEVER EXECUTED — a mutation that reads as a pass. Lowering the pin with it did not help for the same reason. Only removing the row, the pin and the row's claim file together produces the tree state "this row was never added": then all 78 cases run and the new case is one of exactly two failures. Case counts moved with the additions, which is the other half of proving the cases run at all: gate-commands 37 -> 39 methods and 39 cases executed, agent-record 75 -> 76 methods and 78 cases executed, both OK after restore. The first draft of the credit case asserted the bare string `scripts/agent-preflight.sh` against a list `runnable_commands` fills with WHOLE commands, so it went red on membership rather than on the spec; the exact strings are asserted now and the reason is in the code. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5-1m [claude-code]
…t run on (#1136) A fresh review of #1132 found the code right and five sentences about it wrong. One of those sentences guarded a real defect, and it is the only behaviour change here. ## The defect: `ResolveModelDeviceType` could remove a working load `model_loader.h` claimed that `ResolveModelDeviceType` "picks what `SelectQueueForModel` will pick, so the two cannot drift". On the AUTO arm they could. `SelectQueueForModel` wraps `GetBackend(dev).CreateQueue()` in a try/catch and falls back to CPU, and its own comment says why: "a platform can be registered while CreateQueue still fails, and CPU must remain reachable". The resolver asked `CurrentPlatform()` and stopped there, so on such a box it answered `kCUDA` while the load ran on the CPU queue — and the #1123 fit refusal, which reads the resolver, then refused a checkpoint by naming device `'cuda'` for a model that previously loaded and served on CPU. Refusing a load that works is the one failure mode worse than the late `cudaMalloc` this row set out to remove. Whether `CreateQueue()` fails is knowable only by calling it, so both callers now call one `ResolveAutoDevice`, which creates the queue and hands it to whichever caller wants one. `ResolveModelDeviceType` destroys the queue it does not use: `vt::Queue` is a NON-OWNING handle (a raw `cudaStream_t`) with no destructor, so dropping the value would leak the stream. The teardown goes through the FREE `vt::DestroyQueue`, not `Backend::DestroyQueue`. That is what this file's only other queue teardown already does for the real load queue, it is what `vt/backend.h` asks of new code so device index and queue cleanup are never ambient, and on CUDA it additionally releases workspaces keyed on the queue. The CREATE side deliberately stays `GetBackend(...).CreateQueue()`: that is the call this arm has always made, and moving it onto the drop-in resource ABI would be a behaviour change to the production queue-selection path, which a repair round must not make. The cost is bounded by where the resolver is called: the GGUF fit check is its only caller outside `SelectQueueForModel`, so a safetensors load pays nothing, an explicitly named device pays nothing, and an auto-arm GGUF load pays one create/destroy pair. Stated in the header rather than hidden, and smaller than removing a working load. Two cases pin it, at the same budget and the same platform, differing only in whether the queue can be created: it refuses when it can, and refuses nothing when it cannot. The first is the POSITIVE CONTROL — without it, "no refusal" in the second could mean the auto arm never selected the fake platform at all. Both also assert the counters, so a resolver that skipped the attempt or leaked the queue is red rather than merely unnoticed. ## The bound CAN over-refuse, and now says so and is tested for it The spec and the previous commit body asserted the footprint was a lower bound "so the refusal can never over-refuse". False. A tensor counted and never staged is a positive over-count, and one is present on every default load: the MTP / `nextn` head is attached only under `speculative_config->method == "mtp"`, so 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 the bound counts)` refuses a weight set that fits. "The under-count dominates" is not an argument that the refusal is safe: the two errors are on different quantities and never cancel. `gguf_device_fit.h` said this correctly from the day it landed. The spec, the commit body and `docs/USAGE.md` did not, and they now use the header's wording. The direction is also EXECUTABLE rather than described: a fixture carrying a `blk.N.nextn.*` tensor asserts that the footprint counts it and asserts both ends of the resulting over-refusal window. Every other case in that file runs on a fixture whose tensors are all staged, so the bound there happened to EQUAL the true staged size and the boundary cases could not tell an exact quantity from an over-counted one. Not closed, 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 the 26-minute-load-then-OOM this row removed, on a device nobody on this fleet has to measure the change against. So it is stated, pinned, documented for operators with `VT_DEVICE_WEIGHT_BUDGET_BYTES` as the way out, and tracked as #1136. ## The false comment, in BOTH places it lived `include/vt/backend.h` still said `DeviceMemoryInfo` is overridden by "ROCm/CUDA" when only ROCm implements it, while the pull request body, the spec's `## Owed` and the issue-index row for #1126 all said the comment had been corrected. It is corrected now — and this round's audit found a SECOND copy of the same claim at the seam's only call site, `gemma4_moe.cpp`. Correcting the header alone would have left the falsehood in the tree, which is exactly the class this row keeps repeating. The CAPABILITY is still #1126: adding the override wakes `Gemma4MoE`'s device-expert LRU, which needs its own measurement. ## The unpinnable assignment, pinned without a CUDA build #1123 recorded "delete `p.device_memory_total_bytes = ...` from `CudaPlatform::residency_policy()`" as an owed mutation, because `src/vllm/platforms/cuda.cpp` compiles only in a CUDA build. The four-line policy assembly now lives in `CudaResidencyPolicy` in `vllm/platforms/interface.h`, which compiles everywhere, and `test_platform.cpp` pins every field it sets. `cuda.cpp` is left holding only the `cudaMemGetInfo` probe it alone can make; that call and the constructor threading are still not mutation-proven and are still said to be. The one-line edit to `cuda.cpp` was syntax-checked with `-fsyntax-only` against a stub `cuda_runtime.h`, and the instrument was verified by a positive control (a deliberate typo in the same expression returns rc 1 naming it). ## Records and anchors `docs/USAGE.md` was wrong in both halves of one bullet: a discrete NVIDIA GPU IS `CudaPlatform` and does get the refusal, and ROCm, Vulkan and Metal answer `needs_weight_staging() == false`, so they have no staging allocation to fail and the refusal is inapplicable rather than owed. The spec carried the same error. Fourteen stale or wrong anchors and names corrected, each verified against the FINAL tree rather than re-quoted (the two this change authored itself are counted in the section below, not here): `cuda_backend.cu:75-81` is `:77-81`; `BuildMoeMarlinResident` is `:6010-6215`, not `:6010-6062`, with its allocations at `:6049-6064` and two temporaries at `:6094-6095`; `gemma4_moe.cpp:439-447`/`:494-506` are `:449-455` and `:506`; `CheckGgufDeviceFit` does not exist and never did — the function is `CheckDeviceWeightFit`; the call-chain table mixed call sites with definition lines and now declares one convention and keeps it (`ExpertMlpKq:5651,5652`, `KqResidentSlice:5114`); and the header's upstream anchor `model_runner.py:504,647` pointed at a draft-model `set_attn` call and a `torch.zeros` — the startup memory profile is `gpu_worker.py:451-495` around `gpu/model_runner.py:682`, and it takes the weight bytes as an INPUT recorded after the load (`gpu/model_runner.py:315`), which is a stronger statement of why upstream has no counterpart than the one that was there. That pair was COPIED from `ENG-EXPERT-STREAM`'s neighbour: `.agents/engine-matrix.md`'s `KV-WARMUP-PROFILE` row carries the same two anchors plus a third, all three stale at the current pin. The rest, less interesting individually and the same class: `ReqFloat` is `qwen3_5_gguf_weights.cpp:843-847`, not `:843-845`; `CurrentPlatform()` is `platform.cpp:91-98`, not `:38-40`, and its walk is seven device types, not the five the comment listed (it omitted `kROCM` and `kTENSTORRENT`); `MoeBlock:6555` is the fp4 BRANCH, while the predicate it is guarded by is `const bool fp4 = !w.expert_gate_fp4.empty()` at `:6548`; and the `DeviceMemoryInfo` comment block is `backend.h:78-93` after this change, not `:79-83`. That is filed as #1139 and NOT fixed here: the fix is one cell in `engine-matrix.md`, which PR #1119 is concurrently bumping alongside the hardcoded count in `check-agent-record.py`, and two edits to that pair is the record-lock AGENTS.md names. The entry-point claim was also loose. Two production entry points reach `LoadedEngine::FromModelDir`: the server binary, on both its embedding lane (`server_main.cpp:959`) and its generative lane (`:1123`), and the C ABI (`vllm_c.cpp:766`). `examples/bench/bench_core.h:586` also calls it and is an example, which AGENTS.md says is not a production entry point. One contradiction was self-inflicted and caught before it shipped: `gguf_device_fit.h` opened by saying the footprint is "built to be WRONG LOW rather than wrong high" and closed, twelve lines later, by saying it can over-refuse. Both halves were about different things, and the fix is to name the scope: the sum is a lower bound on staging EVERY TENSOR IN THIS FILE, which is not a lower bound on what one load stages, because a load may stage a subset. The per-tensor term is exact-or-low; the SET is exact-or-high. That is where the over-count lives, and it is why the field keeps the name `lower_bound_bytes`. The "Release (what CI uses)" claim was also false: `build-test-cpu` passes no `CMAKE_BUILD_TYPE`, so CI's test lane has asserts live. This change was built and gated in that configuration. ## The anchors this change itself moved Inserting ~45 lines near line 100 of `model_loader.cpp` shifts every absolute line citation below it. Measured between `e7d0a1f7c` and this head by comparing the TEXT at each cited line: **203 moved line references over 109 citing sites in 45 files**, 10 unmoved. Two of them were authored by this change — a `model_loader.cpp:135-141` in the reachability gate and a `:1411-1412` in the arithmetic gate — and both are corrected against the final tree and added to this round's anchor verifier, which is the trap the round before hit with `platforms/cuda.cpp:67`. The other 107 are NOT swept, and not for effort: several were already stale at `e7d0a1f7c` (`model-matrix.md:197` cites `:184-223` as the "live loader" while line 184 there is `static const bool once = [] {`), so rewriting them from the current tree would launder pre-existing debt into a clean-looking record. Filed as #1143 with the measurement and three candidate fixes, and listed under `## Owed`. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [Claude Code]
|
Addendum to the round summary above, at
Three cases now make both marks executable, in the idiom the neighbouring credits for #117, #606, #633, #632 and #81 already use: the row is named in the engine matrix with the count required to agree; the One of the two mutations took three attempts to become VALID, and that is the part worth recording. Dropping the Case counts moved with the additions, which is the other half of proving the cases run: gate-commands 37 → 39 methods, 39 executed; agent-record 75 → 76 methods, 78 executed; both OK after restore. Also in this addendum: the first draft of the credit case asserted the bare string Gate re-run on |
…F16 default record Bring the branch onto ac50579 so the trailer and commit-style gates, which scope themselves to origin/main..HEAD, keep examining this tree rather than skipping. Records and documents only on the incoming side: no source or test file changes, so nothing this row compiled is invalidated. One conflict, in docs/ENVIRONMENT.md, and it is a KEYED record rather than an append-only log, so it is resolved per key rather than by taking a side. Both branches rewrote the same three adjacent table rows. `VT_GGUF_KEEP_F16` takes MAIN's row, because rewriting it is the whole point of ac50579. `VT_GGUF_MMAP` and `VT_GGUF_PREFAULT` take THIS branch's, because this row closes #1109 by correcting `VT_GGUF_PREFAULT`'s documented default from off to ON and main's copy still carries the wrong one. Each of the three is asserted byte-identical to the side it came from, and `check-env-doc` reports all 342 production variables documented or classified. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5-1m [claude-code]
|
Second addendum, at
The merge changed no file under
Merge order against #1132 was re-checked after the two merges: still exactly one conflict, still The build tree was removed. |
…, by name (#1123) (#1132) A 370 GiB checkpoint loaded for 26 minutes on `--device cuda`, reported ready, and then died on the first request. It now refuses at load, by name, and says what is missing. `Qwen3.8-2.4T-A95B UD-Q1_0` serves correctly on `--device cpu` on a DGX Spark: TTFT 667.0 s, steady decode 44.2 s/token, coherent output. On `--device cuda`, same box and same binary, it reached a serving state after 26 minutes and then died inside the EngineCore busy loop with `vt cuda: cudaMalloc: out of memory` (#1123). 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. ## The allocation, named and sized The log line could not 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)` (`:48-52`), where `what` is a compile-time literal. `bytes` is in scope and discarded. So the answer came from the code and the checkpoint rather than from the message. **It is `d.b.Alloc(nb)` in `ResidentWeight`, `qwen3_5.cpp:1010-1011`**, with `nb = w.bytes.size()` — for a routed-expert weight, the whole STACKED `[E*N,K]` keep-quant 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: | Configuration | Path | |---|---| | default (grouping on) | `MoeBlock:6615,6616,6620` → `KqGrouped:5694` → `ResidentWeight` | | `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()` (`:5578`), so on a device platform it falls through before the store is even constructed. Ruled out by reading rather than by assumption: `cuda_moe.cu` and `cuda_glue.cu` contain no allocation at all, and `BuildMoeMarlinResident` (`:6010-6215`, whose per-expert allocations are `:6049-6064` plus two repack temporaries at `:6094-6095`) is not on this path, because `MoeBlock` takes the fp4/Marlin arm at `:6555` under `const bool fp4 = !w.expert_gate_fp4.empty()` (`:6548`), and a GGUF populates `expert_*_kq`. The size came from re-censusing both GGUF tensor tables at revision `567d3e6ac26c5474b18311e619c04350fb9a5556` over all ten shards by HTTP range request, no tensor data downloaded: **1702 records parsed against the 1702 declared in `split.tensors.count`**, which is the coverage claim. | Tower | Bytes | | Count | |---|---|---|---| | IQ1_XXXS `ffn_{gate,up,down}_exps` | **1,275,068,416** | 1.1875 GiB | 276 | | Q2_K, block 92 (the `nextn` MTP block) | **2,818,572,288** | 2.6250 GiB | 3 | | all `*_exps` | **360,374,599,680** | **335.62 GiB** | 279 | Cross-check that the arithmetic and the running lane agree on the same weight: `1,275,068,416 / 512 = 2,490,368`, exactly the `slot_bytes=2490368` the row's W4 banner printed. **The budget needed the right instrument.** `nvidia-smi --query-gpu=memory.total,memory.free,memory.used` answers `[N/A], [N/A], [N/A]` on a GB10, and the `rc` fleet label records `vram=[N/A]M`. `cudaMemGetInfo` answers honestly — measured on `dgx:gpu0` under an `rc` hold, through `libcudart.so.13`: ``` 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. So 335.62 GiB of tower staging is 2.8x the whole machine; the load survives only because a borrowed tower costs zero anonymous bytes, and staging exhausts the pool after roughly `(119.631 - 62) / 1.1875 = 48` towers, partway through layer 16 of 93. This corrects the hypothesis in the issue body, which had the mechanism right and the quantity wrong: it is not ~6.5 GB of per-token staging that fails, it is the whole tower set being made device-resident once. Per-token slicing never runs. ## The refusal Keyed on the measured condition, never on "CUDA + GGUF" and never on an architecture name: refuse <=> needs_weight_staging AND budget != 0 AND needed > budget so a GGUF that fits the pool still loads on `--device cuda`, and every `--device cpu` load is byte-identical because a non-staging platform returns before a footprint is even computed. Three polarities are deliberate. **The footprint is a PER-TENSOR lower bound, and the sum is wrong in BOTH directions.** Per tensor it is `min(gguf_bytes, elems * model_dtype_bytes)`: a kept-quantized weight is staged verbatim, an expanded one at the model dtype, and which happens is a per-tensor loader policy this module does not try to predict. Each term is therefore a true lower bound on that tensor's staged size. **The sum is not a lower bound on the load**, and the first version of this pull request claimed it was ("so the refusal can never over-refuse"). It can over-refuse: a tensor counted and never staged is a positive over-count, and one is present on every default load — the MTP / `nextn` block, block 92, 20 tensors, 8,940,488,704 of 397,245,341,184 bytes, **8.33 GiB, 2.2506 %**. A budget in `[what a default load stages, what the bound counts)` rejects a weight set that fits. It also under-counts, by everything that is not a weight, which is larger. The two errors are on DIFFERENT quantities and never cancel, so "the under-count dominates" is an argument about a number the refusal does not compare. Both directions are named in `gguf_device_fit.h`, the over-count is now pinned executably rather than described, and both remainders are owed — the over-count to #1136, the under-count to `KV-WARMUP-PROFILE`. **The budget is the pool TOTAL, not the free bytes**, so the verdict does not move with contention. **A budget of 0 means UNKNOWN, which is not a verdict.** A caller that cannot learn the budget declines to decide, because refusing a load on an unknown budget would break every device whose budget nothing probes. That is the opposite polarity from `gemma4_moe.cpp:506`, which refuses a device allocation on unknown, and the difference is deliberate: there a hung `hipMalloc` is worse than a host fallback, here a false refusal is worse than a late failure. **Which platforms this covers**, stated precisely because the first version got it wrong in both halves. The two predicates coincide on exactly one platform: `needs_weight_staging()` is true only on `CudaPlatform` (`platforms/cuda.cpp:71`) and only `CudaPlatform` probes a budget. So EVERY NVIDIA GPU this build runs on — discrete or GB10 — 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, `platforms/rocm.cpp:74`): they read the mapping where it lies, so there is no staging allocation to fail and the refusal is inapplicable rather than owed. What is owed on ROCm is the separate `Backend::DeviceMemoryInfo` capability (#1126). The budget arrives as NEW data on `ResidencyPolicy`, probed once by `CudaPlatform`'s registrar. It is deliberately **not** `Backend::DeviceMemoryInfo`, whose comment claimed "ROCm/CUDA override with hipMemGetInfo/cudaMemGetInfo" while only ROCm does (`src/vt/rocm/rocm_backend.hip:338-345`) — so `Gemma4MoE`'s device-expert LRU is dead on every CUDA device today, and waking it is a behaviour change with its own measurement. That comment is corrected here, in **both** places that carried it: `include/vt/backend.h:78-93` and `gemma4_moe.cpp:440-448`. The capability is #1126. `ResolveModelDeviceType` is extracted from `SelectQueueForModel` because the check must know the target device before any weight I/O while the load's queue is not created until after the weights load. One description, not two — and on the AUTO arm that description CREATES A QUEUE, for the reason in the next section. Reachable from two production entry points, both of which funnel through `LoadedEngine::FromModelDir`: the server binary, on its embedding lane (`server_main.cpp:959`) and its generative lane (`:1123`), and the C ABI (`vllm_c.cpp:766`). `examples/bench/bench_core.h:586` also calls it and is an example, which `AGENTS.md` says is not a production entry point. ## Round two: the review repairs (#1136) A fresh review at `e7d0a1f7c` returned FAIL on one gate and five claims. One of those claims guarded a real defect. **`ResolveModelDeviceType` could remove a working load.** The header claimed it picks what `SelectQueueForModel` picks, so "the two cannot drift". On the AUTO arm they could. `SelectQueueForModel` wraps `CreateQueue()` in a try/catch and falls back to CPU, and its own comment says why: "a platform can be registered while CreateQueue still fails, and CPU must remain reachable". The resolver asked `CurrentPlatform()` and stopped, so on such a box it answered `kCUDA` while the load ran on CPU — and the refusal then rejected a checkpoint by naming `'cuda'` for a model that previously loaded and served on CPU. Whether `CreateQueue()` fails is knowable only by calling it, so both callers now share one `ResolveAutoDevice` that creates the queue and hands it to whichever caller wants one; `ResolveModelDeviceType` destroys the queue it does not use, because `vt::Queue` is a non-owning handle with no destructor and dropping it would leak the stream. The cost is one extra stream created and destroyed per auto-arm GGUF load, stated in the header rather than hidden. **`CudaPlatform`'s policy assembly is now pinnable without a CUDA build.** #1123 recorded "delete the `device_memory_total_bytes` assignment" as an owed mutation because `platforms/cuda.cpp` compiles only in a CUDA build. The four-line assembly moved to `CudaResidencyPolicy` in `vllm/platforms/interface.h`, which compiles everywhere, and `test_platform.cpp` pins every field. `cuda.cpp` keeps only the `cudaMemGetInfo` probe it alone can make; that call and the constructor threading are still unproven and still said to be. Its one-line edit was checked with `g++ -fsyntax-only` against a stub `cuda_runtime.h`, and the instrument was itself verified: a deliberate typo in the same expression returns rc 1 naming it. **Fourteen stale or wrong anchors and names**, each re-derived against the final tree rather than re-quoted (the two this change authored itself are counted in the next paragraph, not here): `cuda_backend.cu:75-81` is `:77-81`; `BuildMoeMarlinResident` is `:6010-6215`, not `:6010-6062`; `gemma4_moe.cpp:439-447` and `:494-506` are `:449-455` and `:506`; `CheckGgufDeviceFit` never existed, the function is `CheckDeviceWeightFit`; the call-chain table mixed call sites with definition lines and now declares one convention and keeps it; and the header's upstream anchor `vllm/v1/worker/gpu/model_runner.py:504,647` pointed at a `DraftModelSpeculator` `set_attn` call and a `torch.zeros`. The startup memory profile is `GPUWorker.determine_available_memory` (`gpu_worker.py:451-495`) around `profile_run` (`gpu/model_runner.py:682`), and it takes the weight bytes as an INPUT recorded after the load completes (`gpu/model_runner.py:315`) — a stronger statement of why upstream has no counterpart than the one that was there. That `:504,647` pair was copied from `.agents/engine-matrix.md`'s `KV-WARMUP-PROFILE` row, which still carries it plus a third stale anchor; filed as **#1139** and NOT fixed here, because the fix is one cell in `engine-matrix.md` that PR #1119 is concurrently bumping alongside the hardcoded count in `check-agent-record.py`. **The anchors this change itself moved.** Inserting ~45 lines near line 100 of `model_loader.cpp` shifts every absolute line citation below it. Measured between `e7d0a1f7c` and this head by comparing the TEXT at each cited line: **203 moved line references over 109 citing sites in 45 files**, 10 unmoved. Two were authored by this change (`model_loader.cpp:135-141` in the reachability gate, `:1411-1412` in the arithmetic gate) and both are corrected and added to this round's anchor verifier — the trap the round before hit with `platforms/cuda.cpp:67`. The other 107 are not swept, and not for effort: several were already stale at `e7d0a1f7c` (`model-matrix.md:197` cites `:184-223` as the "live loader" while line 184 there is `static const bool once = [] {`), so rewriting them from the current tree would launder pre-existing debt into a clean-looking record. Filed as **#1143** with the measurement and three candidate fixes. **One contradiction was self-inflicted and caught before it shipped.** `gguf_device_fit.h` opened by calling the footprint "WRONG LOW rather than wrong high" and closed twelve lines later by saying it can over-refuse. The fix is naming the scope: the sum is a lower bound on staging EVERY TENSOR IN THIS FILE, which is not a lower bound on what one load stages, since a load may stage a subset. The per-tensor term is exact-or-low; the SET is exact-or-high. That is where the over-count lives, and it is why the field keeps the name `lower_bound_bytes`. **The false comment had a second copy.** The correction to `include/vt/backend.h` is the finding; the audit for it found the same "(ROCm/CUDA)" claim at the seam's only call site in `gemma4_moe.cpp`. Correcting the header alone would have left the falsehood in the tree. **`docs/USAGE.md` and `docs/ENVIRONMENT.md`** now state the platform coverage correctly and warn an operator about the over-count window. **The "Release (what CI uses)" claim was false.** `build-test-cpu` (`ci.yml:903-911`) passes no `CMAKE_BUILD_TYPE`, so CI's test lane defines no `NDEBUG` and has asserts live. This round was configured and gated in exactly that configuration. ## Gates Configured as CI's `build-test-cpu` does — `-DVLLM_CPP_BUILD_TESTS=ON -DCMAKE_EXPORT_COMPILE_COMMANDS=ON`, no `CMAKE_BUILD_TYPE`, so `assert` is live. `origin/main` advanced to `a7583ac75` during the round and is merged in with an AUTHORED merge commit, because `agent-preflight.sh` SKIPS both trailer gates when `origin/main` is not an ancestor of HEAD — and a skip still exits 0, so a reader who trusts the return code reads two unexecuted gates as green. That is documented at `agent-preflight.sh:35-42` and `--fail-on-skip` is the opt-in for it; every preflight figure below is from a post-merge run with that flag. Build and gate were re-run in full after the merge. **That merge silently corrupted a record, and a gate caught it.** Both sides had appended after the same `#837 / #1117 / #1118` tail of `.agents/issue-index.md` — this branch `#1136 / #1139 / #1143`, main `#1129 / #1130 / #1134` — and the `merge=union` driver emitted the shared three-row block TWICE, once as each side's context. `git merge` exited 0 and reported no conflict. `check-agent-record.py` and `test_agent_record` both failed with "issue #837 listed twice. Under `merge=union` a duplicate is what two branches appending the same issue look like". Deleting the second copy was not enough, and a SECOND gate said why: `issue-index append-only` stayed red, reporting main's `#1129` and `#1130` as REMOVED, because the two files disagree on where the shared block sits. Main's tail runs `#1129, #1130, #837, #1117, #1118, #1134`; this branch's ran `#1127, #837, #1117, #1118, #1136, ...`. Any resolution that keeps this branch's ordering reads as a reorder, and a reorder of an append-only file is a deletion plus an insertion however the rows are spelled. Resolved the way `AGENTS.md` prescribes for a keyed record — take the COMPLETE target-branch version, then apply the scoped edit again: `origin/main`'s file verbatim with the seven rows only this branch has appended at the end. Verified rather than assumed: every line of main's version is present, `git diff --numstat origin/main` reports **`7 0`** (seven insertions, zero deletions), 322 rows, no duplicate id. `merge=union` makes concurrent appends merge without a conflict; it does not make a RELOCATED append safe, and the duplicate is what a relocation looks like on the way out. Two gates were needed to see the whole of it, which is the argument for both. | Gate | Result | |---|---| | configure into an empty build dir, then full build (re-run after the merge) | rc 0, 508 `Built target` lines, **0 ENOSPC**, 0 `error:`, 0 `warning:` | | `test_gguf_device_fit` | 8 cases, 61 assertions, 0 failed, rc 0 (was 7/53) | | `test_gguf_device_fit_reach` | 5 cases, 23 assertions, 0 failed, rc 0 (was 3/14) | | `test_platform` | 12 cases, 95 assertions, 0 failed, rc 0 (was 11/86) | | `test_device_selection` (regression on the resolver) | 2 cases, 11 assertions, 0 failed, rc 0 | | `ctest --test-dir build` (serial, exactly CI's command) | **506 tests, 100 % passed, 0 failed, rc 0**, 293 s, on the merged tree; 2 expected skips (`test_modelopt_mixed_precision_checkpoint`, `test_voxtral_e2e`). 506 rather than 505 because the merge brings `test_backend_cross_device`, so the count moving is itself the proof the merged test is registered and ran. The gate was run THREE times over the round and every run was 100 %, at host load averages of 16-134, with no starvation flakes | | `scripts/agent-preflight.sh --staged --fail-on-skip` | rc 1 — **80 of 81 gates ok, ZERO skipped, one failure, and it is not this change**: see below | The single preflight failure is `commit-trailers`, on the two unauthored MERGE commits already on the task branch (`058f5f31a`, `e7d0a1f7c`): both carry the default `git merge` subject and an empty body, and `scripts/check-commit-trailers.py:331,336` walks the merge-base range with no merge exclusion. This change's own commit passes the contract on its own range (rc 0), and `main` is squash-only so neither merge commit ever reaches it. Repairing them means rewriting history that is already pushed, which `AGENTS.md` forbids doing by force, so it is a branch-history decision rather than a code one. Every other gate — all 80 — is `ok`, including `check-agent-record` (`ENGINE=157`, no collision with #1119), `issue-index append-only`, `check-env-doc`, `check-public-doc-tables` and `commit-style`. **RED captured first, on a compiling stub**, because a red that fails to BUILD reads as a pass: | Red | Result | |---|---| | the two AUTO-arm cases, against the unfixed resolver | 5 cases, **2 failed, 4 assertions red**, rc 1 — one because the refusal fired for a load that runs on CPU, one because `queues_created` stayed 0, which is the divergence itself | | the `CudaResidencyPolicy` case, against a stub returning a default policy | 12 cases, **1 failed, 7 assertions red**, rc 1, and the case count MOVED 11 → 12 | The case count is asserted to move because a test class that never compiles into the binary prints a clean pass: the first run of `test_platform` here read 11 cases and 86 assertions — the pre-change numbers — because the build had already compiled that file before the case was added. The `test_gguf_device_fit` over-count case has **no red-first**, and that is stated rather than papered over: it characterises behaviour the tree already had, so its only evidence is mutation M4 below. ### Mutations Seven, each printing its `applied` proof (sha256 before/after plus a non-empty `git diff --stat`), its build rc, and a NON-ZERO doctest case count taken from the LAST `test cases:` line. A mutation that did not apply, or did not build, is reported INVALID and never as a pass. Every file was restored from a byte backup — never `git checkout --`, which would discard an uncommitted repair — and the restore verified by sha256. Every literal was asserted to match EXACTLY once before it was applied, so a mutation cannot silently hit the wrong site or no site. | # | Mutation | applied | compile | Result | |---|---|---|---|---| | M1 | the resolver goes back to the bare `CurrentPlatform()` query (the pre-fix code) | `16389220`→`8396036f`, 20 lines | rc 0 | reach: 5 cases, **2 failed** — CAUGHT. `test_device_selection` 2/2 green, which is correct: it covers only the explicitly-named arm | | M2 | the probe queue is dropped instead of destroyed | `16389220`→`54e7ace6`, 1 line | rc 0 | reach: 5 cases, **1 failed** — CAUGHT | | M3 | `CudaResidencyPolicy` stops carrying the probed budget | `5e51dfc0`→`483449d7`, 1 line | rc 0 | platform: 12 cases, **1 failed** — CAUGHT | | M4 | the footprint EXCLUDES `*.nextn.*`, i.e. the over-count is removed | `4edabac3`→`fcbcf384`, 1 line | rc 0 | fit: 8 cases, **1 failed** — CAUGHT. This is the over-count case's only evidence, because it has no red-first | | M5 | delete the production call site (`if (fit.refuse) throw ...`) | `16389220`→`22879f14`, 1 line | rc 0 | reach: 5 cases, **2 failed** — CAUGHT | | M6 | strict `>` becomes `>=` | `4edabac3`→`8b409c96`, 1 line | rc 0 | fit: 8 cases, 2 failed; reach: 5 cases, 1 failed — CAUGHT | | M7 | drop the non-staging early return | `4edabac3`→`035a886e`, 1 line | rc 0 | fit: 8 cases, 1 failed; reach: 5 cases, 2 failed — CAUGHT | Every row above was produced in ONE pass against the head being merged, after the last repair, so none of it is evidence about an earlier tree. `git status` is empty afterwards and each file's sha256 matches its pre-mutation value byte for byte. **Still owed and still said to be**: the `cudaMemGetInfo` call in `platforms/cuda.cpp` and the constructor that threads its value are not mutation-proven, because no CUDA toolkit is reachable from this host. #1123 owed the whole policy assembly; this round moved the assembly out and pinned it, so what remains owed is the probe alone. Its file is syntax-checked, with a verified positive control. ## Owed, filed rather than folded in - **#1124** — the device-side expert slot store this refusal stands in for. Four concrete pieces, sized at 2790 slices per token times 2,490,368 bytes = 6.95 GB per token. Not started here because W7 owns the pluggable backing store and the CPU arm's own decode bandwidth is still VOID. - **#1126** — `CudaBackend::DeviceMemoryInfo`, and the Gemma4 measurement that has to come with it. Both false comments about it are corrected here; the capability is not built. - **#1127** — moving `VT_DEVICE_WEIGHT_BUDGET_BYTES` into the `vllm_cpp` config namespace once #1119 lands. - **#1136** — the bound's over-count direction. Not closed because 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 an exclusion that is wrong under-counts toward zero, restoring the exact failure this row removed — on a device nobody on this fleet has to measure the change against. - **#1139** — `KV-WARMUP-PROFILE`'s three stale upstream anchors, blocked on the `engine-matrix.md` / `check-agent-record.py` record lock that #1119 holds. - **#1143** — `model_loader.cpp`'s 109 line-number citations across 45 files, and the three candidate fixes for the class. Needs a row of its own. All six are listed under `## Owed` in [`expert-streaming.md`](.agents/specs/expert-streaming.md). Coordination note for #1119: the two changes overlap on `src/vllm/entrypoints/model_loader.cpp`, `include/vllm/entrypoints/model_loader.h`, `docs/USAGE.md`, `docs/ENVIRONMENT.md` and `.agents/issue-index.md`. This one adds no roadmap row and touches neither `.agents/engine-matrix.md` nor `scripts/check-agent-record.py`, specifically so it cannot collide with #1119's `ENGINE_ROWS` bump; `check-agent-record.py` still reports `ENGINE=157` here. The issue-index rows are appends. Closes #1123. Closes #1136. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [Claude Code] --------- Co-authored-by: Ettore Di Giacinto <mudler@localai.io>
The base moved while round three was under repair, including #1132, which adds a VT_DEVICE_WEIGHT_BUDGET_BYTES row to the same environment table this row rewrites. Merging rather than rebasing, because a force-push would destroy the merge base the reviewed head is anchored to. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [Claude Code] # Conflicts: # docs/ENVIRONMENT.md
The tier that makes a 370 GiB checkpoint fit in 119 GB was environment-only. It is now a config key, the mirror is untouched, a typo in that key is refused at every level of the document, and a second engine in one process can install a partial document without either failing or silently dropping the first engine's fields.
vllm.cpp offloads weights at two tiers and only one was reachable from the
config surface.
--offload-config '{"uva":{"cpu_offload_gb":N,"cpu_offload_params":["experts"]}}'runs from the server flag through
EngineParamstoLoadedEngine::FromModelDir.The tier below it, the one that makes
Qwen3.8-2.4T-A95B UD-Q1_0serve on asingle 119 GB GB10 because ~330 GiB of experts stay borrowed in the mapping, was
VT_GGUF_MMAP,VT_GGUF_PREFAULT,VT_MOE_EXPERT_STREAM,VT_MOE_EXPERT_STREAM_SLOTS,VT_MOE_EXPERT_STREAM_SLOT_BYTESand nothing else.Qwen35ExpertStreamRequestedread onegetenvand had no other input. Theasymmetry was the more awkward for the first tier's config already being
expert-aware, so a user reaches for
--offload-configto control expertplacement, gets the GPU tier, and finds nothing for the tier the big-model case
needs. Closes #1110.
The five knobs are now also keys under a namespaced
vllm_cppobject inside thatsame flag:
include/vllm/config/offload.his unchanged. It transcribesvllm/config/offload.pyline for line and upstream has no disk tier at the pin(
OffloadBackendisLiteral["auto","uva","prefetch"],offloader/uva.py:21isCPU-blanket UVA,
offloader/prefetch.py:557-560is cpu-only, and nothing reads aweight off a file at inference time), so the extension is vllm.cpp-original by
construction rather than by preference. One flag still covers one user-facing
concept, two parsers read one string, and the key is literally named
vllm_cppso nothing about it reads as upstream. The C ABI needed no new field for the same
reason:
vllm_model_params.offload_configis one string carrying both halves, andinclude/vllm.hnow documents which half of it moves weights.The schema's own rule, applied to both ends of the install
An absent field means UNCHANGED. That binds two pieces of code, and the first two
shapes of this row broke it in both directions, each with its own user-visible
failure on a legal two-model process. They are one root cause and are fixed as one
thing (#1133 H1, H2).
FrozenFieldstreated absent as a change. It comparedin.expert_streamagainst the stored optional, and
nullopt != engagedis true, so once a processhad latched the streaming decision,
{"vllm_cpp":{"mmap":{"enabled":true}}}on asecond engine threw
std::logic_errorout ofLoadedEngine::FromModelDirand cameback as
VLLM_ERR_MODEL_LOAD. The earlier narrowing had moved that failure, notremoved it. The message made it worse by asserting that "accepting this would
record a configuration the engine is not running" — for a document asking for
exactly what the process had resolved, the engine was running it.
The install treated absent as a clear.
g.config = configreplaced wholesale,so a partial second document turned
expert_stream=onwith 8000 slots into OFFwith 64, with no diagnostic, and the slot store reads those values lazily — on the
first slice taken, which can be long after the second engine loaded. The wholesale
replace predates the per-field narrowing; the narrowing widened its reach, because
before it any differing document after a decision threw, so the drop could not
happen once anything had been decided.
Both now mean what the schema says. A field is scored frozen only when the
document SETS it. The comparison is against the decision actually taken — the
cached streaming answer, the geometry the store was built with — resolved through
the same function the production resolver calls, so a document that agrees with
the running engine installs, and so does one the environment overrides anyway,
while one that would make a resolver return something else still throws. The
install merges field by field, which makes the empty document a no-op by
construction rather than by a special case. The refusal message quotes the
decision rather than the stored document, because after the merge the stored
document is not what is in force either.
The streaming flag carries its own answer to make that comparison possible: one
tri-state atomic (nothing decided / decided off / decided on) rather than a bool
beside a value, so a reader takes one relaxed load and there is no ordering
question between a fact and a value in two places. The geometry's numbers move
from a file-static into
Globalundermu, beside the flag's writer. Both storesare idempotent and both are monotonic, and that is what carries
relaxedhere —not a synchronises-with edge, which a relaxed store outside
mudoes not give.A typo is refused, at every level of the document
parse_offload_config_jsonlooks its three keys up by name and never enumeratesthe document, which is what makes a namespaced sibling key workable at all — and
it is also the hazard, because a misspelling is then invisible to both parsers. So
the extension parser enumerates the whole document.
That enumeration first closed only the top level and the inside of
vllm_cpp,while
include/vllm.hpromised an unknown key anywhere. Measured:{"uva":{"cpu_offload_gbb":1}},{"uva":{"cpu_offload_GB":10}}and{"prefetch":{"offload_groupsize":8}}were all ACCEPTED, each giving a 0 GiBbudget or a group size of 0 under a document the operator believes configures
offloading (#1133 H3). The claim is made true rather than scoped down, because
refusing is mirror-faithful here too:
UVAOffloadConfig(
vllm/config/offload.py:15-16@555967922) andPrefetchOffloadConfig(
:47-48) each carry@config, whose body setsConfigDict(extra="forbid")(
vllm/config/utils.py:68-69), so upstream refuses a nested typo and the tolerancewas the deviation. The enumeration stays in the extension parser and lists NAMES
only, so
parse_offload_config_jsonis untouched and keeps sole ownership of thosefields' types, defaults and bounds.
A silently ignored
{"vllm-cpp":…}or{"vllm_cpp":{"mmapp":…}}starts a serverrunning this tier at its defaults — prefault ON, streaming OFF, the two settings
the big-model case exists to change — which the operator meets as an out-of-memory
kill rather than as an error. It now says
offload config: unknown key "vllm-cpp" (expected one of: offload_backend uva prefetch vllm_cpp)before the multi-GB load.
Precedence, and what latches
Precedence is environment variable, then config, then built-in default. Those
variables exist so a benchmark arm is switchable without restarting the server
with a new document and an A/B in flight depends on it, so an override that could
not turn a configured knob back off would be useless:
VT_X=0beats a configtrue. The install prints one line naming the fields of the document it installedand a second naming every variable that would win over one of them, because a
document silently overridden by something exported weeks ago is the one way that
polarity hurts. The first line reports what was asked for and not what the engine
resolves, and exactly one of the five is the reason: the streaming answer is
cached the first time it is asked, so resolving it at install would move that
decision ahead of the load.
prefaultandslotscould be resolved there, andmmap/slot_bytesneed a default only their caller has, so reporting the documentfor all five is a consistency choice on top of the one real constraint.
Two of the five knobs latch a decision; three do not.
ResolveExpertStreamRequestedcaches its answer in a function-local static(
Qwen35ExpertStreamRequestedis the model-side name and a pure delegation to it),and the slot store's
slots x slot_bytesreservation is fixed when the store isbuilt.
SetWeightResidencyConfigthrows for exactly those two, and only for adocument that would change them.
mmapdoes not latch —GgufLoadPolicy::FromEnv()runs per load — and neither does
prefault, whose site drops its static altogetherhere: one
getenvper span is nothing beside the megabytes of pages the functionthen reads, and it buys a config that cannot be missed plus an A/B whose two arms
actually differ. The loader installs in
FromModelDir's first statement block,ahead of the offloader and every path and weight operation.
The five knobs get one named resolver each, and that is not decoration. The
polarities differ and one is deliberately odd:
VT_GGUF_MMAPandVT_GGUF_PREFAULTcompare the whole value against""/0/false/off, whileVT_MOE_EXPERT_STREAMexamines only the first character, soVT_MOE_EXPERT_STREAM=falseis ON anddocs/ENVIRONMENT.mdsays so. Routing allfive through one helper would have normalised that silently, and a row whose
subject is where a value comes from must not change what a value means.
VT_MOE_EXPERT_STREAM_STATS_EVERYstays environment-only by decision: it changesa diagnostic cadence rather than what the process reserves, so it is the
instrument and not the configuration, and the config surface refuses it as an
unknown key rather than accepting and dropping it.
Also fixes #1109 in flow.
docs/ENVIRONMENT.mdgaveVT_GGUF_PREFAULT's defaultas off while the code has always defaulted it on. This change writes that default
into a resolver and a config key, so shipping it beside a table stating the
opposite would put the contradiction inside one commit. The direction matters to a
user and not only to a document: prefault on is exactly what a model larger than
memory has to turn off.
What the test shape was hiding
Eight mutations and a fresh review missed both behaviour defects, and the reason is
worth a case shape rather than another sentence. Every latch case installed a COPY
of the first document —
mmap_too = cfg,mmap_only = sizes— or the empty one,so the second install always restated the frozen field at the value it already
had, and that is the one shape in which neither defect can arise.
Two DIFFERENT partial documents in one process is the shape that distinguishes
them. It is now a case at the unit level and again through
FromModelDirandvllm_engine_load: engine A installs the full document and takes both decisions,engine B brings
mmapalone, engine C bringsprefaultalone, and each mustinstall without dropping what came before. The reach case has to read the global
rather than the return code, because a refused install and a missing checkpoint
both leave
vllm_engine_loadasVLLM_ERR_MODEL_LOAD;mmapasked for FALSEagainst an installed TRUE is the discriminator. Neither new case assumes the
streaming answer's value — it is a per-process static, so the cases read it and
assert relative to it, or a single-case
-tcrun would pass or fail on test orderrather than on the code. Two more cases pin what the refusal must NOT do: accept a
document that agrees with the decision while the stored document is empty (the
shape the environment produces), and accept one the environment overrides.
Mutations
Sixteen for this round, each applied alone, with
git diff --statand the file'ssha256 printed before and after so a never-applied edit cannot read as a pass, the
build's own exit status and an ENOSPC count printed beside every result so a
non-building mutation is INVALID rather than a pass, a non-zero doctest case count
required, the LAST
test cases:match read, every rc captured directly rather thanthrough a pipe, and the tree restored by byte copy with the sha256 compared. Three
results are worth naming.
R6andR12first came back INVALID on-Werror=unused-function, which is anon-building mutation and not a passing test; both were re-run in a form that keeps
the call and discards its result.
R12then came back GREEN, and that was theassertion's fault rather than the code's: the message check was
Mentions(e.what(), "expert_stream_slots=8000"), and at that point the storeddocument also held 8000, so quoting either produced the substring. The check moved
to the one moment the two differ — a refusal taken while nothing is stored, where
quoting the document produced
environment/default— and goes red there.R15(the reset leaving a stale built geometry) is GREEN and recorded as GREEN: nothing
can observe it, because the numbers are read only while the geometry latch is set
and the reset clears that too. The code says so where the line is, rather than
letting a reader assume a gate exists.
The reachability mutation deletes the install call site in
FromModelDirand turnsthe reach suite red (4 failed) and the server suite red (3 failed).
Records
TEN stale or imprecise
path:lineanchors in this row's own spec were found bychecking anchors last against the final tree, and ALL TEN were already wrong at the
reviewed head. For
qwen3_5.cppandinclude/vllm.hthat was verified by reading thesame line numbers out of
git show HEAD~1:<path>; for the other three files it followsfrom this change not touching them, so their lines cannot have moved. The Port map's
four
VT_MOE_EXPERT_STREAM*rows pointed at a comment and at three lines of thestatistics printer rather than at the constructor beside them,
GgufLoadPolicy::FromEnvwas cited 85 lines past itself,include/vllm.h:436fellmid-sentence, and the
FromModelDirrange started one line late and ended 30 linesearly. Three stale anchors belonging to OTHER rows are named in the spec and left
alone, because correcting them would widen this diff into three unrelated
subsystems.
## Owed's unreached-entry-point gap named #1122 as its owner — the very issue thispull request closes, so on landing the gap would have had no open issue. Filed as
#1135 and pointed there, in the spec and in
docs/USAGE.md.Gate
CPU build, documented recipe, on the merged tree.
origin/mainmoved twice duringthis work, to
d1e5e9bc0and then toa7583ac75, and both are merged in. That mattersfor more than currency: the trailer and commit-style gates scope themselves to
origin/main..HEAD, so before the first merge they reported nothing at all about thisbranch.
d1e5e9bc0merge, 974translation units: exit 0, zero compiler warnings, zero ENOSPC lines. Rebuilt
incrementally on the
a7583ac75merge: exit 0, zero warnings, zero ENOSPC. The diskran between 89% and 98% full throughout, and a full disk produces a link error that
reads exactly like broken code, so an ENOSPC count is printed beside every build and
mutation result here.
ctest --test-dir build -j 3on the final head, over a second clean rebuild from anempty build directory (974 translation units, exit 0, zero warnings, zero ENOSPC):
508 of 508 passed, exit 0, two
pre-existing environment skips (
test_modelopt_mixed_precision_checkpoint,test_voxtral_e2e). The same command on thed1e5e9bc0merge was 507/507 exit 0,twice; the extra test comes from
a7583ac75.test_engine_core_procfailed one ofthose runs, and a 30-run sweep of it alone gives 29 pass / 1 fail. It is
#1052, open since before this work,
also recorded at
.agents/environment.md:540-547;grep -cfor this row's symbols inthat suite returns 0 and this change touches no file under
src/vllm/v1/. A duplicateissue was opened for it during this gate and closed as a duplicate; the measurement
went to test_engine_core_proc's immediate-shutdown case is load-dependent: a FIXED 1000-frame budget racing an unbounded producer, and no issue names it #1052 as a comment. What is NOT claimed: that it was red before this change,
because no
ctestwas run on the pre-change tree.test_serve_low_toolsalso starvedonce at
-j 6under load average 47 and passes serially.scripts/agent-preflight.sh --staged --fail-on-skip: exit 0, all gates green,no gate skipped —
commit-trailersandcommit-styleboth ran and both ok, andeach was additionally run by hand over
origin/main..HEAD.test cases:match:test_weight_residency_config17 cases / 250 assertions,test_weight_residency_reach7 / 76,test_expert_stream_latch1 / 9,test_serve_residency_config6 / 64 (1 skipped, the re-exec'd child).pr-sizewas already red on the reviewed head, and neither earlier round saidso. It is a required check, and it had been failing since this row's first commit
because
scripts/check-pr-size.pyclasses any change toscripts/check-*.pyas agovernance_checkerand demands a non-empty edit in the paired test file — and thisrow changed
ENGINE_ROWS157 -> 158 incheck-agent-record.pyand oneRUNNABLE_BASELINEentry incheck-gate-commands.pywith no test edit at all. Whatslipped past three reviews is that both are MARKS with no branch to invert; what the
contract is right about is that the constant was then the only artifact. Three cases
make both marks executable, in the idiom the neighbouring credits for Binary release? #117, vllm-serve aborts on --enable-auto-tool-choice and --trust-remote-code, so 89 of 157 official recipe commands fail to start #606,
vLLM-Omni has no parity pin: H3 W3+, LTX-2.5 and ~40 omni-only architectures (the whole TTS family included) cannot be gated against any oracle #633, Recorded line anchors are never validated against what they name, and ACTIVE rows are not anchor-checked at all #632 and Speculative decoding: MTP k>1, then dynamic and adaptive speculation depth #81 already use. Mutation-proven, and one mutation took three attempts
to become VALID: removing the row from the engine matrix aborted the suite's
setUpClasson a dirty baseline, so 51 of 78 cases ran and the new case neverexecuted — a mutation that reads as a pass. Only removing the row, the pin and the
row's claim file together yields a runnable tree, and then all 78 cases run and the
new case is one of exactly two failures.
check-pr-size.pygoes from two ERRORs toOK.
merge=unionon.agents/issue-index.mdconcatenates both sides of a changed region withoutdeduplicating, so merging
a7583ac75doubled ROCm Gemma-4 hyp B: GetBlas single TLS destroys hipBLAS handle on peer-MoE device hop #837, LTX-2.5: A2VidPipelineTwoStage has no recipe row, so audio-to-video rides the distilled two-stage trajectory #1117 and LTX-2.5: LoRA adapters fuse once at load, so no recipe can put the distilled adapter on stage 2 alone #1118 byte-identicallyand
check-agent-record.pyrefused it by name — correctly, because under union aduplicate is exactly what two branches appending the same issue look like. Fixed the
way AGENTS.md prescribes for a keyed record: take main's complete file, re-apply this
branch's scoped append.
a7583ac75's version is asserted to be a BYTE-EXACT PREFIX ofthe result, all 315 of its rows textually unchanged, with this branch's five rows
after it.
pushed and then amended twice — once to correct claims this session's own audit found
wrong — which made the local tip a sibling of the remote one and the push a
non-fast-forward. The pushed commit is merged back in as a parent rather than
overwritten; the resulting tree is byte-identical to the pre-merge one (
git diffofthe two is empty), and the reviewed head
c7fa7084bandorigin/mainare both stillancestors, each verified with
git merge-base --is-ancestor.row/ENG-EXPERT-STREAM-DEVICE-FIT, which alsoedits
model_loader.{cpp,h},docs/USAGE.md,docs/ENVIRONMENT.mdand the index)produce the SAME single conflict, in
docs/ENVIRONMENT.md, where that branch appendsa
VT_DEVICE_WEIGHT_BUDGET_BYTESrow against the expert-stream table rows this branchrewrote; the resolution is to keep this branch's rows and append theirs. Everything
else auto-merges,
model_loader.cppincluded — and a cleanmerge-treeis not a treethat builds, so the merged tree was materialised and compiled: library exit 0, zero
warnings, and
test_weight_residency_config17/17,test_weight_residency_reach7/7and that branch's own
test_gguf_device_fit_reach3/3 all pass on it. The trial mergewas then aborted and the head verified identical to the pushed SHA with a clean tree.
Spec committed before the implementation. Row
ENG-RESIDENCY-CONFIGstaysACTIVE, notDONE, and two things keep it there, both under the spec's## Owed.The 370 GiB checkpoint has not been driven through the JSON form on a box that can
hold it: everything here is CPU-local, and GPU work on this fleet goes through a
lease rather than a shell. And the config form reaches the generate server path and
the C ABI but not
vllm-cli, nor the server's pooling and transcription paths,which build their engine parameters without the offload document at all — the
mirrored
uva/prefetchhalf included, and since before this key existed. That gapis now #1135, and
docs/USAGE.mdsays so beside the config form rather than leavinga reader to discover it.
Closes #1110. Closes #1109. Closes #1122. Closes #1133.
FOLLOWING_AGENTS_PROTOCOL
Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:claude-opus-5-1m [claude-code]