diff --git a/.agents/benchmark-record.md b/.agents/benchmark-record.md index 91b3578a5..845aa63e5 100644 --- a/.agents/benchmark-record.md +++ b/.agents/benchmark-record.md @@ -22960,3 +22960,102 @@ for CUDA runtime version`. The Jetson 540.4.0 driver cannot run a CUDA 13 runtim untried route is a CUDA 12.x toolkit for that driver; it was not pursued, because the dgx gate had already answered the question orin was there to support. No lease held; `orin:gpu0` returned to ready. + +## ENG-EXPERT-STREAM-DEVICE W0e/W0f — `--device cuda` decodes a 369.97 GiB checkpoint on a 119.631 GiB GB10, and the token gate fails on a near-tie (2026-08-19, `row/ENG-EXPERT-STREAM-DEVICE-W0F`, source `9c783a8be`, #1299) + +**Setup.** One `rc hold` on `dgx:gpu0` (GB10, `sm_121a`, CUDA 13.0.1 in +`vllmcpp-build:gb10`, driver 580.173.02). `Qwen3.8-2.4T-A95B UD-Q1_0` +(369.97 GiB, 10 shards) from the host at `/home/mudler/ckpt/qwen3.8-q1_0`. +Streaming ON, 4000 slots (9.28 GiB arena), greedy, 32 tokens, prompt ids +`760,6511,314,9338,369`. Both arms on the SAME binary and the SAME lease, page +cache dropped between them. Harness `w0e_gen`, logs under +`/home/mudler/work/es-w0e/logs` on `dgx.casa`. + +**Two VOID runs first, and why they were void.** The first two CUDA attempts +reproduced #1299 exactly (guard trip, zero decode steps) and looked like a W0f +result. They were not: the build ran `cmake --build build --target vllm`, which +is the STATIC library, while the harness links `build/libvllm.so`. That file was +still the previous day's pre-W0f build, `LIB_EXIT=0` all the same. W0e's own +`build.sh` had it right with `--target vllm-cli`. The corrected script records +the shared object's mtime and sha256 before and after and greps the built binary +for a string that exists only in the new code: `87c58eec` to `cf771cec`, marker +count 0 to 1. A build that does not relink is now reported as STALE rather than +as a pass. + +**G0-LIVE: PASS.** + +| Observable | CUDA | CPU | +|---|---|---| +| load | 266.330 s | 253.504 s | +| RSS after load | 61.20 GiB | 62.45 GiB | +| decode steps | 32 | 32 | +| `exhausted` at step 1 / step 32 | 6077 / 6077 | 6074 / 6074 | +| decode-phase `exhausted` delta | **0** | **0** | +| peak RSS | 97.75 GiB | 92.19 GiB | +| swap used at peak | 0 | 0 | +| container exit | `W0E_DOCKER_RC=0` | `W0E_DOCKER_RC=0` | + +**What W0f moved, counted rather than inferred.** An RSS curve cannot separate +"the branch declined and staged", "the branch re-homed and the pages did not come +back" and "something else allocated", so `MakeHostBytesDeviceAliasable` reports +its outcome per weight and `ResidentWeight` prints the split every 4 GiB on +`VT_LOAD_STATS`. First-forward totals, at the point re-homing plateaus (call +1361): + +| Outcome | Bytes | +|---|---| +| re-homed into an aligned host block, then aliased | **60.793 GiB** | +| declined, misaligned GGUF borrow, still staged | ~9.2 GiB | +| aliased in place (already 256-aligned) | 0.02 GiB at that point | + +On the CPU arm the same counter reads **0 calls**, which is the live control that +the branch is platform-gated rather than an argument that it is. + +**G0-CORRECT: FAIL, on a measured near-tie.** + +``` +CPU 11751,13,11751,369,264,3177,7172,303,279,17631,919,314,9338,11,383,279,... +CUDA 11751,13,11751,369,264,3177, 303,9338, 13, 9338,369,264,3046,303,4357,13,... + ^ first divergence, step 7 +``` + +The CPU arm on this binary reproduces its four-times-recorded ids byte for byte, +so the divergence is between the arms and not W0f. An instrumented CPU run +printing the top-2 logits per step gives the reason: + +| step | top-1 | logit | top-2 | logit | margin | +|---|---|---|---|---|---| +| 5 | 264 | 18.954491 | 279 | 18.668240 | 0.286251 | +| 6 | 3177 | 19.375208 | 6037 | 18.425795 | 0.949413 | +| **7** | **7172** | **18.779411** | **303** | **18.514702** | **0.264709** | +| 8 | 303 | 20.953234 | 383 | 18.930481 | 2.022753 | +| **9** | 279 | 19.850554 | 9338 | 19.827751 | **0.022802** | + +At the divergent step the CPU arm's own runner-up IS the token CUDA emitted, +1.4 % behind; one step later the margin is 0.1 %. The two arms rank the same +candidates and disagree about a coin flip. The declared gate still fails and the +wave still stops. + +**G0-SPEED: VOID and NOT claimed**, because a speed number behind a failing +correctness gate is the #912 F1 shape. Taken for the record only, over the 31 +DECODE steps of each arm (step 1 is prefill and is excluded), interleaved on one +lease: + +| arm | n | min | median | max | +|---|---|---|---|---| +| CUDA | 31 | 3.012 | **4.598** | 126.456 | +| CPU | 31 | 7.857 | **9.055** | 23.174 | + +The medians are the figures; the maxima are the first decode step after prefill, +when the slot cache is cold, and quoting either end of the range would be +quoting the least representative number in it. The implied ratio is 1.97x and it +is NOT a result: the correctness gate that would license it fails. Note also +that this CPU arm is FASTER than the 11.05 s/token previously recorded for the +CPU arm at 4000 slots, so the same-lease interleaved denominator here and that +earlier figure are not the same measurement and should not be mixed. + +**Owed from this run.** The CUDA arm's own top-2 margin: the scratch instrument +that reads `logits` in the completion callback SIGSEGVs on that arm +(`SCRIPT_EXIT=139`), almost certainly because the pointer is not host memory +there. And a ratified gate for a two-arm comparison whose greedy path is this +finely balanced, which `AGENTS.md` reserves as an explicit operator decision. diff --git a/.agents/issue-index.md b/.agents/issue-index.md index 6eb473143..db0809c6c 100644 --- a/.agents/issue-index.md +++ b/.agents/issue-index.md @@ -377,3 +377,5 @@ rather than merged. `scripts/check-agent-record.py` gates both. | [#1226](https://github.com/mudler/vllm.cpp/issues/1226) | `ENG-CUDAGRAPH-DEDUP` | The owed device A/B for #1162 RAN on `dgx:gpu0` 2026-08-18 (GB10, driver 580.173.02, nvcc 13.0.88, `rc` job f88d484b, gated commit `72de552c8`, whose four dedup sources are byte-identical to the merged `2a976eb9f`) and it SPLIT. CORRECTNESS PASSED and closes [#1184](https://github.com/mudler/vllm.cpp/issues/1184): 12/12 cells exit 0, zero `invalid device function` and zero `engine-fatal` where the pre-fix head `e4ce5571a` died after exactly one replay, ON replays as often as OFF (60=60, 33=33, 43=43), and `--output-token-ids` is IDENTICAL over 10/10 comparisons with the three OFF/OFF controls passing FIRST and three workloads hashing to three DIFFERENT values. THE BENEFIT IS REFUTED for the case the row was filed for: `N == M` in every `VT_CUDA_GRAPH_DEDUP=1` cell — 3 graphs to 3 execs on sizes [24 16 8], 2 to 2 on [16 8], 2 to 2 on [32 24] — with the registry count CLIMBING 1→1, 2→2, 3→3, so more than one capture reached it and the 1:1 is a measurement rather than the single-capture artefact of the first attempt. The cause is structural and was pre-registered before the run: `AppendKernelPayload` hashes (`func`, `gridDim.{x,y,z}`, `blockDim.{x,y,z}`, `sharedMemBytes`) at `src/vt/graph_dedup_runtime.h:121-128` and the memcpy payload hashes the copy extent, so the padded batch dimension is IN THE KEY, no candidate group forms and `cudaGraphExecUpdate` is never attempted. That contradicts the row's own premise, and SGLang keys the same fields (`cuda_graph_dedup_mixin.py:105-114`), so whatever folds upstream is not decode buckets either. NO throughput or memory number is recorded — clocks unpinned AND the ON arm allocated exactly as many executables as OFF. `VT_CUDA_GRAPH_DEDUP` therefore stays OFF and the default flip is NOT JUSTIFIED on this evidence. THE OPEN HYPOTHESIS this issue owns, deliberately not decided by the record that filed it: a COARSER key keeping the function addresses and topology but dropping the launch dimensions and memcpy extents would let two padded buckets form a candidate group at all, and the probe-before-fold design means a wrong grouping costs one wasted `cudaGraphExecUpdate` probe and a private executable rather than a wrong replay, so it is a cost question needing its own spec, a red-first discrimination test and a device probe-refusal rate. Also owed here: whether ANY other capture site or two models sharing the process-singleton registry can fold (only the Qwen3 dense decode driver was exercised), and a supporting `orin:gpu0` leg, BLOCKED because the Jetson 540.4.0 driver cannot run a CUDA 13 runtime (`cudaGetDeviceCount err=35`; a CUDA 12.x toolkit is the untried route). Honest gaps carried into the record: per-shape replay counts are unavailable, so B's ~30-per-shape is arithmetic over a printed TOTAL; the driver's "N captured size(s)" counts SLOTS not captures; the container's own cuBLASLt was never re-tested at CUDA 13.0. Full evidence [`benchmark-record.md`](benchmark-record.md) entry `ENG-CUDAGRAPH-DEDUP W4`; spec [`eng-cudagraph-dedup.md`](specs/eng-cudagraph-dedup.md) `## Outcome` | perf | | [#1229](https://github.com/mudler/vllm.cpp/issues/1229) | `ENG-CUDAGRAPH-DEDUP` | `origin/main` went RED at `fe24a3029` because `tests/scripts/test_check_gate_commands.py:670` pinned the literal string `Device byte-identity A/B (owed` inside `.agents/specs/eng-cudagraph-dedup.md`, and that commit legitimately rewrote the sentence from recording the device A/B as owed to recording that it RAN and what it found. The record edit is correct and the assertion described a state the record left, so every branch that merged `origin/main` inherited the failure. The protected property is stated by the assertion's own comment: the row is credited in `RUNNABLE_BASELINE` for `ctest -R test_graph_dedup` and `./scripts/agent-preflight.sh`, both CPU-tier, while gate item 6 of the spec contributes NO runnable command, so the credit covers the CPU tier and the record must not go silent about the leg it does not cover. That property is STATE-INDEPENDENT: `owed` was the leg's value on the day, not the property, and silence is the defect whether the leg was never run or the line was deleted after it ran. Asserting the NEW wording would only move the pin from `(owed` to `RAN 2026-08-18, PASS`, and the rerun this row already owes under [#1226](https://github.com/mudler/vllm.cpp/issues/1226) against a coarser signature key would red it again. FIXED by keying on structure: `gate_items`, `item_lead` and `gate_disposition` land in `scripts/check-gate-commands.py`, and the test now asserts that exactly one gate item names the device byte-identity A/B as its SUBJECT, that `runnable_commands` extracts nothing from it, and that its bold lead declares a disposition from a closed vocabulary. Lead-scoping is measured rather than assumed: a whole-item search credits item 5 with a disposition it does not declare, because its body contains `fail` in the phrase "proving the instrument can fail". Both historical wordings of item 6 are pinned as fixtures, which is what proves the new assertion is state-independent rather than re-pinned to today's sentence. The rule is NOT swept over every gate item and the reason is measured, not assumed: items 1 and 5 yield no command and declare no disposition, so a sweep is red on arrival, and widening the vocabulary until they pass would make it match ordinary English. That general rule is carried under `## Owed`. Also established and NOT a defect: `ERROR: these baseline rows left the gated population: ROW-THAT-IS-NOT-THERE` on the same run is the expected negative-path output of `test_check_mode_is_never_silently_swallowed_by_json`, confirmed by running that fixture alone. Spec [`fix-gate-commands-prose-pin.md`](specs/fix-gate-commands-prose-pin.md) | bug | | [#1240](https://github.com/mudler/vllm.cpp/issues/1240) | `MODEL-FP8-BLOCK-WEIGHT` | `test_qwen27_dense_forward` and `test_qwen27_dense_forward_glue_fuse_off` are RED on `origin/main` at `aba8d5ffb`. Both throw the block-wise-FP8 refusal `09597106e` (#1228, issue #1189) added at `src/vllm/model_executor/models/qwen3_5_dense_weights.cpp:428`: `'layer.linear_attn.in_proj_qkv.weight_scale_inv' is present ... but the checkpoint's quantization_config declares no weight_block_size`. ATTRIBUTED by measurement, not inference: `git log -1` on that file is `09597106e`, and `git show --stat 09597106e` changes it (+219) while touching neither test — the producer moved and its consumer did not. It hides behind the doctest thrown-case shape: the run reports 563 assertions with NONE failed beside 9 cases with ONE failed, because a case that THROWS records no failed assertion, so the assertions line reads clean and only `Status: FAILURE!` and the exit code tell the truth. NOT fixed in flow, deliberately: the refusal is correct in its own terms (guessing 128x128 is exactly what #1166 asked not to happen), so whether the synthetic fixture is wrong to emit `weight_scale_inv` without `weight_block_size`, or the refusal is too broad for a checkpoint that is not block-wise FP8 at all, is a semantics choice belonging to the owning row with its own red-first evidence — not a fixture edit that makes the message go away. Found by `ENG-EXPERT-STREAM-DEVICE` W0 running the full ctest suite after merging `origin/main` (523 of 525 passed; these two the only failures, and W0 touches neither file) | bug | +| [#1299](https://github.com/mudler/vllm.cpp/issues/1299) | `ENG-EXPERT-STREAM-DEVICE` | With W0's lane on, `Qwen3.8-2.4T-A95B UD-Q1_0` (369.97 GiB) LOADS on `--device cuda` on a 119.631 GiB GB10 — 61.20 GiB resident, ~265 s, the [#1123](https://github.com/mudler/vllm.cpp/issues/1123) refusal correctly silent — and then exhausts the machine inside its FIRST forward, **zero decode steps across seven attempts**. Diagnosed by the measurements rather than by reading: a **0.15 GiB** slot arena died exactly where an 18.55 GiB one did, so the arena is not the cost; a 1-token prompt (protected set 2,790 slices, no in-place fallback) behaved identically to a 5-token one, so prefill protection is not either; growth was ANONYMOUS (`RssAnon` 8.1 -> 61.4 GB) while file-backed stayed flat, refuting "the GPU pins the mapping"; and host anon + swap ~65 GB against system `used` ~119 GB leaves ~42 GB of device memory unified memory does not charge to RSS. So **the non-expert weights are resident twice** — once as the host `OwnedTensor`, once as the `ResidentWeight` device staging copy — and on a part where device memory IS host memory the second copy buys nothing. About 39 GiB of the 61.20 is `attn_qkv` + `ssm_out`, which the GDN V-head reorder makes `kTransformedWeight` and therefore expands to bf16 in OWNED host buffers. FIXED as W0f: where `Platform::host_memory_is_device_addressable()` — the same probed predicate W0c already serves every expert slice on — `ResidentWeight` returns a tensor over `w.bytes.data()` instead of allocating and copying. Safety is by ALIGNMENT rather than by enumerating kernels: `MakeHostBytesDeviceAliasable` guarantees the aliased pointer meets `kDeviceAliasAlignment = 256`, which is what `cudaMalloc` returns, so no consumer can distinguish the two pointers; an OWNED buffer below that (glibc gives 16) is re-homed once into an aligned block, and a BORROW below it declines and stages, because copying a clean file-backed mapping into anonymous memory would ADD the residency this change removes. A DISCRETE device answers the predicate false and is byte-identical to before, asserted by its own case. Red-first and mutation-proven in `tests/vllm/model_executor/test_resident_weight_host_addressable.cpp` | bug | +| [#1320](https://github.com/mudler/vllm.cpp/issues/1320) | `ENG-EXPERT-STREAM-DEVICE` | `VT_CPU_QUANT_REPACK` rewrites a Q8_0 weight into the `block_q8_0x4` i8mm interleave at load, and only the CPU `MatmulBTKernel` understands that layout — the CUDA quant dot reads plain `block_q8_0` and never consults `Tensor::repacked`, so a repacked weight uploaded to a device produces WRONG TOKENS rather than a crash. Its sibling `elem_kn_repack` has BOTH a CPU-platform gate in the loader policy (`gguf_keep_quant.cpp`, `CurrentPlatform().device_type() == kCPU`) and a refusal in `ResidentWeight`; `quant_repack` had NEITHER, because it rides `vt::cpu::QuantRepackActive()`, a HOST-CPU Arm i8mm probe that says nothing about where the weight will execute. An aarch64 box doing `--device cuda` — which is exactly `dgx:gpu0`, a Grace-Blackwell GB10 — satisfies it. Currently silent rather than firing, and that is measured: eligibility is Q8_0 with `n % 4 == 0 && k % 32 == 0`, `Qwen3.8-2.4T-A95B UD-Q1_0` has exactly one Q8_0 tensor at 0.01% of parameters, and the instrumented load of that checkpoint recorded `quant_repack = 0`. Found while implementing W0f ([#1299](https://github.com/mudler/vllm.cpp/issues/1299)), which changes the same function and whose G0-CORRECT gate compares CUDA tokens against CPU tokens — precisely the instrument that would have reported this as a W0f defect. FIXED IN FLOW as the tripwire its sibling already has, placed beside it so it covers the staging branch AND W0f's aliasing branch (where the bytes live was never the question; the kernel that reads them is), gated red-first on both arms and mutation-proven. `VT_CPU_QUANT_REPACK=0` is the operator's way past it. NOT fixed here, deliberately: adding the missing CPU-platform gate to `p.quant_repack` itself is a loader-policy semantics change owned by `QUANT-GGUF-KEEPQ-LOADER`, carried under `## Owed` in [`expert-stream-device-slots.md`](specs/expert-stream-device-slots.md) | bug | diff --git a/.agents/specs/expert-stream-device-slots.md b/.agents/specs/expert-stream-device-slots.md index a9df1b6c4..097a4057d 100644 --- a/.agents/specs/expert-stream-device-slots.md +++ b/.agents/specs/expert-stream-device-slots.md @@ -10,9 +10,58 @@ platform may read it. ## Now -`ACTIVE`. **W0b, W0c and W0d are implemented and unit-gated on the CPU tier. -W0a and W0e are QUEUED behind a four-hour lease on `dgx:gpu0` and neither has -run, so no GPU number exists yet and none is claimed.** +`ACTIVE`. **`--device cuda` DECODES this checkpoint on a GB10, and the +correctness gate that would let us publish a number does not pass.** W0e ran on +2026-08-19 inside one `rc hold` on `dgx:gpu0` at source `9c783a8be`. + +* **G0-LIVE: PASS.** 32/32 steps where seven previous attempts produced ZERO; + decode-phase `exhausted` delta **0** (6077 at step 1 and at step 32; the total + is the structural prefill number this spec predicted); `W0E_DOCKER_RC=0`, no + guard trip, peak RSS **97.75 GiB** with swap untouched. +* **G0-CORRECT: FAIL as declared, and MEASURED to be a near-tie rather than a + disagreement about the model.** The 32 ids match the CPU arm for six tokens and + diverge at the seventh — `...,264,3177,7172,...` on CPU against + `...,264,3177,303,...` on CUDA. Both continuations are coherent. Three things + were then established rather than assumed: + 1. **It is not W0f.** The CPU arm was re-run on the SAME binary and the SAME + lease and reproduced the recorded answer byte for byte. W0f cannot reach + that arm at all — `ResidentWeight` returns at the `is_cpu()` branch above + it, and the instrument counted `w0f-alias` calls **0** on the CPU arm, + which is a live control and not an argument. + 2. **The two arms rank the same two candidates.** An instrumented CPU run + printing the top-2 logits per step shows that at the divergent step + (`lp_call=7`) the CPU arm's top-1 is `7172` at 18.779411 and its **top-2 is + `303` at 18.514702** — `303` being exactly the token CUDA emitted. The + **margin is 0.264709 logits**, 1.4 % of the winning logit. + 3. **This decode is full of ties that narrow.** `lp_call=9` has a margin of + **0.022802**, about 0.1 %. A greedy path this finely balanced flips on any + arithmetic difference, and the two arms run genuinely different GEMM + kernels. + So the declared gate fails and the wave stops, which is correct. What the + failure means is a different question, and it is now answered with numbers: + the arms agree about the distribution and disagree about a coin flip. + Whether a token-exact cross-arm gate is the right instrument for a path with + no oracle is a decision for the operator, not something this row may assume. +* **G0-SPEED: VOID, by this row's own stop condition.** It was measured over + the 31 DECODE steps of each arm (step 1 is prefill and is excluded), + interleaved on one lease: CUDA median **4.598 s/token** (min 3.012, max + 126.456), CPU median **9.055 s/token** (min 7.857, max 23.174). Both maxima + are the first decode step, with the slot cache cold. It is NOT claimed, + because a speed number behind a failing correctness gate is exactly the shape + #912 F1 was. Two cautions if anyone is tempted to quote it anyway: the implied + 1.97x rests on a token comparison that FAILED, and this CPU arm is faster than + the 11.05 s/token previously recorded at 4000 slots, so the two are not the + same measurement and must not be mixed. + +**What W0f did, measured rather than inferred.** The instrument added for this +run counts **60.793 GiB** of dense weight aliased instead of duplicated into +device memory, against **~9.2 GiB** that declined (misaligned GGUF borrows) and +still stages. That is the whole difference between zero decode steps and 32. + +W0a remains unrun as a standalone probe, and it no longer blocks anything: the +load succeeds only when `host_memory_is_device_addressable()` answers true (W0d's +conditional refusal is keyed on it), so a completed run is itself the +measurement W0a was going to take. What that means precisely, because "W0 landed" would overstate it: @@ -39,6 +88,17 @@ What that means precisely, because "W0 landed" would overstate it: `test_gguf_device_fit_reach` (8 cases / 36) for the production reach. **The lane-off bound is byte-identical**, asserted three ways against literal values. +* **W0f — the dense half, and the reason W0 still produced no token.** With + W0b-W0d in the tree the checkpoint LOADS on `--device cuda` and then exhausts + the box inside the first forward, zero decode steps over seven attempts + (issue #1299). The lane was doing its job; the DENSE weights were resident + twice, once as the host `OwnedTensor` and once as `ResidentWeight`'s device + staging copy, and on a part where device memory IS host memory that doubling + is what runs it out. W0f gives `ResidentWeight` the same branch W0c gave + `KqExpertSlice`, on the same probed predicate. Gated in the new + `test_resident_weight_host_addressable` (9 cases / 45 assertions) over the same + fake staging, host-addressable platform, plus one defect the work uncovered + and fixed in flow (#1320). * **W0a — the probe.** Submitted to `rc` FIRST, before any tree change, and still queued. It is the experiment that decides whether the W0b design is sound at all: if `cudaDevAttrPageableMemoryAccess` is 0 on GB10, the CUDA leg @@ -220,6 +280,7 @@ Nothing is ported; there is no upstream. This is the local change map. | W0b | `include/vllm/platforms/interface.h`, `src/vllm/platforms/cuda.cpp`, `src/vllm/platforms/rocm.cpp` | new `virtual bool host_memory_is_device_addressable() const { return false; }` beside `is_integrated_gpu`; CUDA overrides from a probe taken once at registration next to the existing `cudaDevAttrIntegrated` probe; ROCm overrides from the `pageable_memory_access` capability it ALREADY probes (`rocm_backend.hip:96-103`) | | W0c | `src/vllm/model_executor/models/qwen3_5.cpp` | `KqExpertSlice` takes the slot arm under `is_cpu()` OR `host_memory_is_device_addressable()`; the slot branch builds its tensor without `ResidentWeight`; a named `VT_CHECK` in `ResidentWeight` refuses a streamed `*_exps` tower reaching device staging | | W0d | `include/vllm/model_executor/model_loader/gguf_device_fit.h`, `src/.../gguf_device_fit.cpp`, `src/vllm/entrypoints/model_loader.cpp` | the fit bound gains an explicit "these tensors are served by the slot lane, and the arena costs this instead" input; the loader passes it when the resolved config says streaming is on and the platform can read host slots | +| W0f | `src/vllm/model_executor/models/qwen3_5.cpp`, `include/vllm/model_executor/models/qwen3_5_weights.h`, `src/vllm/model_executor/models/qwen3_5_weights.cpp` | `ResidentWeight` returns a tensor over `w.bytes.data()` where `host_memory_is_device_addressable()`, instead of `Alloc` + `Copy` into `w.d_dev`; `MakeHostBytesDeviceAliasable` + `kDeviceAliasAlignment` make that pointer indistinguishable from the `cudaMalloc` one it replaces; a named `VT_CHECK` refuses an i8mm-repacked weight reaching device residency on EITHER branch (#1320) | | W1 | `include/vllm/model_executor/expert_streamer.h`, new `include/vllm/model_executor/device_expert_slot_store.h` | `CommitSlot(int32_t, size_t)` on `ExpertSlotStore` (no-op on the host store); `DeviceExpertSlotStore` allocating slots through `vt::Backend::Alloc` with one pinned host staging slot, `SlotForWrite` returning staging and `CommitSlot` doing the H2D; correct the two false sentences in `expert_streamer.h` | | W2 | `expert_streamer.h`, `host_expert_slot_store.h`, `device_expert_slot_store.h`, `qwen3_5.cpp` | `virtual uint8_t* SlotForRead(int32_t)`; `Qwen35ExpertStream::store_` becomes `std::unique_ptr`; `:5381` and `:5437` read through the virtual; the store is selected from the platform | @@ -241,6 +302,7 @@ mutation: | `tests/vllm/model_executor/test_gguf_device_fit.cpp` (extend) | with the lane on, the bound excludes `*_exps` and adds the arena; with it off, the bound is byte-identical to today | make the exclusion unconditional | | `tests/vllm/entrypoints/test_gguf_device_fit_reach.cpp` (extend) | the loader reaches the conditional refusal from the production entry point | delete the production call site | | a `qwen3_5` slot-arm unit gate | the slot branch never calls `ResidentWeight`, and a streamed tower reaching device staging throws by name | remove the `VT_CHECK`; restore the `ResidentWeight` call | +| `tests/vllm/model_executor/test_resident_weight_host_addressable.cpp` (new, W0f) | `ResidentWeight` aliases the host bytes on a host-addressable staging platform and allocates NOTHING; the aliased pointer meets `kDeviceAliasAlignment`; a discrete platform stages byte-identically to today; a MISALIGNED BORROW declines and stages rather than being copied into anonymous memory; the three refusals fire on the aliasing branch too | delete the aliasing branch; make the predicate unconditional; delete each `VT_CHECK`; drop the `borrowed()` guard; claim alignment without providing it; re-home without copying the bytes | ## Gates @@ -346,6 +408,33 @@ where the wave ENDS, not where it degrades quietly into the next one. taking a general per-tensor staging POLICY (the shape #1136 explicitly refuses to invent), stop and return `NEEDS_DECISION`. The lane's tensor set is `*_exps` and is knowable; a general policy input is not. +* **W0f — the dense half.** Discovered by W0e's first seven attempts and scoped + by them, not by reading: the checkpoint loads and then exhausts the box with + zero decode steps, and the four measurements in #1299 rule out the arena, the + prefill fallback, and a pinned mapping in turn. Give `ResidentWeight` the same + branch W0c gave `KqExpertSlice`. + **Why an alignment contract and not a kernel survey.** The staging branch is a + verbatim byte copy, so the ONLY thing a consumer can notice about the + substitution is the pointer's alignment. `cudaMalloc` returns 256; a + `std::vector` returns 16, because a large glibc block is an mmap chunk + landing at page+16. Matching the allocator therefore settles every consumer at + once, and the alternative — deriving a floor from the widest load any kernel + performs — does not close: the widest hand-written one is a 16-byte `cp.async` + granule whose gate checks the SHAPE and assumes the base, and cuBLASLt is + separately PROMISED 256 by a preference default this tree never sets. + **Why a borrow is not re-homed.** It owns no anonymous pages. Copying a clean, + file-backed GGUF mapping into an aligned anonymous block would create exactly + the residency this row exists to remove, and would break a tied + `token_embd`/`lm_head` pair's single keep-alive. + **Gate:** `test_resident_weight_host_addressable`, mutation-proven. + **Stop condition:** if any weight on this path needed a device layout DIFFERENT + from its host bytes, that weight could not skip the copy and W0f would need a + per-tensor answer instead of a branch. It does not: `ResidentWeight` copies + bytes verbatim and returns the same dtype, shape and (dropped) marker set on + both arms, so there is no device layout to preserve. The layout-bearing + markers are handled instead — `elem_kn_repacked` and `repacked` are refused by + name, and `q8_0_aligned` is a load-time rewrite of the HOST bytes that no + Qwen3.5 path sets. * **W0e — the measurement.** G0-CORRECT, G0-LIVE, G0-SPEED, on one lease. **Stop condition:** a token mismatch, `steps == 0`, or a non-zero decode-phase `exhausted` delta stops the wave and voids the number. @@ -417,5 +506,9 @@ re-derived here. | **G-DISCRETE: validate W1/W2 on a discrete NVIDIA GPU.** The measurement: on a device with VRAM V and `host_memory_is_device_addressable() == false`, load a GGUF whose `*_exps` towers exceed V, with the lane on, and gate (i) token-exactness against the CPU arm on the same checkpoint, (ii) decode-phase `exhausted` delta 0, (iii) peak device allocation <= non-expert remainder + arena. | No discrete NVIDIA GPU is reachable from this project. `dgx:gpu0` is a GB10 where device memory IS host memory, so a device store there exercises the plumbing and not the thing W1 exists for. Recorded rather than implied, because a gate nobody can run is not a gate. | | **A zero-copy device filler (GPUDirect Storage / `cuFile`).** | W1 ships the staging bounce by choice, for the reasons in its design note. The measurement that would justify replacing it — a device-arm decode where the H2D leg is a measurable fraction of fill time — does not exist until W1 has run somewhere. | | **The CPU arm's streaming decode figure is still VOID.** `docs/BENCHMARKS.md:8` records it as VOID (#912 F1) with a re-measure owed. | Owned by `ENG-EXPERT-STREAM` and arranged separately by the operator. It is the DENOMINATOR for G0-SPEED, not a precondition for G0-CORRECT or G0-LIVE. | +| **A ratified gate for a two-arm comparison whose greedy path is a coin flip.** The measurement that would settle it: over N prompts, the distribution of top-2 margins at each step, and the fraction of steps whose margin is below the arms' measured arithmetic spread. | W0e MEASURED the margin at the divergent step (0.264709 logits, 1.4 %) and one step later (0.022802, 0.1 %), so the token-exact gate is failing on ties rather than on a defect. Ratifying a distributional gate is exactly the decision `AGENTS.md` reserves for an explicit act — "use an explicitly ratified distributional gate only when the oracle's greedy decode is non-deterministic" — and it is the operator's, not this row's. Until it is taken, G0-CORRECT stays FAILING and G0-SPEED stays VOID, which is the conservative reading and the one that cannot publish a wrong number. | +| **The CUDA arm's own top-2 margin at the divergent step.** | The scratch instrument that reads `logits` in the completion callback SIGSEGVs on the CUDA arm (`SCRIPT_EXIT=139`), almost certainly because the pointer it is handed there is not host memory on that arm. The CPU arm's margin is enough to establish the near-tie — `303` is its own runner-up — but the symmetric number is not in hand and is cheap to take once the callback's pointer residency is known. | +| **The family-wide copy of this change: `include/vllm/model_executor/models/dense_attn_block.h`'s `ResidentWeight` still stages unconditionally.** The measurement: on a host-addressable staging platform, load any of the ~50 models that include that header and show peak resident bytes falling by the model's weight size, with tokens unchanged. | W0f deliberately changes only `qwen3_5.cpp`'s PRIVATE copy, which is the one that governs `Qwen3.8-2.4T-A95B UD-Q1_0` (that file kept its own helper; the header's copy is not on the Qwen3.5 path). The header's version is reached from `ModelRegistry::Forward` for every model that includes it, so extending it is not dead code — but nothing on a CPU tier can drive one of those forwards on a staging platform, so the extension would land with its reachability argued rather than gated, across ~50 architectures at once. That is a scope and a review question, not a line of code, and it gets its own row. | +| **The missing CPU-platform gate on `p.quant_repack` itself ([#1320](https://github.com/mudler/vllm.cpp/issues/1320)).** The measurement: `elem_kn_repack` is resolved with `CurrentPlatform().device_type() == kCPU` and `quant_repack` is not, so a device load can still perform a CPU-only transform and be caught afterwards instead of never doing it. | W0f fixes the CONSEQUENCE in flow — a named refusal on both arms of `ResidentWeight`, red-first and mutation-proven — because that is the small and clear part. Moving the gate into the loader policy changes what a GGUF load DOES on a device rather than what it refuses, which is `QUANT-GGUF-KEEPQ-LOADER`'s semantics and needs its own red-first evidence. | | **`.agents/specs/expert-streaming.md`'s `## Owed` entry for #1124 still names no owning row ID.** | Not edited here on purpose; PRs #1200 and #1216 both edit that file. One-line follow-up once both land. | | **W1 may land UNREACHED if it is split from W2.** | The recommendation is one pull request. If a split is chosen, the commit body and the PR body must name what is unreached and name W2 as the owning wiring, per `## Nothing lands dead`. | diff --git a/docs/BENCHMARKS.md b/docs/BENCHMARKS.md index f4f2307df..ec051c436 100644 --- a/docs/BENCHMARKS.md +++ b/docs/BENCHMARKS.md @@ -6,7 +6,7 @@ |---|---|---|---| | **Weight load (`ENG-LOAD-DIRECT-UPLOAD`, #150)** | Qwen3.6-27B bf16 (50.098 GiB), GB10 Vulkan, same binary both arms, interleaved under one GPU lock | **Load phase 1.54x warm** (19.27 -> 12.48 s), **1.61x cold** (52.62 -> 32.75 s); bytes moved **100.196 -> 81.260 GiB**. Every ON leg beat every OFF leg | byte-identical; 6/6 token-exact | | **Expert streaming vehicle (`ENG-EXPERT-STREAM`, [#912](https://github.com/mudler/vllm.cpp/issues/912))** | Qwen3.8-2.4T-A95B `UD-Q1_0` (370 GiB) on ONE GB10, 119 GiB, c1 | **Loads and generates** (streaming OFF): resident **62 GiB**, decode **66.7 s/tok**. Experts BORROW the mmap; ~6.7 GB/token as 4 KiB faults, **~50x** off NVMe. Streaming-ON decode **VOID** (#912 F1); re-measure owed | correct answer; no oracle runs this | -| **Device expert slices (`ENG-EXPERT-STREAM-DEVICE`, [#1124](https://github.com/mudler/vllm.cpp/issues/1124))** | Qwen3.8-2.4T-A95B `UD-Q1_0`, `--device cuda` vs `cpu`, one GB10, greedy, 32 tokens | **PENDING; nothing measured.** W0b-W0d gated on CPU; W0a and W0e queued on `dgx:gpu0`. No speed floor: a CUDA arm slower than CPU is a real result. [spec](../.agents/specs/expert-stream-device-slots.md) | correct answer; no oracle runs this | +| **Device expert slices (`ENG-EXPERT-STREAM-DEVICE`, [#1299](https://github.com/mudler/vllm.cpp/issues/1299))** | Qwen3.8-2.4T-A95B `UD-Q1_0` on `cuda` vs `cpu`, one GB10, greedy, 32 tokens | **IT RUNS: G0-LIVE PASS, G0-CORRECT FAIL, G0-SPEED VOID.** 32/32 steps after seven zero-step tries, `exhausted` delta 0, peak RSS 97.75 GiB, step-7 divergence is a near-tie. [record](../.agents/benchmark-record.md) | correct; no oracle | | **Structured state record (active)** | v1 scalar + relational + Git-history contracts | No benchmark. At `776c56f1`: 157 imports = 3,231,342 exact bytes; append preserved all 156 wrappers/rows. 95 tests: validator/core 44 (checker 20 + core 24), NOW 18, migration 22, cutover 11. New raw-row mutation guard. | n/a | | **Binary release (ACTIVE; Windows pre-alpha pending)** | v0.0.2 shipped eight primary archive/checksum/provenance triplets + two indexes (26 assets) from source SHA `7020de93652ca920424a10ac5255b34810dd2f24`, run `31466516224` | Windows W14-W16 implemented. **PENDING:** native hosted gates, merged-SHA ten-tuple dry run, matching-hardware evidence, v0.0.3-pre.1 publication, 32-asset audit | W12 optional/non-primary | | **Container images (ACTIVE; arm64 cuda verified on GB10 + Orin 2026-08-11)** | `ENG-RELEASE-CONTAINERS` ([spec](../.agents/specs/container-images.md)) | cpu amd64 783 MB; cuda arm64 **1.71 GB**. GB10 `sm_121a`: `/health`+`/version`+SIGTERM on `--gpus all`. Orin `sm_87` (Tegra): Qwen3-0.6B **generates**, GPU **GR3D 95-97%** | n/a | diff --git a/include/vllm/model_executor/models/qwen3_5_weights.h b/include/vllm/model_executor/models/qwen3_5_weights.h index c3fd66c42..a3401f588 100644 --- a/include/vllm/model_executor/models/qwen3_5_weights.h +++ b/include/vllm/model_executor/models/qwen3_5_weights.h @@ -20,6 +20,7 @@ // toggles retain the split residents. #pragma once +#include // size_t, for kDeviceAliasAlignment #include #include #include @@ -220,6 +221,140 @@ struct OwnedTensor { // behavior (house convention for a default-on residency change). void AdoptDeviceBytesAsHost(vt::Backend& backend, const OwnedTensor& w); +// The alignment a HOST pointer must meet before a device kernel may be handed it +// in place of the `Backend::Alloc` pointer it would otherwise have received. +// +// 256, because that is cuBLASLt's documented +// `CUBLASLT_MATMUL_PREF_MIN_ALIGNMENT_A_BYTES` DEFAULT, which this tree never +// sets, and it dominates every explicit pointer gate in the tree. Measured, not +// assumed: `grep -rn MIN_ALIGNMENT src/vt/` finds nothing, so the 256 default +// applies to every cuBLASLt matmul this tree issues; and the only genuine +// POINTER-alignment gate in the CUDA kernels is +// `src/vt/cuda/cuda_matmul_nvfp4.cu`'s `reinterpret_cast(prow) & 0xf`, +// which asks for 16. The other `% 32` and `% 64` tests nearby read as alignment +// gates and are not: they check a DIMENSION (`d`, `dv`), not an address. It is also what `cudaMalloc` +// returns in practice, though CUDA guarantees only "suitably aligned" and +// current devices return more — so "indistinguishable from a `cudaMalloc` +// pointer" is the intuition, and "at least what every consumer is promised" is +// the claim. +// +// WHAT ALIGNMENT DOES AND DOES NOT BUY. It makes the substitution CORRECT: no +// kernel can fault or mis-vectorise on this pointer that would not have on the +// other. It does not make the two pointers indistinguishable in every respect, +// and two in-tree facts say so. `src/vllm/model_executor/models/laguna.cpp` +// records a MEASURED GB10 penalty for reading system-allocated memory from the +// GPU rather than a `cudaMalloc` allocation, worst on a long-K low-parallelism +// GEMV — a consumer telling them apart by BANDWIDTH, which is why +// `VT_QWEN35_ALIAS_HOST_WEIGHTS` exists below. And the Vulkan and Metal backends +// resolve a tensor pointer against their own allocation tables and throw if it +// is outside them, telling them apart by IDENTITY; harmless only because neither +// overrides `host_memory_is_device_addressable()`, so this argument is scoped to +// backends that take raw pointers. Deriving a smaller number would mean +// enumerating every kernel that ever binds a weight and being right about all of +// them, and the enumeration does not close: the widest thing any of them +// dereferences is a 16-byte `cp.async` granule +// (`src/vt/cuda/cuda_matmul_nvfp4.cu`, whose shape gate assumes an aligned base +// rather than checking it), but cuBLASLt is PROMISED 256 — +// `CUBLASLT_MATMUL_PREF_MIN_ALIGNMENT_A_BYTES` defaults to 256 and this tree +// never sets it (`src/vt/cuda/cuda_matmul.cu`), so the heuristic may pick an +// algorithm on the strength of a promise a 16-aligned pointer breaks. Matching +// the allocator instead of the consumers makes the whole question go away, and +// it costs one memcpy that REPLACES the host->device copy it removes. +inline constexpr size_t kDeviceAliasAlignment = 256; + +// Make `w.bytes` safe to hand to a device kernel directly, and say whether it +// worked. On return `true`, `w.bytes.data()` is non-null and aligned to +// `kDeviceAliasAlignment`. On `false` the caller must fall back to staging, and +// nothing has changed. +// +// THREE CASES, and the middle one is the point (ENG-EXPERT-STREAM-DEVICE W0f, +// issue #1299). +// +// * ALREADY ALIGNED — true, and nothing is copied. A GGUF mmap borrow lands +// here whenever its tensor offset happens to be a multiple of 256; GGUF's +// `general.alignment` guarantees only 32, so this is luck rather than a +// contract, and the fallback below is what makes that acceptable. +// * OWNED AND MISALIGNED — the bytes are moved into a `kDeviceAliasAlignment` +// allocation and `w.bytes` is re-pointed at it, keeping the new block alive +// the way `AdoptDeviceBytesAsHost` keeps the device block alive. A plain +// `std::vector` from glibc is 16-byte aligned and no more (a large +// block is an mmap chunk, so it lands at page+16), which is exactly what the +// GDN V-head reorder's ~44.6 GiB of bf16-expanded `attn_qkv` / `ssm_out` +// arrive as. Without this they could never be aliased and W0f would move no +// bytes at all. +// * BORROWED AND MISALIGNED — false. A borrow owns no anonymous pages: it is a +// clean, file-backed GGUF mapping or a tied pair's single shared expansion. +// Copying it would CREATE the anonymous residency this change exists to +// remove, and would break the tie. Staging is the right answer for it. +// +// Logically const, like the lazy device residency beside it: only where the +// bytes live changes, never what they are. +// The outcomes, so a caller and a log can say WHICH one happened. +enum class HostAliasOutcome { + kAliasedInPlace, // already aligned; nothing allocated and nothing copied + kRehomed, // an OWNED misaligned buffer moved into an aligned block + kDeclinedBorrow, // a misaligned BORROW; the caller must stage + kDeclinedEmpty, // no host bytes at all + kDeclinedDisabled, // VT_QWEN35_ALIAS_HOST_WEIGHTS=0 +}; + +bool MakeHostBytesDeviceAliasable(const OwnedTensor& w, + HostAliasOutcome* outcome = nullptr); + +// Bytes seen by `MakeHostBytesDeviceAliasable`, split by outcome, since process +// start. +// +// WHY A COUNTER AND NOT AN INFERENCE FROM RSS. W0f's first device attempt was +// read only through `free -m`, and what it showed — about 47 GB appearing in +// 30 seconds at the first forward — is equally consistent with "the branch +// declined and staged as before", with "the branch re-homed and the old pages +// did not come back", and with "something else allocated". Those three call for +// three different changes, and no amount of staring at an RSS curve chooses +// between them. This says how many bytes took each outcome. It is printed +// PERIODICALLY rather than at exit, because the process it measures is one the +// memory guard kills before any exit handler runs. +struct HostAliasStats { + uint64_t aliased_in_place_bytes = 0; + uint64_t rehomed_bytes = 0; + uint64_t declined_borrow_bytes = 0; + uint64_t declined_other_bytes = 0; + uint64_t calls = 0; +}; +HostAliasStats HostAliasSnapshot(); + +// The same-binary A/B back to the staging behaviour, per the house convention +// for a default-on residency change that `VT_ADOPT_DEVICE_BYTES` and +// `VT_MOE_HOST_FREE` already follow. `VT_QWEN35_ALIAS_HOST_WEIGHTS=0` makes +// every call decline, so one build can measure both arms — which matters more +// here than usual, because `src/vllm/model_executor/models/laguna.cpp` records +// a MEASURED GB10 penalty for reading system-allocated memory from the GPU +// rather than a `cudaMalloc` allocation, worst on a long-K low-parallelism +// GEMV. This branch installs exactly that retag by default, and without a knob +// W0e could not tell a decode regression from the workload. +bool HostWeightAliasEnabled(); + +// May the host mirror of `w` be released, because a DEVICE copy exists to be +// authoritative in its place? +// +// THE INVARIANT A USE-AFTER-FREE TAUGHT US (issue #1299). `MoeBlockBf16Cuda` +// captures `ResidentWeight(...).data` for every expert into a device-resident +// pointer table, uploads the table once, and then releases the host mirrors. It +// justified that with "once the device copy exists it is authoritative and +// nothing reads the host bytes again", which was true while `ResidentWeight` +// had two behaviours. It has three: on a host-addressable platform it ALIASES, +// so the captured pointers ARE `w.bytes.data()` and releasing them frees memory +// the resident table still points at, for the model's lifetime and from inside +// captured graphs. A fresh review demonstrated it with a scratch case that takes +// SIGSEGV. +// +// The question is therefore not "did we upload" but "is there something else to +// read", and `d_dev` already answers it: null on exactly the arm that aliases, +// non-null on every arm that staged. Named rather than inlined so the release +// sites state the invariant they depend on, and so a gate can mutate it. +inline bool HostMirrorIsRedundant(const OwnedTensor& w) { + return w.d_dev != nullptr; +} + // Lazily-built per-weight DEVICE-RESIDENT state, OWNED BY THE WEIGHT (issue // #237). // diff --git a/src/vllm/model_executor/models/qwen3_5.cpp b/src/vllm/model_executor/models/qwen3_5.cpp index 2a40b4806..f52f5b478 100644 --- a/src/vllm/model_executor/models/qwen3_5.cpp +++ b/src/vllm/model_executor/models/qwen3_5.cpp @@ -1044,6 +1044,39 @@ std::vector WeightF32(const OwnedTensor& w) { // model's lifetime. On CPU the bytes are already host-resident, so a direct view // avoids the copy. The weight is a read-only matmul-B / norm / embed operand, so // the const_cast is safe. `shape` defaults to the owned shape. +// Print what the W0f aliasing branch has actually done, every 4 GiB of weight it +// has seen, on the same `VT_LOAD_STATS` switch the loader's byte counters use. +// +// WHY PERIODIC AND NOT AT EXIT. The `[vt load] bytes@exit` line is registered +// with `std::atexit`, and the run this instruments is one a memory guard +// SIGKILLs — no exit handler runs, so the one number that would have explained +// the run is the one number the run cannot print. W0f's first device attempt was +// read from an RSS curve for exactly that reason, and an RSS curve cannot tell +// "declined and staged" from "re-homed and the pages did not come back". +void ReportHostAliasResidency() { + static const bool on = [] { + const char* e = std::getenv("VT_LOAD_STATS"); + return e != nullptr && e[0] != '0'; + }(); + if (!on) return; + const vllm::HostAliasStats s = vllm::HostAliasSnapshot(); + const uint64_t total = s.aliased_in_place_bytes + s.rehomed_bytes + + s.declined_borrow_bytes + s.declined_other_bytes; + static uint64_t last = 0; + constexpr uint64_t kStep = 4ULL << 30; + if (total < last + kStep && last != 0) return; + last = total; + const double gib = 1024.0 * 1024.0 * 1024.0; + std::fprintf(stderr, + "[vt load] w0f-alias calls=%llu aliased_in_place=%.3f GiB " + "rehomed=%.3f GiB declined_borrow=%.3f GiB declined_other=%.3f GiB\n", + static_cast(s.calls), + static_cast(s.aliased_in_place_bytes) / gib, + static_cast(s.rehomed_bytes) / gib, + static_cast(s.declined_borrow_bytes) / gib, + static_cast(s.declined_other_bytes) / gib); +} + Tensor ResidentWeight(Dev d, const OwnedTensor& w, std::vector shape = {}) { if (shape.empty()) shape.assign(w.shape, w.shape + w.rank); // HOST-POINTER ALIASING IS A CPU PROPERTY, NOT A "NOT-CUDA" PROPERTY (issue @@ -1110,6 +1143,76 @@ Tensor ResidentWeight(Dev d, const OwnedTensor& w, std::vector shape = "expert-stream lane serves its slices from host slot storage and the " "whole tower must never be uploaded (ENG-EXPERT-STREAM-DEVICE W0c, " "issues #1123 and #1124)"); + // The SAME invariant as the elem_kn_repacked guard above, for the i8mm + // interleave, and it was missing until now (issue #1320). The CUDA + // quant dot reads `block_q8_0`; `VT_CPU_QUANT_REPACK` rewrites the buffer to + // `block_q8_0x4` at load and only the CPU MatmulBTKernel understands that. + // Unlike `elem_kn_repack`, whose policy IS gated on the CPU platform + // (gguf_keep_quant.cpp), `quant_repack` rides `QuantRepackActive()` alone — + // a HOST-CPU i8mm probe — so an aarch64 box doing `--device cuda` can repack a + // Q8_0 weight and then upload it verbatim to a kernel that misreads it. That is + // silent wrong tokens, not a crash. Measured harmless on the target checkpoint + // (one Q8_0 tensor, 0.01% of parameters, and the instrumented load recorded + // `quant_repack = 0`), which is why it is a tripwire here rather than a + // campaign; `VT_CPU_QUANT_REPACK=0` is the operator's way past it. + VT_CHECK(!w.repacked, + "qwen3_5: an i8mm-repacked (block_q8_0x4) weight reached device " + "residency; VT_CPU_QUANT_REPACK is a CPU-only load transform and the " + "device quant kernels read plain block_q8_0"); + // ENG-EXPERT-STREAM-DEVICE W0f (issue #1299). THE SECOND COPY THIS ROW EXISTS + // TO PREVENT, at the one line that makes it. + // + // Everything below this branch is a VERBATIM byte copy: `Alloc(w.bytes.size())`, + // `Copy`, then a tensor with the same dtype, the same shape and the same + // (dropped) marker set as the source. Nothing about the bytes changes, which is + // exactly why a token gate cannot see the cost — and the cost is a second full + // resident copy of every dense weight. On a discrete device that copy is the + // whole point: the kernel cannot follow a host pointer. On a part whose kernels + // CAN, it buys nothing and comes out of the same RAM the first copy did. + // + // MEASURED (#1299, `dgx:gpu0`, seven runs). `Qwen3.8-2.4T-A95B UD-Q1_0` loads + // on `--device cuda` at 61.20 GiB resident and then exhausts a 119.631 GiB box + // inside the FIRST forward, zero decode steps, every time. A 0.15 GiB slot + // arena died exactly where an 18.55 GiB one did, so the arena is not the cost; + // growth was anonymous while file-backed stayed flat, so the mapping is not + // pinned. About 39 GiB of that 61.20 is `attn_qkv` (21.56) and `ssm_out` + // (17.25), which the GDN V-head reorder makes `kTransformedWeight` and + // therefore expands to bf16 in OWNED host buffers — the split is measured in + // `.agents/specs/expert-streaming.md`, not derived here. The CPU arm pays that + // once and serves. This branch is what stops the CUDA arm paying it twice. + // + // WHY THE SAME PREDICATE AS W0c AND NOT A NEW ONE. `KqExpertSlice` already + // hands this platform a host pointer for every expert slice it serves; a dense + // weight is the same question about a different tensor. `is_cpu()` is what the + // early return above answers, `needs_weight_staging()` is true on CUDA + // everywhere and would gate nothing, and `is_unified_memory()` answers the + // opposite question — GB10 reports unified while a `cudaMalloc` pointer is + // still not host-dereferenceable (vt/backend.h). A DISCRETE device answers + // false here, falls through, and gets byte-for-byte what it gets today. + if (vllm::platforms::GetPlatform(d.q.device.type) + .host_memory_is_device_addressable()) { + // A weight whose host bytes are gone cannot be aliased, and the staging + // branch would not notice: it would `Alloc(0)`, copy nothing, and hand out a + // pointer to nothing. `ReleaseHost()` is the operation that takes them away + // and it is not reachable for the dense weights this branch serves, so this + // states the precondition rather than handling a case that exists. + VT_CHECK(!w.bytes.empty(), + "qwen3_5: a weight reaching device residency has no host bytes; its " + "host mirror was released and there is nothing to alias or upload"); + const bool aliased = MakeHostBytesDeviceAliasable(w); + ReportHostAliasResidency(); + if (aliased) { + // NOT `load_stats::AddDeviceUpload`: nothing was uploaded. Issue #150's + // counter measures bytes moved host->device, and this branch moves none. + return MakeTensor(const_cast(w.bytes.data()), w.dtype, d.q.device, + shape); + } + // A MISALIGNED BORROW, or the `VT_QWEN35_ALIAS_HOST_WEIGHTS=0` A/B, reaches + // here. A borrow's pages are clean and file-backed, so staging copies them + // without adding anonymous residency. Falling through is deliberate and is + // not a failure; `ReportHostAliasResidency` above says how often it happens + // and for how many bytes. + } if (!w.d_dev) { const size_t nb = w.bytes.size(); void* p = d.b.Alloc(nb); @@ -6700,11 +6803,32 @@ DBuf MoeBlockBf16Cuda(Dev d, const MoeBlockWeights& w, const HfConfig& cfg, /*committed_compute_path=*/MoeBf16FastEnabled(), /*host_free_env=*/host_free_on)) { d.b.Synchronize(d.q); // all E x 3 H2D uploads complete before any free + // W0f (#1299) FALSIFIED THIS BLOCK'S PREMISE, AND THIS IS THE REPAIR. + // + // The paragraph above says "once the device copy exists it is + // authoritative and nothing reads the host bytes again", and it was true + // while `ResidentWeight` had exactly two behaviours. It has three now: on + // a platform whose kernels can dereference host storage the function + // ALIASES, `d_dev` is never populated, and the pointers captured into + // `gp/up/dp` above ARE `w.bytes.data()`. Releasing the host mirror then + // frees the memory the resident device pointer table points at, and the + // grouped GEMM keeps reading it for the model's lifetime — including from + // inside a captured graph. A fresh review caught it with a scratch case + // that replays this exact sequence and takes SIGSEGV. + // + // The condition is therefore not "did we upload" but "IS THERE A DEVICE + // COPY TO BE AUTHORITATIVE", asked per weight, which is what `d_dev` + // already answers. It is `nullptr` on precisely the arm that aliases, and + // non-null on every arm that staged, so the discrete behaviour this + // paragraph was written for is unchanged. for (int64_t e = 0; e < E; ++e) { const size_t se = static_cast(e); - w.expert_gate[se].ReleaseHost(); - w.expert_up[se].ReleaseHost(); - w.expert_down[se].ReleaseHost(); + if (vllm::HostMirrorIsRedundant(w.expert_gate[se])) + w.expert_gate[se].ReleaseHost(); + if (vllm::HostMirrorIsRedundant(w.expert_up[se])) + w.expert_up[se].ReleaseHost(); + if (vllm::HostMirrorIsRedundant(w.expert_down[se])) + w.expert_down[se].ReleaseHost(); } } mr.ready = true; @@ -7583,9 +7707,9 @@ vt::Tensor detail::ExpertSliceForTest(vt::Queue& q, const OwnedTensor& w, return KqExpertSlice(d, w, N, K, row_off, expert); } -void detail::StageWeightForTest(vt::Queue& q, const OwnedTensor& w) { +vt::Tensor detail::StageWeightForTest(vt::Queue& q, const OwnedTensor& w) { Dev d{vt::GetBackend(q.device.type), q}; - (void)ResidentWeight(d, w); + return ResidentWeight(d, w); } void detail::EndExpertStreamStep() { Qwen35ExpertStream::EndStepIfActive(); } diff --git a/src/vllm/model_executor/models/qwen3_5_internal.h b/src/vllm/model_executor/models/qwen3_5_internal.h index 5fbcf0a29..08f2cc16e 100644 --- a/src/vllm/model_executor/models/qwen3_5_internal.h +++ b/src/vllm/model_executor/models/qwen3_5_internal.h @@ -526,7 +526,7 @@ void ExpertStreamSetForceFallback(bool on); // whole value is that it never fires in production. vt::Tensor ExpertSliceForTest(vt::Queue& q, const OwnedTensor& w, int64_t N, int64_t K, int64_t row_off, int64_t expert); -void StageWeightForTest(vt::Queue& q, const OwnedTensor& w); +vt::Tensor StageWeightForTest(vt::Queue& q, const OwnedTensor& w); // End one decode step for the streamed-expert cache. The Qwen3.5 MoE forward // runs this from its own layer driver; a SECOND full-attention MoE model diff --git a/src/vllm/model_executor/models/qwen3_5_weights.cpp b/src/vllm/model_executor/models/qwen3_5_weights.cpp index 1f94c70aa..aa7b8f71a 100644 --- a/src/vllm/model_executor/models/qwen3_5_weights.cpp +++ b/src/vllm/model_executor/models/qwen3_5_weights.cpp @@ -9,6 +9,7 @@ #include #include #include +#include // over-aligned operator new/delete (kDeviceAliasAlignment) #include #include #include @@ -123,6 +124,149 @@ void ReleaseDirectUploadSource(const OwnedTensor& w) { } // namespace +namespace { + +// Drop the resident anonymous pages of `[p, p + nb)` without touching the +// allocator's boundary metadata. Interior WHOLE pages only, for the reason +// spelled out in `ReleaseHost` above: glibc's `free()` alone often leaves the +// pages resident on the sbrk arena's free list, and the whole subject here is +// the RSS. +void DropResidentInteriorPages(const uint8_t* p, size_t nb) { +#if defined(__unix__) || defined(__APPLE__) + if (p == nullptr || nb == 0) return; + const long ps_l = ::sysconf(_SC_PAGESIZE); + const auto ps = static_cast(ps_l > 0 ? ps_l : 4096); + const auto begin = reinterpret_cast(p); + const uintptr_t end = begin + nb; + const uintptr_t page_begin = (begin + ps - 1) & ~(ps - 1); + const uintptr_t page_end = end & ~(ps - 1); + if (page_end > page_begin) { + ::madvise(reinterpret_cast(page_begin), + static_cast(page_end - page_begin), MADV_DONTNEED); + } +#else + (void)p; + (void)nb; +#endif +} + +} // namespace + +namespace { + +HostAliasStats& AliasStats() { + static HostAliasStats s; + return s; +} + +} // namespace + +HostAliasStats HostAliasSnapshot() { return AliasStats(); } + +bool HostWeightAliasEnabled() { + static const bool on = [] { + const char* e = std::getenv("VT_QWEN35_ALIAS_HOST_WEIGHTS"); + return !(e != nullptr && e[0] == '0'); + }(); + return on; +} + +bool MakeHostBytesDeviceAliasable(const OwnedTensor& w, + HostAliasOutcome* outcome) { + HostAliasStats& st = AliasStats(); + ++st.calls; + const uint64_t nbytes = static_cast(w.bytes.size()); + auto report = [&](HostAliasOutcome o, bool ok) { + if (outcome != nullptr) *outcome = o; + switch (o) { + case HostAliasOutcome::kAliasedInPlace: + st.aliased_in_place_bytes += nbytes; + break; + case HostAliasOutcome::kRehomed: + st.rehomed_bytes += nbytes; + break; + case HostAliasOutcome::kDeclinedBorrow: + st.declined_borrow_bytes += nbytes; + break; + default: + st.declined_other_bytes += nbytes; + break; + } + return ok; + }; + if (!HostWeightAliasEnabled()) + return report(HostAliasOutcome::kDeclinedDisabled, false); + // Nothing to alias. The caller refuses this by name rather than handing a + // kernel a null pointer; see ResidentWeight. + if (w.bytes.empty()) return report(HostAliasOutcome::kDeclinedEmpty, false); + if (reinterpret_cast(w.bytes.data()) % kDeviceAliasAlignment == 0) { + // ISSUE #150's WINDOWED RELEASE STILL HAS TO HAPPEN, and this branch is the + // third path that can skip it. A safetensors direct-upload borrow whose + // offset happens to be a multiple of 256 is aliased here and never reaches + // `AdoptDeviceBytesAsHost`, which is the only other caller of + // `ReleaseDirectUploadSource` — so its consumed source pages would stay + // resident, data-dependently, for roughly one borrow in eight. The comment + // on the ordering inside `AdoptDeviceBytesAsHost` insists the release + // happens on EVERY path including the `VT_ADOPT_DEVICE_BYTES=0` arm, and + // this honours the same rule. A no-op unless `mmap_src` is set, which the + // GGUF loader never does. + ReleaseDirectUploadSource(w); + return report(HostAliasOutcome::kAliasedInPlace, true); + } + // A borrow owns no anonymous pages, so re-homing it would ADD residency + // instead of removing it, and a tied pair's shared expansion must keep its one + // keep-alive. Same reasoning, and the same answer, as `ReleaseHost`'s and + // `AdoptDeviceBytesAsHost`'s borrowed branches. + // + // ONE EXCEPTION TO "A BORROW OWNS NO ANONYMOUS PAGES" NOW EXISTS, AND IT IS + // THE ONE THIS FUNCTION CREATES. The block below turns an OWNED buffer into a + // BORROWED one whose keep-alive is an over-aligned `operator new` block — + // anonymous memory. That does not change the answer here (such a buffer is + // already aligned and returns above), but it does mean the sentence is no + // longer universally true, and the three places that reason from it — + // `ReleaseHost`'s borrowed branch, `AdoptDeviceBytesAsHost`'s, and this one — + // are now reasoning about GGUF mappings and tied expansions specifically. The + // consequence worth naming: `ReleaseHost()` on a re-homed weight drops the + // keep-alive rather than madvising, which frees the block through the deleter + // and is correct, but it does not take the `MADV_DONTNEED` path. Nothing calls + // `ReleaseHost` on a dense weight today; `HostMirrorIsRedundant` is what keeps + // the one caller that could from doing it to an aliased one. + if (w.bytes.borrowed()) + return report(HostAliasOutcome::kDeclinedBorrow, false); + + // SINGLE-THREADED BY PRECONDITION, stated rather than enforced. The + // re-pointing below destroys the source vector, so two threads reaching it for + // the same weight would race — and unlike `AdoptDeviceBytesAsHost`, which + // hides behind `if (!w.d_dev)`, this branch has no memo and re-tests alignment + // on every call. It is safe because the first touch of every weight happens + // inside one forward on one thread, which is the same assumption the `d_dev` + // memo two branches down has always made. A model that ever builds residents + // from several threads must add a `call_once` here and there. + auto& self = *const_cast(&w); + const size_t nb = self.bytes.size(); + // Over-aligned `operator new` rather than `aligned_alloc`/`posix_memalign`: + // it is the one spelling that is standard C++17 AND available on MSVC, which + // this tree still compiles for, and it does not require the size to be a + // multiple of the alignment. + void* p = ::operator new(nb, std::align_val_t{kDeviceAliasAlignment}); + // The keep-alive is built IMMEDIATELY, before anything that can throw, so the + // block is owned from the instant it exists. `shared_ptr`'s own control-block + // allocation is the throwing step, and holding a raw `p` across it is how an + // allocation leaks on a path nobody tests. + std::shared_ptr keep(static_cast(p), [](const void* q) { + ::operator delete(const_cast(q), std::align_val_t{kDeviceAliasAlignment}); + }); + std::memcpy(p, self.bytes.data(), nb); + // ORDER: release the OLD pages while they are still mapped, then re-point. + // The assignment below destroys the vector that owns them, and madvise'ing a + // range after it has been unmapped is at best a silent no-op and at worst + // discards whatever mapped into the hole first — the same trap + // `AdoptDeviceBytesAsHost` documents at length. + DropResidentInteriorPages(self.bytes.data(), nb); + self.bytes = OwnedBytes::Borrow(static_cast(p), nb, std::move(keep)); + return report(HostAliasOutcome::kRehomed, true); +} + void AdoptDeviceBytesAsHost(vt::Backend& backend, const OwnedTensor& w) { if (w.d_dev == nullptr) return; // ENG-LOAD-DIRECT-UPLOAD: a direct-upload borrow is the ONE borrow that may be diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt index 8b4129f1e..986bbffb0 100644 --- a/tests/CMakeLists.txt +++ b/tests/CMakeLists.txt @@ -1186,6 +1186,14 @@ target_include_directories(test_expert_stream_mixed_slot PRIVATE ${CMAKE_SOURCE_ # other suite. vllm_cpp_add_test(test_expert_stream_device_slot vllm/model_executor/test_expert_stream_device_slot.cpp) target_include_directories(test_expert_stream_device_slot PRIVATE ${CMAKE_SOURCE_DIR}/src) +# ENG-EXPERT-STREAM-DEVICE W0f (#1299): which branch `ResidentWeight` takes on a +# staging platform whose kernels can dereference host storage. Its own binary for +# the same reason as its sibling above -- it registers a fake backend and a fake +# platform in the process-global kXPU slot, which must not leak into any other +# suite -- and separately because the two disagree about the backend's +# UnifiedMemory(), which is a process-global registration and not a per-case flag. +vllm_cpp_add_test(test_resident_weight_host_addressable vllm/model_executor/test_resident_weight_host_addressable.cpp) +target_include_directories(test_resident_weight_host_addressable PRIVATE ${CMAKE_SOURCE_DIR}/src) # ENG-EXPERT-STREAM (#1091): the step clock at the MoE entry points ForwardLayers # does NOT cover, and the final statistics line. Its own binary because it is the # only one that can observe `steps == 0` -- the store is a process-lifetime diff --git a/tests/vllm/model_executor/test_expert_stream_device_slot.cpp b/tests/vllm/model_executor/test_expert_stream_device_slot.cpp index d4529cccd..0b06026db 100644 --- a/tests/vllm/model_executor/test_expert_stream_device_slot.cpp +++ b/tests/vllm/model_executor/test_expert_stream_device_slot.cpp @@ -316,15 +316,44 @@ TEST_CASE("a STREAMED tower that reaches device staging is refused BY NAME") { CHECK(tower.d_dev == nullptr); } -TEST_CASE("an unclaimed tower still stages normally, so the refusal is not a blanket") { +TEST_CASE("an unclaimed tower is served normally, so the refusal is not a blanket") { // The negative control for the case above. A refusal that fired for every // tower would pass that case and break every model, so the same helper must - // succeed on a tower the lane never touched. + // SUCCEED on a tower the lane never touched. + // + // WHAT IT MEANS BY "NORMALLY" CHANGED UNDER IT, and the change is the point of + // W0f (issue #1299). This case used to assert `d_dev != nullptr` — that the + // tower was STAGED. On a host-addressable platform `ResidentWeight` no longer + // stages anything: it hands back the tower's own host bytes, because the + // second copy bought nothing on a part where device memory IS host memory and + // cost the whole box. So the assertion here is the property this case was + // always about — the refusal did NOT fire and a usable tensor came back — + // stated without pinning WHICH residency serves it. The discrete arm, where + // "normally" still means a staged copy, is asserted in the case above and in + // `test_resident_weight_host_addressable`. Platform_().host_addressable = true; const OwnedTensor plain = MakeTower(/*tag=*/5); Queue q = XpuQueue(); CHECK_FALSE(plain.expert_streamed); - CHECK_NOTHROW(vllm::detail::StageWeightForTest(q, plain)); - CHECK(plain.d_dev != nullptr); + Tensor t; + CHECK_NOTHROW(t = vllm::detail::StageWeightForTest(q, plain)); + REQUIRE(t.data != nullptr); + // Served from the tower's own host bytes, and not staged (W0f). + CHECK(plain.d_dev == nullptr); + CHECK(t.data == static_cast(plain.bytes.data())); + CHECK(std::memcmp(t.data, plain.bytes.data(), plain.bytes.size()) == 0); + + // And on a DISCRETE platform the same helper still stages, which is what keeps + // "the refusal is not a blanket" true on both arms rather than only the one + // W0f changed. + Platform_().host_addressable = false; + const OwnedTensor discrete = MakeTower(/*tag=*/6); + const int allocs_before = Fake().allocs; + Tensor dt; + CHECK_NOTHROW(dt = vllm::detail::StageWeightForTest(q, discrete)); + CHECK(discrete.d_dev != nullptr); + CHECK(Fake().allocs == allocs_before + 1); + CHECK(dt.data == discrete.d_dev.get()); + Platform_().host_addressable = true; } diff --git a/tests/vllm/model_executor/test_resident_weight_host_addressable.cpp b/tests/vllm/model_executor/test_resident_weight_host_addressable.cpp new file mode 100644 index 000000000..5ee550ace --- /dev/null +++ b/tests/vllm/model_executor/test_resident_weight_host_addressable.cpp @@ -0,0 +1,481 @@ +// ENG-EXPERT-STREAM-DEVICE W0f (issue #1299): on a platform whose kernels can +// DEREFERENCE host storage, does `ResidentWeight` hand back the host bytes +// instead of allocating a second copy of them? +// +// WHAT THE DEFECT IS, AND WHY A TOKEN GATE CANNOT SEE IT. `ResidentWeight`'s +// staging branch is a VERBATIM byte copy: `d.b.Alloc(w.bytes.size())`, +// `d.b.Copy(...)`, then a tensor over the copy with the same dtype, the same +// shape and the same (dropped) marker set. Nothing about the bytes changes. On a +// discrete device that copy is the whole point — the kernel cannot follow a host +// pointer. On a part where device memory IS host memory it buys nothing and +// costs a second full resident copy of every dense weight, out of the same RAM +// the first one came from. The tokens are identical either way, which is exactly +// why this is asserted as an ALLOCATION and never as an output. +// +// MEASURED (issue #1299, `dgx:gpu0`, GB10, seven runs). With W0's lane on, +// `Qwen3.8-2.4T-A95B UD-Q1_0` loads on `--device cuda` — 61.20 GiB resident, +// ~265 s — and then exhausts a 119.631 GiB box inside the FIRST forward, zero +// decode steps, every time. A 0.15 GiB slot arena died exactly where an +// 18.55 GiB one did, so the arena is not the cost; the growth was anonymous +// (`RssAnon` 8.1 -> 61.4 GB) while file-backed stayed flat, so the mapping is not +// pinned. About 39 GiB of that 61.20 is `attn_qkv` (21.56) and `ssm_out` (17.25), +// which the GDN V-head reorder makes `kTransformedWeight` and therefore expands +// to bf16 in OWNED host buffers — the split is measured, not derived, in +// `.agents/specs/expert-streaming.md`. The CPU arm pays that once and serves. The +// CUDA arm paid it twice and could not. +// +// THE PLATFORM THIS NEEDS DOES NOT EXIST ON A CPU TIER. The branch under test is +// selected by `needs_weight_staging() && host_memory_is_device_addressable()`, +// and exactly one machine this project can reach answers that pair (a GB10). A +// fake platform in the otherwise-unused kXPU slot supplies the bit, over a fake +// backend whose `UnifiedMemory() == true` and `DeviceMemoryIsHostAddressable() +// == false` — which is not an arbitrary pair but the GB10 CUDA backend's own +// answers (`src/vt/cuda/cuda_backend.cu:113` and the base default in +// `include/vt/backend.h:76`; a `cudaMalloc` pointer there is still not +// host-dereferenceable even though host and device address the same RAM). +// +// WHAT THIS FILE DOES NOT CLAIM. It does not claim to prove that a forward +// reaches `ResidentWeight`. `test_expert_stream_wiring` proves that through +// `Qwen3_5Model::Forward`, and the reachability mutation for this change is +// stated against that binary, not this one. What is proved here is which BRANCH +// the function takes, and that the predicate — not the device name, not the +// staging flag — is what selects it. +#include + +#include +#include +#include +#include +#include +#include +#include + +#include "vllm/model_executor/models/owned_bytes.h" +#include "vllm/model_executor/models/qwen3_5_internal.h" +#include "vllm/model_executor/models/qwen3_5_weights.h" +#include "vllm/platforms/interface.h" +#include "vt/backend.h" +#include "vt/device.h" +#include "vt/dtype.h" +#include "vt/tensor.h" + +namespace { + +using vllm::OwnedTensor; +using vt::Device; +using vt::DeviceType; +using vt::DType; +using vt::Queue; +using vt::Tensor; + +// A backend over ordinary host memory, standing in for a device allocator. The +// point of malloc is that a STAGED weight is then a real, inspectable allocation +// at an address that differs from the weight's own bytes — which is how a case +// tells "staged" from "aliased" without a GPU. +class HostBackend final : public vt::Backend { + public: + void* Alloc(size_t bytes) override { + ++allocs; + alloc_bytes += bytes; + return std::malloc(bytes == 0 ? 1 : bytes); + } + void Free(void* p) override { std::free(p); } + void Memset(Queue&, void* p, int v, size_t bytes) override { + std::memset(p, v, bytes); + } + void Copy(Queue&, void* dst, const void* src, size_t bytes) override { + std::memcpy(dst, src, bytes); + } + Queue CreateQueue() override { return Queue{Device{DeviceType::kXPU, 0}, nullptr}; } + void DestroyQueue(Queue&) override {} + // TRUE, matching GB10's CUDA backend (cuda_backend.cu:113): host and device + // address the same physical RAM. + bool UnifiedMemory() const override { return true; } + // FALSE, also matching GB10: a `cudaMalloc` pointer is NOT host-dereferenceable + // there, which is the asymmetry this whole change turns on. It is also what + // keeps `AdoptDeviceBytesAsHost` inert, so a staged weight stays visibly + // staged instead of being folded back onto its host buffer by a second + // mechanism and confusing what this file measures. + bool DeviceMemoryIsHostAddressable() const override { return false; } + + // Counted, so a case can say HOW MANY allocations it observed rather than only + // that a pointer was or was not null. + int allocs = 0; + size_t alloc_bytes = 0; +}; + +HostBackend& Fake() { + static HostBackend b; + return b; +} + +// The GB10 shape: a platform that STAGES its ordinary weights and whose kernels +// can nevertheless dereference host storage. `host_addressable` is a settable +// field rather than a second registered platform because the platform registry +// is process-global — two registrations would fight, and one flag lets a case +// prove the predicate SELECTS by moving only the bit under test. +class FakeDevicePlatform final : public vllm::platforms::Platform { + public: + DeviceType device_type() const override { return DeviceType::kXPU; } + vt::Backend& backend() const override { return Fake(); } + vllm::platforms::DeviceCapability get_device_capability() const override { + return {12, 1}; + } + std::vector supported_dtypes() const override { return {DType::kBF16}; } + vllm::platforms::ResidencyPolicy residency_policy() const override { return {}; } + bool needs_weight_staging() const override { return true; } + bool host_memory_is_device_addressable() const override { + return host_addressable; + } + + bool host_addressable = true; +}; + +FakeDevicePlatform& Platform_() { + static FakeDevicePlatform p; + return p; +} + +struct Registrar { + Registrar() { + vt::RegisterBackend(Device{DeviceType::kXPU, 0}, &Fake()); + vllm::platforms::RegisterPlatform(DeviceType::kXPU, &Platform_()); + } +}; +const Registrar kRegistrar; + +Queue XpuQueue() { return Queue{Device{DeviceType::kXPU, 0}, nullptr}; } + +constexpr int64_t kN = 6; +constexpr int64_t kK = 8; + +// A plain bf16 [N,K] projection — the shape every dense weight in the Qwen3.5 +// forward arrives as, and the shape the ~39 GiB of transformed `attn_qkv` / +// `ssm_out` arrives as on the target checkpoint. +OwnedTensor MakeWeight(uint8_t tag) { + OwnedTensor t; + t.dtype = DType::kBF16; + t.rank = 2; + t.shape[0] = kN; + t.shape[1] = kK; + t.nk = true; + std::vector b(static_cast(kN * kK) * 2); + for (size_t i = 0; i < b.size(); ++i) + b[i] = static_cast((i * 7 + tag) & 0xFF); + t.bytes = vllm::OwnedBytes(std::move(b)); + return t; +} + +// Restore `host_addressable` on EVERY exit path, including a REQUIRE that +// aborts the case body. Three cases below move the bit; a case that fails +// halfway used to leak `false` into every later case in this binary and turn one +// finding into a cascade of confusing ones. +struct PlatformArm { + explicit PlatformArm(bool on) : prev(Platform_().host_addressable) { + Platform_().host_addressable = on; + } + ~PlatformArm() { Platform_().host_addressable = prev; } + bool prev; +}; + +} // namespace + +TEST_CASE("a host-addressable staging device ALIASES the weight and allocates nothing") { + const PlatformArm arm(true); + const OwnedTensor w = MakeWeight(/*tag=*/1); + // A copy of the bytes taken BEFORE the call, because the call may move them. + const std::vector expect(w.bytes.data(), w.bytes.data() + w.bytes.size()); + Queue q = XpuQueue(); + const int allocs_before = Fake().allocs; + + const Tensor t = vllm::detail::StageWeightForTest(q, w); + + // THE #1299 ASSERTION, and it is an allocation rather than an output. `d_dev` + // is set by, and only by, the staging branch; null means the second resident + // copy of this weight does not exist. + CHECK(w.d_dev == nullptr); + CHECK(Fake().allocs == allocs_before); + + // ...and the tensor the kernel binds points at the weight's OWN host bytes. + // Both halves matter: a null `d_dev` with a pointer somewhere else would be a + // different defect wearing the same green. + REQUIRE(t.data != nullptr); + CHECK(t.data == static_cast(w.bytes.data())); + + // THE SUBSTITUTION IS INDISTINGUISHABLE, which is the whole safety argument. + // The pointer this branch hands a kernel is at least as aligned as the + // `cudaMalloc` pointer it replaces, so no consumer — not the 16-byte + // `cp.async` granule, not cuBLASLt's 256-byte minimum-alignment promise — can + // tell that the staging copy is gone. A plain `std::vector` does NOT + // satisfy this on its own (glibc gives 16), so a green here is also the + // statement that the re-homing ran. + CHECK(reinterpret_cast(t.data) % vllm::kDeviceAliasAlignment == 0); + + // ...and the bytes did not change while being re-homed. Aliasing the WRONG + // bytes is the one defect that would otherwise pass every assertion above. + CHECK(std::memcmp(t.data, expect.data(), expect.size()) == 0); + CHECK(w.bytes.size() == expect.size()); + + // The tensor is otherwise EXACTLY what the staging branch produced: same + // dtype, same device, same shape, same strides. This branch is that branch + // minus the copy, and that is the whole claim. + CHECK(t.dtype == w.dtype); + CHECK(t.device.type == DeviceType::kXPU); + CHECK(t.rank == 2); + CHECK(t.shape[0] == kN); + CHECK(t.shape[1] == kK); + CHECK(t.stride[0] == kK); + CHECK(t.stride[1] == 1); +} + +TEST_CASE("the alias is stable across calls and still allocates nothing") { + // `d_dev` is the staging branch's memo, so a branch that does not set it must + // not become a per-call allocation instead. Two calls, one address, zero + // allocations. + const PlatformArm arm(true); + const OwnedTensor w = MakeWeight(/*tag=*/2); + Queue q = XpuQueue(); + const int allocs_before = Fake().allocs; + + const Tensor a = vllm::detail::StageWeightForTest(q, w); + const Tensor b = vllm::detail::StageWeightForTest(q, w); + + CHECK(Fake().allocs == allocs_before); + CHECK(w.d_dev == nullptr); + CHECK(a.data == b.data); + CHECK(a.data == static_cast(w.bytes.data())); +} + +TEST_CASE("a device that CANNOT read host memory stages exactly as before") { + // THE DISCRETE ANSWER, and it must be byte-identical to today. A GPU whose + // kernels cannot follow a host pointer gets the allocation, the copy, the + // memo, and a tensor over the copy — the pre-W0f behaviour, unchanged. This is + // the case that proves the predicate SELECTS rather than that the branch was + // taken unconditionally. + const PlatformArm arm(false); + const OwnedTensor w = MakeWeight(/*tag=*/3); + Queue q = XpuQueue(); + const int allocs_before = Fake().allocs; + const size_t bytes_before = Fake().alloc_bytes; + + const Tensor t = vllm::detail::StageWeightForTest(q, w); + + REQUIRE(w.d_dev != nullptr); + CHECK(Fake().allocs == allocs_before + 1); + CHECK(Fake().alloc_bytes == bytes_before + w.bytes.size()); + CHECK(t.data == w.d_dev.get()); + CHECK(t.data != static_cast(w.bytes.data())); + // The staged copy holds the same bytes, which is what makes the aliasing arm + // above a legitimate substitution rather than a different weight. + CHECK(std::memcmp(t.data, w.bytes.data(), w.bytes.size()) == 0); +} + +TEST_CASE("the aliasing branch keeps the elem_kn_repacked refusal") { + // VT_CPU_ELEM_KN_REPACK transposes the buffer to [K,N] and ONLY the CPU + // MatmulBTKernel honours the marker. Where the bytes live has nothing to do + // with which kernel reads them: on a host-addressable device the reader is + // still the DEVICE kernel, which would read transposed bytes as [N,K] and + // produce garbage silently. The refusal therefore has to survive the new + // branch, and this is the case that says so. + const PlatformArm arm(true); + OwnedTensor w = MakeWeight(/*tag=*/4); + w.elem_kn_repacked = true; + Queue q = XpuQueue(); + + CHECK_THROWS_WITH_AS( + vllm::detail::StageWeightForTest(q, w), + doctest::Contains("an elem_kn_repacked ([K,N]) weight reached device staging"), + std::runtime_error); +} + +TEST_CASE("an i8mm-repacked weight reaching device residency is refused BY NAME") { + // A DEFECT FOUND WHILE WRITING W0f (issue #1320), fixed in the same flow, and + // it predates this change on both branches. + // + // `VT_CPU_QUANT_REPACK` rewrites a Q8_0 weight into the `block_q8_0x4` i8mm + // interleave at load. Only the CPU `MatmulBTKernel` understands that layout; + // the CUDA quant dot reads plain `block_q8_0`. Its sibling transform + // `elem_kn_repack` has BOTH a CPU-platform gate in the loader policy and the + // refusal above — `quant_repack` had NEITHER. It rides `QuantRepackActive()`, + // which probes the HOST CPU for Arm i8mm, so an aarch64 box doing + // `--device cuda` (which is precisely the target box) can repack a weight and + // then hand it to a kernel that misreads it. The result is wrong tokens, not a + // crash, and no gate in this tree could see it. + // + // Harmless on the target checkpoint as measured — one Q8_0 tensor, 0.01% of + // parameters, and the instrumented load recorded `quant_repack = 0` — which is + // why this is a tripwire beside its sibling rather than a campaign. + const PlatformArm arm(true); + OwnedTensor w = MakeWeight(/*tag=*/9); + w.repacked = true; + Queue q = XpuQueue(); + + CHECK_THROWS_WITH_AS( + vllm::detail::StageWeightForTest(q, w), + doctest::Contains("an i8mm-repacked (block_q8_0x4) weight reached device residency"), + std::runtime_error); + + // ...and on the DISCRETE arm too, because the kernel that misreads it is the + // same kernel either way. Where the bytes live was never the question. + { + const PlatformArm discrete(false); + OwnedTensor d = MakeWeight(/*tag=*/10); + d.repacked = true; + CHECK_THROWS_WITH_AS( + vllm::detail::StageWeightForTest(q, d), + doctest::Contains("an i8mm-repacked (block_q8_0x4) weight reached device residency"), + std::runtime_error); + } +} + +TEST_CASE("the aliasing branch keeps the streamed-tower refusal") { + // Same reasoning as the case above, for the W0c tripwire. A tower the expert + // lane claimed must never be consumed WHOLE, and "whole" is a statement about + // the tower, not about the allocator: reaching here at all means the lane was + // defeated. Cheap on this platform and catastrophic on the other, so it fails + // by name on both. + const PlatformArm arm(true); + OwnedTensor w = MakeWeight(/*tag=*/5); + w.expert_streamed = true; + Queue q = XpuQueue(); + + CHECK_THROWS_WITH_AS( + vllm::detail::StageWeightForTest(q, w), + doctest::Contains("a STREAMED expert tower reached device staging"), + std::runtime_error); +} + +TEST_CASE("an ALREADY-ALIGNED buffer is aliased in place, with no second copy") { + // The free case, and it needs its own assertion because the re-homing above + // would satisfy every other check in this file while quietly copying a weight + // that did not need copying. A GGUF mmap borrow whose tensor offset happens to + // be a multiple of 256 lands here, and so does the SECOND call for any weight + // the first call re-homed. + const PlatformArm arm(true); + OwnedTensor w = MakeWeight(/*tag=*/7); + const size_t nb = w.bytes.size(); + // An aligned block, borrowed, standing exactly where a lucky mmap offset + // would. `Borrow` needs a keep-alive, and the block IS the keep-alive. + void* aligned = ::operator new(nb, std::align_val_t{vllm::kDeviceAliasAlignment}); + std::memcpy(aligned, w.bytes.data(), nb); + std::shared_ptr keep( + static_cast(aligned), [](const void* p) { + ::operator delete(const_cast(p), + std::align_val_t{vllm::kDeviceAliasAlignment}); + }); + w.bytes = vllm::OwnedBytes::Borrow(static_cast(aligned), nb, keep); + REQUIRE(reinterpret_cast(w.bytes.data()) % + vllm::kDeviceAliasAlignment == 0); + Queue q = XpuQueue(); + const int allocs_before = Fake().allocs; + + CHECK(vllm::MakeHostBytesDeviceAliasable(w)); + // The SAME address: nothing was moved and nothing was allocated. + CHECK(w.bytes.data() == static_cast(aligned)); + + const Tensor t = vllm::detail::StageWeightForTest(q, w); + CHECK(w.d_dev == nullptr); + CHECK(Fake().allocs == allocs_before); + CHECK(t.data == static_cast(aligned)); +} + +TEST_CASE("a MISALIGNED BORROW is not re-homed, and stages instead") { + // THE CASE THAT KEEPS THIS CHANGE FROM BACKFIRING. A borrow owns no anonymous + // pages — it is a clean, file-backed GGUF mapping, or a tied + // token_embd/lm_head pair's single shared expansion. Copying it into an + // aligned anonymous block to satisfy the alias would CREATE the residency this + // row exists to remove, and would break the tie. GGUF guarantees only 32-byte + // tensor alignment, so this is a real population and not a hypothetical. + // + // The correct answer is to decline, and let the (unchanged) staging branch + // copy it into device memory, where its file pages stay reclaimable. + const PlatformArm arm(true); + OwnedTensor w = MakeWeight(/*tag=*/8); + const size_t nb = w.bytes.size(); + auto backing = std::make_shared>(nb + vllm::kDeviceAliasAlignment); + // Deliberately off by 8: aligned enough for the element type, nowhere near 256. + uint8_t* base = backing->data(); + uint8_t* off = base + (vllm::kDeviceAliasAlignment - + (reinterpret_cast(base) % + vllm::kDeviceAliasAlignment)) + 8; + std::memcpy(off, w.bytes.data(), nb); + w.bytes = vllm::OwnedBytes::Borrow( + off, nb, std::static_pointer_cast(backing)); + REQUIRE(w.bytes.borrowed()); + REQUIRE(reinterpret_cast(w.bytes.data()) % + vllm::kDeviceAliasAlignment != 0); + Queue q = XpuQueue(); + const int allocs_before = Fake().allocs; + + CHECK_FALSE(vllm::MakeHostBytesDeviceAliasable(w)); + // The borrow is UNTOUCHED — same address, still borrowed. + CHECK(w.bytes.data() == off); + CHECK(w.bytes.borrowed()); + + const Tensor t = vllm::detail::StageWeightForTest(q, w); + REQUIRE(w.d_dev != nullptr); + CHECK(Fake().allocs == allocs_before + 1); + CHECK(t.data == w.d_dev.get()); + CHECK(std::memcmp(t.data, off, nb) == 0); +} + +TEST_CASE("an ALIASED weight's host mirror is NOT redundant, so nothing may free it") { + // THE USE-AFTER-FREE A FRESH REVIEW CAUGHT (#1299). `MoeBlockBf16Cuda` + // captures `ResidentWeight(...).data` for every expert into a DEVICE-resident + // pointer table, uploads the table once, and then releases the host mirrors. + // Its own comment justified that with "once the device copy exists it is + // authoritative and nothing reads the host bytes again", which held while + // `ResidentWeight` had two behaviours. It has three: this branch ALIASES, so + // the captured pointers ARE `w.bytes.data()`, and the release frees memory the + // resident table still points at for the model's lifetime, from inside + // captured graphs. The reviewer demonstrated it with a scratch case that takes + // SIGSEGV. + // + // THIS CASE DOES NOT DEREFERENCE FREED MEMORY, deliberately: a segfault is a + // red that also destroys the rest of the binary's report, and a gate should + // fail by assertion. It asserts the DECISION instead, on both arms, which is + // the thing the production site now asks. + const PlatformArm arm(true); + const OwnedTensor aliased = MakeWeight(/*tag=*/11); + Queue q = XpuQueue(); + const Tensor t = vllm::detail::StageWeightForTest(q, aliased); + + REQUIRE(aliased.d_dev == nullptr); + REQUIRE(t.data == static_cast(aliased.bytes.data())); + // There is no device copy, so the host bytes are the ONLY copy and releasing + // them would free what the kernel reads. + CHECK_FALSE(vllm::HostMirrorIsRedundant(aliased)); + + // The discrete arm is the other half: a staged weight DOES have an + // authoritative device copy, and the release that predates W0f stays correct + // for it. Without this half the invariant could be satisfied by refusing every + // release, which would silently undo a measured host-memory lever. + { + const PlatformArm discrete(false); + const OwnedTensor staged = MakeWeight(/*tag=*/12); + const Tensor dt = vllm::detail::StageWeightForTest(q, staged); + REQUIRE(staged.d_dev != nullptr); + CHECK(dt.data == staged.d_dev.get()); + CHECK(vllm::HostMirrorIsRedundant(staged)); + } +} + +TEST_CASE("a weight whose host bytes are GONE is refused by name, not aliased to null") { + // THE LIFETIME PRECONDITION, stated in code. The aliasing branch hands out + // `w.bytes.data()` and keeps no reference of its own, so it is correct only + // while the weight owns those bytes. `ReleaseHost()` is the one operation that + // takes them away — it is reachable for the routed-expert fp4/Marlin mirrors + // (`ShouldReleaseHostWeights`, qwen3_5.cpp) though not for the dense weights + // this branch serves. If that ever changes, the failure without this check is + // a null weight pointer inside a kernel, which is a segfault at best and wrong + // tokens at worst. With it, it is one legible sentence. + const PlatformArm arm(true); + OwnedTensor w = MakeWeight(/*tag=*/6); + w.ReleaseHost(); + REQUIRE(w.bytes.empty()); + Queue q = XpuQueue(); + + CHECK_THROWS_WITH_AS( + vllm::detail::StageWeightForTest(q, w), + doctest::Contains("has no host bytes"), std::runtime_error); +} diff --git a/tests/vllm/test_load_direct_upload.cpp b/tests/vllm/test_load_direct_upload.cpp index 4bdd79a56..0b8eea8c0 100644 --- a/tests/vllm/test_load_direct_upload.cpp +++ b/tests/vllm/test_load_direct_upload.cpp @@ -522,6 +522,41 @@ TEST_CASE("adopt: the source pages are released even where device memory is NOT CHECK(m.addr[0] == 0); } +TEST_CASE("alias: an ALIASED direct-upload borrow still releases its source pages") { + // ENG-EXPERT-STREAM-DEVICE W0f (#1299), found by a fresh review. W0f gave + // `ResidentWeight` a third residency: on a platform whose kernels can + // dereference host storage it ALIASES the bytes and never populates `d_dev`, + // so `AdoptDeviceBytesAsHost` is never called for that weight. That function + // is the only other caller of `ReleaseDirectUploadSource`, so the alias branch + // became a third path past a release whose own ordering comment insists it + // happens on EVERY path, including the `VT_ADOPT_DEVICE_BYTES=0` arm. + // + // A borrow reaches the alias branch when it is already + // `kDeviceAliasAlignment`-aligned, and an `mmap` return always is, so this is + // not a corner case: it is every direct-upload borrow on such a platform. + // Without the fix the mapping's consumed pages stay resident for the life of + // the process, silently, and issue #150's measurement stops being true. + ForcedResidencyArm arm; + ObservableMapping m; + vllm::OwnedTensor w = BorrowedWeight(m, 2 * PageSize()); + REQUIRE(w.bytes.borrowed()); + REQUIRE(w.mmap_src == static_cast(m.addr)); + // Asserted, not assumed: the case proves nothing if the borrow declines. + REQUIRE(reinterpret_cast(w.bytes.data()) % + vllm::kDeviceAliasAlignment == 0); + + CHECK(vllm::MakeHostBytesDeviceAliasable(w)); + + // The borrow is UNTOUCHED, because an alias moves nothing... + CHECK(w.bytes.borrowed()); + CHECK(static_cast(w.bytes.data()) == static_cast(m.addr)); + CHECK(w.mmap_src == static_cast(m.addr)); + CHECK_FALSE(m.dropped); + // ...and the consumed source pages were released all the same. Zero means the + // madvise ran; the source pattern would mean it did not. + CHECK(m.addr[0] == 0); +} + TEST_CASE("adopt: VT_ADOPT_DEVICE_BYTES=0 moves ONLY the adoption, not the page release") { ForcedResidencyArm arm; ScopedEnvVar adopt_off("VT_ADOPT_DEVICE_BYTES", "0");