diff --git a/.agents/benchmark-record.md b/.agents/benchmark-record.md index 4822cc7ba..1e7f4eae0 100644 --- a/.agents/benchmark-record.md +++ b/.agents/benchmark-record.md @@ -25349,3 +25349,219 @@ row's process and it was left alone. And the FIRST arm started seconds after the previous holder's four-hour render released the device, with the one-minute load average still at 17.5, which is one reason the first 8000-slot run is the noisiest of the four. Every later arm started from a quiet box. + +## ENG-EXPERT-STREAM-DEVICE 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) + +**Read the W0e section above first; this one is a different run of the same +harness on a different tree.** W0e measured the CPU arm and found the CUDA arm +loading without generating, at source `95883dcae`, which is W0f's parent. This +section is the re-run WITH W0f, at `9c783a8be`. The CPU figures therefore do not +agree between the two sections and must not be mixed: W0e's **11.05 s/token at +4000 slots** is the standing CPU number, taken on its own lease against a live +cache, and the CPU column below is a same-lease interleaved control for the CUDA +arm rather than a second attempt at that measurement. Where they disagree, W0e's +is the one `docs/BENCHMARKS.md` carries. + +**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. +**That does NOT by itself acquit W0f, and an earlier draft of this entry said it +did.** Both of the grounds first offered constrain only the arm W0f cannot reach: +the CPU arm returns at `ResidentWeight`'s `is_cpu()` early return +(the `is_cpu()` branch of `qwen3_5.cpp`'s `ResidentWeight`) roughly ninety lines above the alias branch (the `host_memory_is_device_addressable()` branch of the same function; no line number, because this change moves it), so +`w0f-alias calls = 0` there is true by construction for every possible state of +W0f, correct or corrupt. It shows the branch is platform-gated. It cannot +discriminate "the arms' GEMM arithmetic differs" from "W0f moved a logit". The +experiment that does is recorded under **Algo identity** below. An instrumented +CPU run printing the top-2 logits per step gives the shape of the disagreement: + +| 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. **NO RATIO IS WRITTEN HERE, and +the omission is deliberate.** A ratio of these two medians is not a result: the +correctness gate that would license one FAILS, so it may not be published, +quoted, or carried into `docs/BENCHMARKS.md`. An earlier revision of this entry +did write the figure out in digits in order to disown it, which is the shape +this repository has watched turn a disowned number into a measured one -- the +digits survive a copy-paste and the disclaimer does not. Nothing is lost by +removing them, because both medians are in the table above and anyone entitled +to the quotient can do the division; what is removed is the pre-computed string +a later reader can lift without its gate. 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 must not be mixed. + +**Algo identity: the discriminating experiment, and it clears W0f.** The grounds +first offered above could not separate "the arms' GEMM arithmetic differs" from +"W0f moved a logit", so the mechanism was measured directly on `thor:gpu0` +(NVIDIA Thor `sm_110`, driver 13020, cudart 13000, cuBLASLt 130101, `rc` job +`c625e836`, 2026-08-19 13:50). Thor answers this branch's own predicate TRUE — +`cudaDevAttrPageableMemoryAccess = 1`, `cudaDevAttrIntegrated = 1` — so it is a +member of the population W0f serves, and a `cudaMalloc` pointer there is a real +device pointer exactly as on GB10. + +The probe transcribes both cuBLASLt formulations out of +`src/vt/cuda/cuda_matmul.cu` at this branch — the row-major NN +`MatmulKernelCuda` (weight is operand B) and the column-major TN +`MatmulBTKernelCuda` (weight is operand A) — over six shapes off the +checkpoint's own `embedding_length = 8192` at M = 1, 5 and 32. Twelve +measurements, `PROBE_FAILURES=0`: + +| Question | Result | +|---|---| +| Does a repeated heuristic call return the same selection? | **12/12 identical** | +| Is the tree's unset preference the documented 256 default? | **12/12** `default == MIN_ALIGNMENT 256` | +| Does STATING a weaker 16-byte promise move the selection? | **12/12 unchanged** (`256 == 16`) | +| `cublasLtMatmul` with weight from `cudaMalloc` vs a 256-aligned HOST block | **12/12 bit-exact**, 0 differing elements, every status `SUCCESS` | +| ...and from a 16-aligned-only host block | 12/12 bit-exact as well | + +The selection is reported in full rather than by id alone. At M=1 N=8192 K=8192, +both layouts: `id=66 tile=573 stages=35 splitK=5 reduction=2 swizzle=0 custom=1 +inner=0 ws=163856 waves=0.8000`, identical across all four queries. **The +instrument discriminates:** five DIFFERENT configurations appear across the six +shapes (tiles 393, 537, 573, 576; workspaces 0 through 5,242,896), so a uniform +answer is not a probe that reports one thing regardless. + +The structural reason sits in the API and needs no lease: +`cublasLtMatmulAlgoGetHeuristic` takes `(handle, operationDesc, Adesc, Bdesc, +Cdesc, Ddesc, preference, count, results, returned)` and **no operand +pointers**, so the only channel by which alignment can reach the heuristic is +`CUBLASLT_MATMUL_PREF_MIN_ALIGNMENT_*_BYTES`, which this tree never sets. + +**Conclusion on Thor: identical algo AND bit-exact output, so W0f cannot move a +logit there.** What that leg could not establish is that it ran on Thor and not +on the GB10 the token gate ran on: same predicate class, not the same silicon. + +**The GB10 leg RAN, on the target silicon, and it agrees.** `rc` job +`7c7a05e9-be87-48f4-94ae-1bbe0340f063` on `dgx:gpu0`, worker `rc-worker-4b8lj`, +2026-08-19 17:47 UTC, output `/workspace/w0f-algo-probe/out-rc-worker-4b8lj-20260819T174748.txt`. +The box re-derived inside the job rather than inherited: `NVIDIA GB10 sm_121`, +driver 580.173.02, cudart 13000, driver 13000, cuBLASLt 130101 resolved from +`/usr/local/cuda-13.0/lib64/libcublasLt.so.13.1.1.3`, headers at +`/usr/local/cuda/include`. It answers this branch's predicate TRUE in the job's +own print: `pageableMemoryAccess=1 integrated=1 hostPageTables=1 uva=1`. The same +six shapes crossed with the same two formulations, `PROBE_EXIT=0`, +`PROBE_FAILURES=0`: + +| Question | Result on GB10 | +|---|---| +| Does a repeated heuristic call return the same selection? | **12/12 identical** (`VERDICT-A repeat-identical=YES`) | +| Is the tree's unset preference the documented 256 default? | **12/12** (`VERDICT-B default==256:YES`) | +| Does STATING a weaker 16-byte promise move the selection? | **12/12 unchanged** (`256==16:YES`) | +| `cublasLtMatmul` with weight from `cudaMalloc` vs a 256-aligned HOST block | **12/12 bit-exact**, `differing=0`, every status `SUCCESS` | +| ...and from a 16-aligned-only host block | 12/12 bit-exact as well, no sync error | + +The pointers are printed per case, so the arms are not assumed to differ: e.g. +`d_w=0x331200000 %256=0 | host256=0xed1243ffd000 %256=0 | host16=0xed123bffa010 +%256=16`. The largest case compared `out_elems=262144` with `differing=0`. + +**The GB10 instrument discriminates too**, and more finely than Thor's: at least +five distinct selections appear across the twelve measurements — `id=23 tile=21 +stages=10 splitK=1 ws=0`, `id=13 tile=0 custom=75 ws=0`, `id=21 tile=11 stages=8 +splitK=64 ws=1920`, `id=13 tile=0 custom=96 splitK=16 ws=32784`, `id=21 tile=5 +stages=19 splitK=5 ws=5120`, `id=21 tile=5 stages=20 splitK=9 ws=737280` and +`id=67 tile=18 stages=35 custom=28 ws=0` — so a uniform verdict on alignment is +not an instrument that returns one answer regardless. The GB10 selections differ +from Thor's, which is expected of a different architecture and is itself evidence +that the heuristic was re-resolved rather than replayed. + +**So the attribution is now measured on the target silicon: W0f cannot move a +logit, and the step-7 divergence is not the alias.** What it does NOT establish +is what the divergence IS. "The two arms' GEMM arithmetic" remains the standing +hypothesis and it is not measured; naming the operation that first differs is +owed. The 16-aligned arm also came back bit-exact on both boxes, which is NOT a +licence to lower `kDeviceAliasAlignment`: twelve shapes is not the enumeration, +and cuBLASLt is still promised 256. + +**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`), and **why it faults is UNMEASURED**. An earlier draft wrote +"almost certainly because the pointer is not host memory there", which is a +hypothesis: nothing printed the pointer, nothing called +`cudaPointerGetAttributes` on it, and no fault address was recorded. In a change +whose central risk is handing device kernels host pointers, a segfault whose +cause was guessed at is the finding that must not be dismissed, so it is carried +as unmeasured. The next lease prints `cudaPointerGetAttributes(logits)` in that +callback first. Also owed: a ratified gate for a two-arm comparison whose greedy +path is this finely balanced, which `AGENTS.md` reserves as an explicit operator +decision. And still owed, now that the alias is excluded on the target silicon: +WHAT the step-7 divergence is. Excluding one cause is not identifying another, +and the next traceable hypothesis is a per-operation two-arm comparison of the +step-7 forward, first differing tensor named. diff --git a/.agents/issue-index.md b/.agents/issue-index.md index 50a1b4a45..d8edbe35d 100644 --- a/.agents/issue-index.md +++ b/.agents/issue-index.md @@ -464,6 +464,7 @@ rather than merged. `scripts/check-agent-record.py` gates both. | [#1394](https://github.com/mudler/vllm.cpp/issues/1394) | `ENG-CUDAGRAPH-BREAK` | The CPU paged attention reads `btab[r * bt_row + (j / block_size) * bt_col]` for every `j < seq_lens[r]` without checking that the block table has that many columns, so a caller with a short table gets an out-of-bounds read, a plausible block index out of it, and attention over the WRONG page — silently. `tests/vllm/models/test_qwen3_5_decode_graph_seam.cpp`'s `SpecAttnMeta` supplies one: hardcoded `block_table_num_cols = 1` against shape C's `seq_lens = 24` at `block_size = 16`. Found while fixing [#1380](https://github.com/mudler/vllm.cpp/issues/1380), whose `DevicePool` change moved the bytes after the table and turned the same read into a SIGSEGV on `thor:gpu0` (`gdb` at `src/vt/cpu/cpu_paged_attn.cpp:224` under `FullAttnBlockPaged`). PRE-EXISTING: the case passes at `origin/main` only because the read landed on bytes that decoded to an in-range index. FIXED IN FLOW in both halves — the kernel refuses a short table with one compare per request outside the token loop, and the helper sizes its table for the sequence length it declares. Owned by row `ENG-CUDAGRAPH-BREAK` | bug | | [#1387](https://github.com/mudler/vllm.cpp/issues/1387) | `ENG-EXPERT-STREAM-DEVICE` | `docs/FEATURES.md`'s routed-expert-streaming row still read "CPU keep-quant towers only" after W0c made the seam take the slot arm on `is_cpu()` OR `host_memory_is_device_addressable()` and W0d made the load-time fit refusal drop those towers from its bound. `AGENTS.md` routes a feature-surface change to that page, and the change that moved the surface did not write it. `scripts/check-doc-checkpoint.py` said so, on commit `939755f99` of `row/ENG-EXPERT-STREAM-DEVICE-W0`: a measurement was appended to `.agents/benchmark-record.md` with no `docs/FEATURES.md` edit beside it. The PAGE is FIXED IN FLOW while repairing the fresh review of [#1377](https://github.com/mudler/vllm.cpp/pull/1377): the row now names both the device arm ([#1124](https://github.com/mudler/vllm.cpp/issues/1124)) and the residency condition ([#1378](https://github.com/mudler/vllm.cpp/issues/1378)), and the observability detail it displaced to stay inside the 220-character cell budget is stated at `docs/USAGE.md:4598-4620`. The GATE is NOT fixed and needs a decision: the checker walks a range one COMMIT at a time, so once a commit is published on a branch that may not be force-pushed, no later commit can make it green, and `scripts/agent-preflight.sh` keeps reporting `doc-checkpoint range` red on this branch until it merges. The squashed commit that lands on `main` carries both paths and passes. Changing the walk is checker semantics and needs its own row, spec and red-first evidence per `AGENTS.md` "Changing the rules or a checker", so it is not folded in here. Spec [`expert-stream-device-slots.md`](specs/expert-stream-device-slots.md) | bug | | [#1386](https://github.com/mudler/vllm.cpp/issues/1386) | — | `tools/bench/gpu_clock_state.py`'s `QUERY_FIELDS` collects nine fields and **none of them is thermal and none is electrical**, so the driver's own `SwThermalSlowdown` label can never be checked against a die reading on any window this helper has ever recorded. The measured consequence is that the nine windows of 2026-08-19 cannot distinguish a load transition from a thermal excursion. The concrete evidence is `clock-c1-r1.samples.json` in `/mnt/nas_share/rc/q38bf16/out/bench-20260819T035148Z/`: ours c1 r1 dips five times on the same period at the same `utilization.gpu = 96` — 48.83 s / 2177 MHz, 80.60 s / 2320 MHz, 109.28 s / 2210 MHz, 137.98 s / 2359 MHz, 166.07 s / 2268 MHz — and **two of those five carry `0x0000000000000000`**, no throttle bit at all (2210 and 2359), while three carry `0x20`. The 2210 MHz unlabelled dip is deeper than two of the three labelled ones, so the driver labels comparable excursions inconsistently and the bit alone cannot decide it. What would settle it: add `temperature.gpu` and `power.draw` to `QUERY_FIELDS`. That changes the clock-record schema, so it owes its own row and spec. Split out of [#1354](https://github.com/mudler/vllm.cpp/issues/1354) and owed under `## Owed` in [lease-clock-pinning.md](specs/lease-clock-pinning.md) | gap | +| [#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 | | [#1426](https://github.com/mudler/vllm.cpp/issues/1426) | `LTX25-GUIDED-VIDEO` | The LTX-2.5 x0 model picks its DiT forward with one ternary on `im.on_device` (`src/vllm/multimodal/ltx2_video.cpp`), and **no test on a box without an accelerator enters the DEVICE branch**, so dropping its `perturbations` argument is undetectable. MEASURED, both mutations built clean at `compile_err=0` and both restored byte-for-byte with the target rebuilt: **D10** (drop `p` from the device branch alone, 1 insertion / 1 deletion) is **GREEN**, `test_ltx2_video` exit 0, 88 cases / 2755 assertions; **D11** (drop it from BOTH branches, 3 insertions / 2 deletions) is **RED**, exit 1, 2 cases / 3 assertions, failing `CHECK(MaxAbsDiffOf(t.video_first_perturbed, t.video_first_cond) > 1e-6 * latent_span)`, the same for `video_first_modality`, and `CHECK(n.video_first_perturbed != n.video_first_cond)`. D11 also has to be written with `(void)p;` to be a RESULT at all, because the naive form leaves the lambda parameter unused and `-Werror=unused-parameter` refuses the build - a mutation that fails to build reads exactly like a passing test. This is a missing RUNNER and not a missing test: `Ltx2VideoEngine::Load` refuses `device != 0` in three places (`ltx2_video.cpp:749-758`, `:779-788`, and the queue-creation check below them) unless a non-CPU platform resolves, registers a backend and does not decline `kLtx2VideoFamily`, and `.github/workflows/ci.yml` has no GPU runner. A SYNTHETIC accelerator is DEFERRED rather than impossible, and the PARTS ALREADY EXIST. `vt::RegisterOp(OpId, DeviceType, void*)` is public (`include/vt/op_provider.h:127`) and is exactly how `src/vt/cpu/cpu_ltx2.cpp:201-205` installs the CPU glue table, so a test can read that pointer back with `vt::GetOp(OpId::kLtx2, kCPU)` and register it for a second device type; `RegisterReferenceTier` loops EVERY OpId (`src/vt/op_provider.cpp:795-807`) and `MaybeInstallReferenceTier` copies `src->fn` verbatim (`:204-225`), which covers the `vt::` ops the staging and the forward need; `Glue()` `VT_CHECK`-refuses BY NAME and never returns nullptr (`src/vllm/model_executor/models/ltx2_device.cpp:81-86`); and the closed 7-value `DeviceType` enum (`include/vt/device.h:16-28`) is handled the way `tests/vllm/multimodal/test_diffusion_device_seam.cpp` ALREADY handles it FOR THIS FAMILY — a SEPARATE EXECUTABLE (`:19-23`, `tests/CMakeLists.txt:317`) that registers a unified-memory `FakeXpuBackend` (`:60-73`), a platform in the XPU and CUDA slots (`:116-121`) and a `device = 1` `kLtx2VideoFamily` load (`:127-133`) — which leaves `test_ltx2_video`'s 88-case baseline byte-for-byte untouched. One design fact goes with that: `Ltx2DeviceKernelsAvailable` reads `vt::OpRegistered`, which EXCLUDES the reference tier deliberately (`op_provider.cpp:677-701`, "a FALLBACK, not a native kernel"), so the glue table needs the direct `RegisterOp` and the tier serves the rest. A fake also DOES prove link B as this row defines it — the x0 model reaching `Ltx2DitForwardDevice(..., p)`, which would turn D10 RED; what it cannot prove is that a GPU runs it, which is the hardware half the leased run below closes. WHAT IT COSTS is that new executable plus a fixture DiT staged through the fake backend (`tests/vllm/multimodal/ltx2_video_fixture.h` already writes a complete reduced model). TWO gates on the route were checked and pass by construction: `vt::RegisterBackend(kXPU, ...)` writes slot 0 (`src/vt/backend.cpp:81-84`) so `:800`'s `TryGetBackend(Device{kXPU, 0})` resolves for `params.device = 1`, and `vt::CreateQueue(Device)` takes its `entry.backend->CreateQueue()` arm (`backend.cpp:146-153`) whose `q.device == device` assertion the fake satisfies. TWO are NOT verified and NOBODY HAS BUILT IT: whether `Ltx2StreamDitToDevice` (`:847`) stages the checkpoint onto that queue, and whether the bf16 render that follows (`:812`) completes through the reference tier. PLAUSIBLE, NOT PROVEN — recorded in those terms because the first draft of this row asserted the opposite three facts as measured. What IS gated: `Ltx2DitForwardDevice` honours all four perturbation fields on the CPU backend, nine mutations RED (`tests/vllm/models/test_ltx2_device.cpp`). Closes on one leased `one_stage --device cuda` run at the model's own guider defaults recording `dit_forwards == 4 * dit_evaluations`. Same shape as [#1131](https://github.com/mudler/vllm.cpp/issues/1131). Listed under `## Owed` in [`ltx25-guided-video.md`](specs/ltx25-guided-video.md) §12.8 link B | bug | | [#1396](https://github.com/mudler/vllm.cpp/issues/1396) | `MODEL-FP8-BLOCK-WEIGHT` | Eight comments and one spec `## Owed` asserted that the block-wise FP8 forward wiring was still owed, after it landed. All were last written by M3 (`09597106e`, MODEL-FP8-BLOCK-WEIGHT) and none was revisited when M4 (`281b4bc76`, the linear method and the dense forward that reads the weight), M5 (`489a9a4c0`, the mainloop-scaled CUTLASS kernel) and M6 (`836c13c35`, the merged `gate_up` and QKV) landed. The tree disagreed with them in two directions: `qwen3_5.cpp` calls the block-scaled GEMMs at TEN sites -- the attention `o_proj`; q/k/v once on the split path, where one `project` lambda serves q, k and v alike, and once more on the merged one; the GDN `in_proj`'s `qkv` and `z` halves; the GDN `out_proj` in each of its three block arms (`GdnBlock`, `GdnBlockPagedMixedSpec`, `GdnBlockPaged`); and the dense MLP's merged `gate_up` and its `down_proj` -- which is eight `MatmulFp8BlockScaledD`, one `MatmulFp8BlockMergedD` and one `Fp8BlockGateUpSwiGLUD`, so `NOTHING CONSUMES THIS YET` was false; and `RefuseUnrunnableQwen3_5DenseFp8Block` had already been narrowed by M4 to refuse a DEVICE with no block-scaled GEMM rather than the weight, so `Deleted by M5` and `Milestone M5 removes this` were false in the other direction -- M5 narrowed the refusal to a CUDA arch outside `VT_CUTLASS_FP8_ARCHS` (12.0a, 12.1a) rather than deleting it. Both available readings of the stale text are costly: re-implement a delivered milestone, or refuse to use a working arm. FIXED IN FLOW, comment and spec text only, no behaviour change. The change deliberately does NOT narrow the real debt, which is unchanged and recorded in [`vt-matmul-fp8-block-cuda.md`](specs/vt-matmul-fp8-block-cuda.md) `## Owed`: the CUDA kernel has never executed on hardware and there is no token gate against `Qwen/Qwen3.8-27B-FP8`. No gate can hold this class -- no checker here compares a comment against the code it annotates -- so the reviewer's check is those ten call sites and the body of `ModelRegistry::Prepare` | bug | | [#1411](https://github.com/mudler/vllm.cpp/issues/1411) | `FIX-FP8-BLOCKWISE-REFUSAL` | Eight more places say the block-wise FP8 CUDA kernel is owed, three of them LIVE refusal messages a user reads. Found by the fresh review of [#1399](https://github.com/mudler/vllm.cpp/pull/1399), the pull request for [#1396](https://github.com/mudler/vllm.cpp/issues/1396), and outside the eight comment anchors that change corrected. The three message strings tell a user on an unsupported CUDA arch to wait for milestone M5, which LANDED at `489a9a4c0`: `layers/quantization/fp8_block_quant.cpp:187-189`, and `models/dense_fp8_block_gemm.h:200-202` in the `MatmulFp8BlockScaledD` guard and `:428-430` in the shared `CheckFp8BlockMergedActivation` helper that BOTH merged entry points call, so that second message also fires on the merged `gate_up` path and not only on QKV. The honest reason is that the arch is outside the `VT_CUTLASS_FP8_ARCHS` cell (12.0a, 12.1a) M5 registered the kernel for, not that no kernel exists; the same build runs the checkpoint on `sm_120a`, on `sm_121a` and on CPU. The five comments are `include/vt/ops.h:1637,1658`, `src/vt/cpu/cpu_ops.cpp:668`, `include/vt/merged_gemm.h:92`, `src/vllm/model_executor/models/qwen3_5_dense_weights.cpp:604` and `tests/vllm/model_executor/models/test_fp8_block_linear.cpp:267`. ONE of those needs care rather than a rewrite: `include/vt/ops.h:1637` names a measurement that is still GENUINELY OWED, so its defect is tense alone. A NINTH place, the `RefuseUnrunnableQwen3_5DenseFp8Block` comment block in `qwen3_5_dense_weights.cpp`, WAS fixed in flow by #1396 and is no longer owed -- no line anchor, because #1396 edits that block itself: it was an entry-point count, saying `MatmulFp8BlockScaledD` reads the weights at each of the ten projections when it reads EIGHT, because `qwen3_5.cpp` reaches those ten through THREE entry points -- `Fp8BlockGateUpSwiGLUD` is the only reader of `gate_proj_fp8_block` and `up_proj_fp8_block`, and `MatmulFp8BlockMergedD` reads q/k/v as one operand. NOT FIXED IN FLOW for the remaining eight, deliberately: #1396 is a comment-and-spec-text change and editing eight more places across `vt`, the quantization layer and a test would widen it past what its fresh review read. The debt itself is NOT narrowed by this issue -- the CUDA kernel has never executed on hardware and there is no token gate against `Qwen/Qwen3.8-27B-FP8`, recorded in [`vt-matmul-fp8-block-cuda.md`](specs/vt-matmul-fp8-block-cuda.md) `## Owed`. No gate can hold this class, because no checker here compares a comment or a message string against the code it annotates. `docs/FEATURES.md` carries the same staleness and is tracked separately by [#1275](https://github.com/mudler/vllm.cpp/issues/1275). Owed under `## Owed` in [`fp8-blockwise-refusal.md`](specs/fp8-blockwise-refusal.md) | bug | diff --git a/.agents/quantization-matrix.md b/.agents/quantization-matrix.md index 7cf40cb35..01259c3d0 100644 --- a/.agents/quantization-matrix.md +++ b/.agents/quantization-matrix.md @@ -35,7 +35,7 @@ otherwise it remains `PARTIAL` or `INVENTORIED` even if parsing works. | `QUANT-GGUF-CIQ-GEMM` | Compute-in-quant GEMM: activation quant (Q8_0/Q8_K) + per-type vec_dot dispatch for Q8_0/Q4_K/Q5_K/Q6_K/Q3_K/Q4_0; portable C++ tier, then x86/Arm SIMD + repack tiers. **G1-G4 landed** — the portable tier-0 path is complete, gated at the OP level, and **ROUTED end to end**: `vt::MatmulBT` dispatches a block-dtype weight to `kMatmulBTQuant`, keep-quant is the production DEFAULT wherever that op is registered, and the six routed encodings compute in quant with **no token movement**. **G6 (2026-07-23)** added the Arm **i8mm mmla `nrc==2` tier** for q8_0/q4_0/q4_K/q6_K (q3_K/q5_K have no upstream mmla → stay portable), 2x2-tiled into `kMatmulBTQuant` at even M,N: op-level q4_K **7–8.4×** / q6_K **3.8–4.5×** / q8_0 ~1.2× over portable, e2e prefill +8.4 % on the q8_0-dominant bench file (1.44× behind llama.cpp), tokens byte-identical. **G7 (2026-07-23)** added q8_0 **repack-at-load** (the `q8_0_4x8` tier `ggml_repack_get_optimal_repack_type` picks on NEON+i8mm): the loader repacks each q8_0 weight once into the `block_q8_0x4` interleave and `kMatmulBTQuant` dispatches a pre-shuffled i8mm gemm/gemv with no per-block register shuffles — op-level q8_0 **3.7–5.9×** over the mmla tier, **E2E prefill 1.92× same-binary → 223.8 t/s vs llama.cpp pp128 177.3 = at/beyond parity** (was ~1.5× behind), decode at parity, tokens byte-identical. **CPU prefill parity reached; the prefill-lever search is closed** (remaining gap = peak RSS 1.39×, loader-bound). G5 (x86) + G8 open. **The FRESH op-dispatch profile this row owed is DONE (2026-08-06, dgx aarch64, `main` @`dfd29060`, same bench file; see `.agents/benchmark-record.md` 'FRESH op-dispatch profile'), and it does NOT support starting G5 next:** `QuantRepackMatmul` is 5.06 % of prefill and 15.99 % of decode on aarch64 where the i8mm tier already landed. The profile re-ranks the CPU levers to (1) threadpool synchronisation at 47 % of decode (`ThreadReady`+`PollForWork`+`Barrier`; M=1 cannot amortise the barrier) and (2) CPU paged attention at ~39 % of prefill, of which 20.68 % is a per-ELEMENT dtype switch in the attention dot loop (`cpu_paged_attn.cpp:29` called from `:143`), the same defect class E1 already removed from the elementwise GEMM. G5 stays a real x86 gap worth closing for x86 users, but it is not the top lever, and the x86 box is VOID for timing so it cannot be speed-gated here | llama.cpp `ggml/src/ggml-cpu/ggml-cpu.c:211-406` traits table, `ggml-cpu/quants.c:174-860` generic vec_dot, `arch/{x86,arm}/quants.c`, `ggml-cpu/repack.cpp:4153-4830` at `237ad9b96` | G1: [block dtypes + geometry](../src/vt/dtype.cpp#L32), [quant traits table](../src/vt/cpu/cpu_quant_traits.cpp#L1), [shared block decoders](../src/vt/cpu/cpu_quant_dequant.cpp#L1), [op surface](../include/vt/quant.h#L1). G2: [activation quant + scratch sizing](../src/vt/cpu/cpu_quant_act.cpp#L1) (`quantize_row_q8_0/q8_K`). G3: [the six generic vec_dot](../src/vt/cpu/cpu_quant_dot.cpp#L1), [block-struct mirror](../src/vt/cpu/cpu_quant_blocks.h#L1), [`kMatmulBTQuant` quantized path + composite fallback](../src/vt/cpu/cpu_quant_gemm.cpp#L1). G4: [the routing point](../src/vt/ops.cpp#L158) — `vt::MatmulBT` sends a block-dtype `b` to `MatmulBTQuant` and is otherwise unchanged, which is sufficient because every model matmul helper already routes an `nk=true` weight there ([qwen3_5.cpp:1067](../src/vllm/model_executor/models/qwen3_5.cpp#L1067)); plus [the default flip + `expand_nk`](../src/vllm/model_executor/model_loader/gguf_keep_quant.cpp#L95) and [the untransposed expand path](../src/vllm/model_executor/models/qwen3_5_gguf_weights.cpp#L194). G6: [Arm i8mm mmla tier](../src/vt/cpu/cpu_quant_dot_arm.cpp#L1) (q8_0/q4_0/q4_K/q6_K `vmmlaq_s32`, HWCAP2_I8MM-probed, `VT_CPU_QUANT_MMLA` defeat) + [2x2 tile in kMatmulBTQuant](../src/vt/cpu/cpu_quant_gemm.cpp#L85), per-file `+i8mm` in CMakeLists | [G1 traits cross-check + fallback units](../tests/vt/test_ops_quant_traits.cpp#L1) — 8 cases / 5,615 assertions green (was 5,694; its composite case now covers Q8_K alone because the six weight types legitimately no longer take that path): vt geometry vs the reader's `GgmlTraits` vs ggml-common.h arithmetic all agree, and the composite equals the loader dequant byte-for-byte. [G2/G3 units](../tests/vt/test_ops_quant_dot.cpp#L1) — 16 cases / 78,052 assertions green: every `vec_dot` gated against an INDEPENDENT f64 dequantize-then-dot reference (tolerance relative to the dot's L1 magnitude, actual agreement ~1e-6) over nblocks {1,2,3,5,7,16} incl. single-block and odd multiples; ragged K throws at every layer; upstream thresholds ported unwidened (test-quantize-fns:17-28, test-backend-ops:4277 NMSE ≤ 5e-4 at M {1,4,32,512} × N {1,7,16}); bit-exact run-to-run and across threads 1/2/4; byte-exact encoder gate pins the rounding rules; 14-mutant battery, 13 caught, the 1 uncaught mutant provably unreachable. [dequant units](../tests/vllm/test_gguf_dequant.cpp#L25) still green after the decoder move. DGX (G2/G3 re-confirmed, each gate STANDALONE, goldens md5 identical before/after): clean CUDA `-Werror` build 0 warnings + full regression set UNCHANGED (27B 235/235, 35B 315/315, Coder 6/6, Qwen3-dense 16/16 on both 0.6B and 4B, OPT 6/6, DeepSeek-V2 8/8) + `test_qwen36_gguf_engine` 28/28 with 16/16 tokens on both APEX files + the new CPU units green on aarch64 with identical counts. **G4 (2026-07-22):** `test_qwen36_gguf_engine` PASSES STANDALONE on a CPU-only dgx build (where keep-quant is live) — 2/2 cases, 16/16 greedy tokens on APEX-Compact AND APEX-Balanced vs the same-file llama.cpp oracle, exercising 5 of the 6 routed encodings end to end; the CUDA regression set is UNCHANGED (27B 235/235, 35B 315/315, Coder 6/6, Qwen3-dense 16/16, OPT 6/6, DeepSeek-V2 8/8, gguf 28/28 incl. `VT_CPU_REF=1`), goldens md5 identical. **Binding CPU A/B** (idle dgx aarch64, one flock, same binary, 3 reps, `Qwen3.5-2B-UD-Q8_K_XL`): decode 2.216 -> 7.650 t/s (**3.45x**), prefill 5.149 -> 21.44 t/s (**4.16x**), peak RSS 7.428 -> 6.401 GiB, output tokens byte-identical across the pre-G4, post-G4 and `VT_CPU_REF=1` arms. Still **3.38x / 8.20x / 2.29x behind llama.cpp** — the projected 9-17x did NOT hold because 60 % of that file's weight bytes are `f16`, which no block encoding covers. **That gap is now CLOSED by `KERNEL-GEMM-CPU-ELEM`** (2026-07-22, same box/recipe/binary discipline): the elementwise kernel went 18-24 -> 69-351 GFLOP/s bit-exactly, taking the CPU position to **decode 1.03x behind (parity within 3.1 %) and prefill 2.34x behind**, tokens unchanged (same md5). Its measured NEGATIVE re-ranks G5-G8 once more: M-blocking the elementwise GEMM bought 1.63x op-level and **0.0 % end-to-end**, so the 95.37 % `kMatmul` attribution these G-rows were ranked against is STALE and a FRESH op-dispatch profile is owed before G5/G6/G7 are started. **G6 (2026-07-23):** [Arm i8mm mmla tier](../src/vt/cpu/cpu_quant_dot_arm.cpp#L1) landed against the refreshed profile (kMatmulBTQuant 50 % + kMatmul 16 % + kMatmulBT 14 % = 80 % of prefill). [test_ops_quant_dot G6 cross-check](../tests/vt/test_ops_quant_dot.cpp#L1) — 19 cases / **78,162** assertions on dgx aarch64: q8_0/q4_0 mmla **BIT-IDENTICAL** to the portable/scalar tier (`vmlaq_f32` non-fused under `-ffp-contract=off`), q4_K/q6_K within NMSE ≤ 5e-4, mmla GEMM bit-identical across threads 1/2/4/20. `test_qwen36_gguf_engine` 2/2 · 16/16 on both APEX files with mmla live (q8_0/q4_K/q6_K at prefill), bench-file token md5 `d235db12f2cd304007530286a1755c95` byte-identical across mmla-OFF/ON/`VT_CPU_REF=1`. Op-level portable→i8mm: q8_0 ~1.2×, q6_K 3.8–4.5×, q4_K 7–8.4×; e2e prefill same-binary 1.084× (1.56×→1.44× behind llama.cpp pp128). CUDA `-Werror` 0-warn, regression set UNCHANGED (27B 235/235, 35B 315/315, Coder 138, Qwen3-dense 184, OPT, DeepSeek-V2 223), goldens untouched. **G7 (2026-07-23):** [q8_0 repack transform](../src/vt/cpu/cpu_quant_repack.cpp#L1) + [i8mm repack gemm/gemv](../src/vt/cpu/cpu_quant_repack_arm.cpp#L1) dispatched from [`kMatmulBTQuant`](../src/vt/cpu/cpu_quant_gemm.cpp#L151) on `b.repacked`; loader repacks via [`OwnGgufQuantBlocks`](../src/vllm/model_executor/models/qwen3_5_gguf_weights.cpp#L20) under `GgufLoadPolicy::quant_repack`, flag carried to the kernel through [`ResidentWeight`](../src/vllm/model_executor/models/qwen3_5.cpp#L702). [test_ops_quant_repack](../tests/vt/test_ops_quant_repack.cpp#L1) — 305 assertions on dgx aarch64: repacked gemm/gemv `memcmp`-equal to plain `kMatmulBTQuant` across decode/leftover/prefill, f32+bf16 out, strided activations, threads 1/2/4/20; interleave matches `make_block_q8_0x4` byte-for-byte (110 on x86, numeric skip). `test_qwen36_gguf_engine` STANDALONE 2/2·16/16 on APEX Compact+Balanced (repack live), token md5 `d235db12f2cd304007530286a1755c95` byte-identical across repack-ON/OFF/`VT_CPU_REF=1`. Binding dgx aarch64 (idle, one flock, 6 interleaved reps): op-level q8_0 3.7–5.9× (518→2401/583→3456/514→1902 GFLOP/s); E2E prefill **1.92×** (1096→572 ms), **223.8 t/s vs llama.cpp pp128 177.3 = 1.26× at/beyond parity**, decode at parity, RSS unchanged; fresh profile q8_0 GEMM 55%→~21%, prefill-lever search CLOSED. CUDA `-Werror` 0-warn, regression set UNCHANGED (27B 235/235, 35B 315/315, Coder 6/6, Qwen3-dense 16/16, OPT 6/6, DeepSeek-V2 8/8, Llama 16/16), goldens content-hash identical . **P0 REGRESSION FOUND + FIXED (2026-08-06, `CLAIM-QUANT-GGUF-CIQ-GROUPED-DTYPE`):** the GROUPED provider `MatmulBTQuantGroupedKernel` was f32-ONLY — it advanced a `float*` by `act.stride[0]` and declared the row `kF32` whatever `act.dtype` said, so a bf16/f16 activation was mis-strode 2x AND mis-decoded. Every prior caller/test passed f32; qwen3_5 W3b `KqGrouped` (bf16 act, `b4f5610a`) was the first non-f32 caller, so CPU-only GGUF 35B decode became all-token-0 while the CUDA gate stayed byte-exact (CUDA always honoured `act.dtype`). Fixed at [`cpu_quant_gemm.cpp:220-268`](../src/vt/cpu/cpu_quant_gemm.cpp) (rows addressed by `SizeOf(act.dtype)`/`SizeOf(out.dtype)`; `repacked`/`q8_0_aligned` now propagate onto the per-expert slice — the CIQ-G7 all-zero mode). Gated per activation dtype + bf16-out by 2 NEW cases in [`test_ops_quant_dot.cpp`](../tests/vt/test_ops_quant_dot.cpp) (RED pre-fix on f16+bf16 for all 12 weight encodings, GREEN after; f32 unaffected either way) | [CIQ GEMM leaf](specs/gguf-compute-in-quant-gemm.md) | `ANCHOR-BACKFILL` | `CLAIM-QUANT-GGUF-CIQ-G7-1` | | `QUANT-GGUF-KEEPQ-LOADER` | Keep-quantized GGUF loader: block-resident 2-D matmul weights ([N,K], no transpose), per-tensor routing, `VT_CPU_REF` dequant-oracle switch, bench-branch `7c91a42` merge. **L1+L2+L3 landed** — block residency, the TOTAL per-tensor routing policy and the `VT_CPU_REF` oracle switch all exist and are gated. **Keep-quant is DEFAULT ON since CIQ G4** wherever the running device has a registered `kMatmulBTQuant` (CPU, and since 2026-07-29 also **CUDA** for the Q8_K family via the `KERNEL-QUANT-CIQ-GEMM-CUDA` kCUDA provider — a CUDA runner now keeps k-quant/i-quant blocks COMPRESSED instead of expanding), with `VT_GGUF_KEEP_QUANT=0` as the opt-out. L4 measured; **L5 LANDED** (mmap in-place residency + tied-head sharing + read-once page release) — peak RSS 6.401 -> **3.884 GiB**, 2.29x -> **1.39x** llama.cpp, byte-identical | llama.cpp `src/llama-model-loader.cpp:1047,1385` (file-typed residency), `:1676` + `ggml/src/llama-mmap.cpp:490` (`unmap_fragment`), `ggml/src/ggml-cpu/repack.cpp:4727` (repack-at-load hook) at `237ad9b96` | L1: dense-arch (`qwen35`) GGUF path on main via the registry — [dense GGUF load](../src/vllm/model_executor/models/qwen3_5_dense.cpp#L60), [arch->registered-ID map](../src/vllm/model_executor/models/qwen3_5_gguf_weights.cpp#L212), [F16/BF16 row dequant](../src/vllm/model_executor/model_loader/gguf_dequant.cpp#L61). L2: [block residency `OwnGgufQuantBlocks`](../src/vllm/model_executor/models/qwen3_5_gguf_weights.cpp#L20) — raw ggml blocks into an `OwnedTensor` with a block `vt::DType`, file `[N,K]` orientation, `nk=true`, no transpose; stacked experts split by byte range. L3: [routing policy + `VT_CPU_REF`/`VT_GGUF_KEEP_QUANT`](../src/vllm/model_executor/model_loader/gguf_keep_quant.cpp#L1) (6 roles, no `default:` label so an unrouted role is a `-Werror=switch` build failure) wired at every loader call site via [`OwnMatmulWeight`/`RequireExpand`](../src/vllm/model_executor/models/qwen3_5_gguf_weights.cpp#L185). **Default now device-derived** (CIQ G4): [`GgufQuantComputeAvailable`](../src/vllm/model_executor/model_loader/gguf_keep_quant.cpp#L95) gates it on `vt::OpRegistered(kMatmulBTQuant, CurrentPlatform().device_type())`, and the same condition drives `expand_nk`, which stops transposing a weight that must expand. **`expand_nk` now also covers the GDN split projections** (2026-07-23, `CLAIM-CPU-GDN-ORIENT-1`): a fresh op-dispatch profile found `LoadGdnGguf`'s `in_proj_qkv/z/b/a` + `out_proj` were the ONE expanded weight family still transposed to [K,N] (nk=false → slow `kMatmul`, 17.9 % of prefill); the new [`gdn_expand_nk` field](../src/vllm/model_executor/model_loader/gguf_keep_quant.cpp#L95) + [`MakeGdnProj`](../src/vllm/model_executor/models/qwen3_5_gguf_weights.cpp#L131) keep them [N,K] nk=true (V-head reorder applied first, orthogonal to orientation; `VT_GGUF_GDN_NK=0` A/B opt-out) → M-blocked `kMatmulBT`, same-binary prefill **1.090×** / decode 1.09×, byte-identical (`kMatmul` 72→0 calls in prefill) | [L2/L3 units](../tests/vllm/test_gguf_keep_quant.cpp#L1) — 17 cases / 5,574 assertions green. **Gate 1 (losslessness) proven PER ENCODING**, one case each for Q4_0/Q8_0/Q3_K/Q4_K/Q5_K/Q6_K: resident bytes `memcmp`-equal to the file span and resident-block dequant BYTE-IDENTICAL to the direct-from-file expansion (f32 and bf16), over pseudo-random block bytes constrained only to finite f16 scales; at loader level the kept weight rehydrates to the expanded `[K,N]` bf16 tensor byte for byte, per weight and per expert, on dense and MoE fixtures. **Totality**: the audit hook proves `routed == the file's complete tensor list` on both fixtures, plus 6 roles × 12 encodings × 6 shapes against a LONGHAND expectation (12 keep / 420 expand, so neither outcome is vacuous). **Gate 2 (oracle stability)**: `VT_CPU_REF=1` keeps nothing quantized and every weight is bit-identical to the historical load; on dgx [`test_qwen36_gguf_engine`](../tests/parity/test_qwen36_gguf_engine.cpp#L143) under `VT_CPU_REF=1` is 28/28 assertions, 16/16 tokens on both APEX files — same as without. 10-mutant battery, 10 caught (the expert-slice-offset mutant survived the first pass, exposed a real coverage hole, and drove the MoE fixture). DGX (each gate STANDALONE, production flags, goldens md5 identical before/after `2965ef5772b556d3f3f86fedf4221b2f`): clean CUDA `-Werror` 0 warnings + regression set UNCHANGED (27B 235/235, 35B 315/315, Coder 6/6, Qwen3-dense 16/16 on both, OPT 6/6, DeepSeek-V2 8/8) + gguf units green on aarch64 with identical counts; full CPU ctest 154/154. **RSS at G4 was 6.401 GiB (2.29x); L5 took it to 3.884 GiB (1.39x)** — binding, idle dgx aarch64, same-binary 3-rep A/B: mmap in-place residency (borrow kept q8_0 blocks out of the mapping, refcounted, -0.998 GiB), tied-head sharing (one bf16 vocab matrix for embed+lm_head, -0.946 GiB), read-once page release (MADV_DONTNEED the expanded tensors' file pages, port of llama.cpp `unmap_fragment`, -0.573 GiB). Decode TPOT 41.7 ms UNCHANGED, prefill TTFT +4% (first-touch faults move into the timed window), output md5 `d235db12f2cd304007530286a1755c95` identical across BEFORE/AFTER/ORACLE. Lifetime safety tested explicitly (borrow outlives the GgufFile AND the on-disk file; shared head freed once either order). **L6 (2026-07-23, `CLAIM-QUANT-GGUF-KEEPF16-L6-1`) implemented keep-f16 residency and REFUTED the "remaining gap is the f16 expansion" attribution above.** New `kKeepF16` residency + [`OwnGgufF16`/`OwnGgufKeptSlice`](../src/vllm/model_executor/models/qwen3_5_gguf_weights.cpp#L79) keep F16 matmul weights (+ F16 embed/tied head, one shared f16 vocab matrix via [`OwnedBytes::KeepAlive`](../include/vllm/model_executor/models/owned_bytes.h)) resident as F16, consumed by the elementwise f16 GEMM. Binding A/B: peak RSS 3.884 → **3.832 GiB (−52 MB, RSS-NEUTRAL)** — L5's page-release ALREADY dropped the f16 file pages, so keep-f16 only swaps an anonymous bf16 buffer for equal-size file-backed f16 pages. smaps attribution: keep-f16 file-backed **2.634 GiB ≈ llama.cpp's 2.68 file** (weight residency AT PARITY), anon 1.20 GiB; the **remaining ~1.08 GiB gap is the engine's ANONYMOUS activation/KV workspace, NOT weights** — the real, separate CPU RSS lever. Also regresses prefill (TTFT 577 → ~1000 ms, first-touch faults into the timed window; decode at parity). Tokens byte-identical (md5 `d235db1…`). Ships DEFAULT OFF at L6. **L7 (2026-07-23, `CLAIM-QUANT-GGUF-RSS-L7-1`) REVERSED L6's refutation and CLOSED the CPU RSS gap to 1.01× llama.cpp.** The profile disproved the "workspace" attribution — DevicePool 20 MiB, whole KV 115 MiB, both ≤ llama.cpp. The 1 GiB residual was a q8_0 repack-source DOUBLE-COUNT: on aarch64 the G7 repack COPIES q8_0 into an anonymous buffer while the f16 borrows keep the mapping alive, so the DEAD source blocks stay file-backed. [`OwnGgufQuantBlocks`](../src/vllm/model_executor/models/qwen3_5_gguf_weights.cpp#L20) now `DropSpanResidency`es the repack source (port of llama.cpp `unmap_fragment`), and [`PrefaultBorrowedSpan`](../src/vllm/model_executor/models/qwen3_5_gguf_weights.cpp#L21) faults borrowed weights at load (port of llama.cpp mmap prefetch), removing L6's prefill regression — so [keep-f16 flips DEFAULT ON](../src/vllm/model_executor/model_loader/gguf_keep_quant.cpp#L168) (`VT_GGUF_KEEP_F16=0` opt-out). Binding A/B (idle dgx aarch64, base-vs-L7 same-binary): peak RSS **3.884 → 2.832 GiB = 1.39× → 1.01× llama.cpp** (File 2.632 → 1.629, the released q8_0 source; anon 1.200 unchanged), prefill **1.18× AHEAD** (204 vs pp128 173.2, denominator SUPERSEDED by #1003), decode ~parity (24.4 vs 25.09), tokens BYTE-IDENTICAL (md5 `809f2d0…` base/L7/oracle). **Against our own keep-f16-off arm the default costs about 9% of prefill (224 → 204 t/s) and about 1.4% of decode (TPOT 40.4 → 40.95 ms) for 1.05 GiB, settled 2026-08-17 as a product decision, NOT by the competitor floor.** Anon 1.200 GiB is IRREDUCIBLE (repacked q8_0 1.06 + KV 0.115 + pool 0.02). Regressions UNCHANGED (27B 235/235, 35B 315/315, Coder 6/6, Qwen3-dense 16/16, OPT 6/6, DeepSeek-V2 8/8, Llama 16/16, GGUF engine 28/28); `test_gguf_keep_quant` 36/36 (+1 L7 prefault byte-transparency case, x86+aarch64) | [keep-quant loader leaf](specs/gguf-keep-quant-loader.md) | `ANCHOR-BACKFILL` | `CLAIM-QUANT-GGUF-RSS-L7-1` | | `QUANT-QWEN38-27B-GGUF-ARM` | The `Qwen3.8-27B-Q4_K_M.gguf` arm end to end: tensor accounting, text decode, the multimodal legs, and this ARTIFACT's own tokenizer and chat template. The standing GGUF k-quant requirement for a model whose bf16 arm is already gated ([#915](https://github.com/mudler/vllm.cpp/issues/915)), and the arm `BACKEND-GATE-CUDA-LLAMACPP` in the [backend matrix](backend-matrix.md) is already recorded as blocked on. **Header-verified 2026-08-18** at `unsloth/Qwen3.8-27B-GGUF`@`fe1e2a23d973adb629709749dc4f6756df66ef10`: GGUF v3, arch `qwen35`, 866 tensors, F32 456 / Q4_K 294 / Q6_K 67 / Q5_K 48 / Q8_0 1, data end == file size 17,106,775,008. **Two facts [#821](https://github.com/mudler/vllm.cpp/issues/821) did not record and which change the scope:** `qwen35.block_count = 65` with `qwen35.nextn_predict_layers = 1`, so block 64 is the MTP/`nextn` DRAFTER (`blk.64.nextn.{eh_proj,enorm,hnorm,shared_head_norm}` plus a full-attention block and an FFN) — exactly the 15-tensor difference from the same model's 851-tensor BF16 GGUF, and a loader that reads `block_count` as decoder depth builds a 65-layer model out of a 64-layer checkpoint plus a drafter; and `tokenizer.ggml.padding_token_id = 248055` against 248044 in the BF16 GGUF and `null` in the official HF config, which is why the tokenizer gate belongs to the ARM. NOT blocked on kernels: every dtype this file carries is already computed natively on BOTH tiers. The CUDA tier really has no prefill/decode split (`LaunchGemm` [cuda_quant_dot.cu:1609](../src/vt/cuda/cuda_quant_dot.cu#L1609) sizes its grid `m*n` and the encoding switch at [:1864](../src/vt/cuda/cuda_quant_dot.cu#L1864) never sees `M`); the CPU tier DOES branch on `M` at [cpu_quant_gemm.cpp:190](../src/vt/cpu/cpu_quant_gemm.cpp#L190), which takes the Arm i8mm `mmla` 2x2 tile only for even `M` and `N` and sends decode (`M=1`) to the portable `nrc==1` path. That is a kernel-TIER split, NOT a coverage split -- no dtype gains or loses support at any `M`, both arms end in the same `BlockVecDot` table -- so the conclusion stands and it is a W3 speed fact rather than a W2 gap | llama.cpp `b10451` = `10bf611e5` ([pin](oracles/llama-cpp.md), `gateable = no`) is the arm's ORACLE and its only comparator — at the vLLM pin `555967922` there is no in-tree GGUF reader (`6635279d8` moved it out of tree) and SGLang's alias table does not reach `qwen3_5` ([#979](https://github.com/mudler/vllm.cpp/issues/979)). llama.cpp is never the MIRROR | the single-file GGUF entry [qwen3_5_gguf_weights.cpp:1474](../src/vllm/model_executor/models/qwen3_5_gguf_weights.cpp#L1474); the CUDA quant enum [cuda_quant_dot.cu:700](../src/vt/cuda/cuda_quant_dot.cu#L700) and CPU [cpu_quant_dot.cpp:787](../src/vt/cpu/cpu_quant_dot.cpp#L787) already cover Q4_K/Q5_K/Q6_K, and Q8_0 has its own path [cuda_quant_dot.cu:1659](../src/vt/cuda/cuda_quant_dot.cu#L1659) | none yet — this row is spec-only. NO Qwen3.5 checkpoint-index accounting gate exists anywhere in `tests/`; the templates are [test_nemotron_h_loader.cpp](../tests/vllm/models/test_nemotron_h_loader.cpp) (18,487 tensors, per-scheme composition) and [muse_glimmer_gguf_manifest.inc](../tests/vllm/models/muse_glimmer_gguf_manifest.inc) (header-only committed manifest). The token gate is PENDING on [#857](https://github.com/mudler/vllm.cpp/issues/857), which owes the llama.cpp gateability measurement | [quantized arms of Qwen3.8-27B](specs/qwen38-27b-quant-arms.md) | `READY` | - | -| `QUANT-QWEN38-27B-NVFP4-ARM` | The `unsloth/Qwen3.8-27B-NVFP4` artifact, which is **not what its name says**. **Its pinned revision is GONE:** `a767244d27bd76589a3e3b2ab4e64032c4ebc7af`, the revision [#821](https://github.com/mudler/vllm.cpp/issues/821) names, answers HTTP 404 and `git ls-remote` reports one ref, `refs/heads/main` = `7d6f8d4d72f56b92b3cdbf22f156b90e1bab0108` — the second in-place re-quantization this publisher has done in this family, after `unsloth/Qwen3.6-27B-NVFP4`. So the user-reported load failure on #821 is CORROBORATED at a different revision, never reproduced. At the live revision (header-verified 2026-08-18, 1953 + 15 tensors, `8 + header_len + max(data_offsets[1])` == file size 22,568,192,096) `quantization_config.format` is `mixed-precision`: `group_0` is FP8 W8A8 with **per-CHANNEL** weight scales and **DYNAMIC per-token** activations over `self_attn.(q\|k\|v\|o)_proj`, `linear_attn.(in_proj_qkv\|in_proj_z\|out_proj)`, `lm_head` and `layers.(56..63).mlp.*`; `group_1` is `nvfp4-pack-quantized` W4A4 over the remaining `mlp.*`; plus an 8-bit static `kv_cache_scheme` and an `ignore` list of **303 entries** -- not just the vision tower: 48 x `linear_attn`, `linear_attn.norm`, `linear_attn.in_proj_b` and `linear_attn.in_proj_a` (the GDN layer count), 27 x 4 vision blocks, 2 mergers, and `re:^mtp.*`. That list is what makes the predicate claim provable rather than asserted: `in_proj_a`/`in_proj_b` are IGNORED while `in_proj_qkv`/`in_proj_z`/`out_proj` are `group_0` TARGETS, so a resolver that reads the groups but not the `ignore` list gets the GDN block wrong in both directions. **`*.input_scale` appears ZERO times in the checkpoint.** Four independent blockers, each anchored in the spec: the unconditional `.input_scale` read, a per-channel BF16 `weight_scale` that `ReadF32Scalar` refuses on BOTH count and dtype, no representation for a dynamic per-token activation scheme, and a scheme that is never read from the config at all. The NVFP4 half is the half CLOSEST to working; the FP8 tower is the blocker | vLLM `555967922` is the MIRROR and the primary oracle — it runs this format, so nothing here may diverge from its compressed-tensors semantics. A generic mixed-precision resolver already exists at [modelopt_mixed_precision.h](../src/vllm/model_executor/layers/quantization/modelopt_mixed_precision.h) and **NO production file includes it** -- the only two includes in the tree are its own two tests, and `nemotron_h_weights.cpp` (which this row previously named) includes `nemotron_h.h`, `nemotron_h_loader.h`, `nvfp4_dequant.h` and `vt/unaligned.h` and reads its quant config inline. A 33,575-byte header reachable only from tests is an `AGENTS.md` §"Nothing lands dead" item, so W4 adopting it is a FIRST production wiring that owes reachability evidence, not a reuse of something proven | the failing read [qwen3_5_weights.cpp:457](../src/vllm/model_executor/models/qwen3_5_weights.cpp#L457), its refusal [dense_weight_loaders.h:164](../include/vllm/model_executor/models/dense_weight_loaders.h#L164) (the count check at `:168`, the dtype check at `:172`; #1258 moved `ReadF32Scalar` down 63 lines and `:101` is now an unrelated `try`/`catch` probe), the scalar-only `Fp8Weight` [qwen3_5_weights.h:342](../include/vllm/model_executor/models/qwen3_5_weights.h#L342), the static activation quant [qwen3_5.cpp:3629](../src/vllm/model_executor/models/qwen3_5.cpp#L3629), and the stale predicate [qwen3_5_dense_weights.cpp:698](../src/vllm/model_executor/models/qwen3_5_dense_weights.cpp#L698) (the `.linear_attn.in_proj_` early-false at `:702`) that declares the GDN input projections never quantized — true for the 3.6 unsloth artifact ([hf_snapshot.h:287](../tests/parity/hf_snapshot.h#L287)) and false for this one | none yet — this row is spec-only. Manifest-capture precedent [minimax_h3_nvfp4_manifest.inc](../tests/vllm/models/minimax_h3_nvfp4_manifest.inc), captured the same way this row's numbers were: an HTTP range read of the file's own header. The token gate is PENDING on [#1185](https://github.com/mudler/vllm.cpp/issues/1185), which records the pinned oracle BUILDING inside an `rc` lease and RUNNING A MODEL as untested | [quantized arms of Qwen3.8-27B](specs/qwen38-27b-quant-arms.md) | `READY` | - | +| `QUANT-QWEN38-27B-NVFP4-ARM` | The `unsloth/Qwen3.8-27B-NVFP4` artifact, which is **not what its name says**. **Its pinned revision is GONE:** `a767244d27bd76589a3e3b2ab4e64032c4ebc7af`, the revision [#821](https://github.com/mudler/vllm.cpp/issues/821) names, answers HTTP 404 and `git ls-remote` reports one ref, `refs/heads/main` = `7d6f8d4d72f56b92b3cdbf22f156b90e1bab0108` — the second in-place re-quantization this publisher has done in this family, after `unsloth/Qwen3.6-27B-NVFP4`. So the user-reported load failure on #821 is CORROBORATED at a different revision, never reproduced. At the live revision (header-verified 2026-08-18, 1953 + 15 tensors, `8 + header_len + max(data_offsets[1])` == file size 22,568,192,096) `quantization_config.format` is `mixed-precision`: `group_0` is FP8 W8A8 with **per-CHANNEL** weight scales and **DYNAMIC per-token** activations over `self_attn.(q\|k\|v\|o)_proj`, `linear_attn.(in_proj_qkv\|in_proj_z\|out_proj)`, `lm_head` and `layers.(56..63).mlp.*`; `group_1` is `nvfp4-pack-quantized` W4A4 over the remaining `mlp.*`; plus an 8-bit static `kv_cache_scheme` and an `ignore` list of **303 entries** -- not just the vision tower: 48 x `linear_attn`, `linear_attn.norm`, `linear_attn.in_proj_b` and `linear_attn.in_proj_a` (the GDN layer count), 27 x 4 vision blocks, 2 mergers, and `re:^mtp.*`. That list is what makes the predicate claim provable rather than asserted: `in_proj_a`/`in_proj_b` are IGNORED while `in_proj_qkv`/`in_proj_z`/`out_proj` are `group_0` TARGETS, so a resolver that reads the groups but not the `ignore` list gets the GDN block wrong in both directions. **`*.input_scale` appears ZERO times in the checkpoint.** Four independent blockers, each anchored in the spec: the unconditional `.input_scale` read, a per-channel BF16 `weight_scale` that `ReadF32Scalar` refuses on BOTH count and dtype, no representation for a dynamic per-token activation scheme, and a scheme that is never read from the config at all. The NVFP4 half is the half CLOSEST to working; the FP8 tower is the blocker | vLLM `555967922` is the MIRROR and the primary oracle — it runs this format, so nothing here may diverge from its compressed-tensors semantics. A generic mixed-precision resolver already exists at [modelopt_mixed_precision.h](../src/vllm/model_executor/layers/quantization/modelopt_mixed_precision.h) and **NO production file includes it** -- the only two includes in the tree are its own two tests, and `nemotron_h_weights.cpp` (which this row previously named) includes `nemotron_h.h`, `nemotron_h_loader.h`, `nvfp4_dequant.h` and `vt/unaligned.h` and reads its quant config inline. A 33,575-byte header reachable only from tests is an `AGENTS.md` §"Nothing lands dead" item, so W4 adopting it is a FIRST production wiring that owes reachability evidence, not a reuse of something proven | the failing read [qwen3_5_weights.cpp:457](../src/vllm/model_executor/models/qwen3_5_weights.cpp#L457), its refusal [dense_weight_loaders.h:164](../include/vllm/model_executor/models/dense_weight_loaders.h#L164) (the count check at `:168`, the dtype check at `:172`; #1258 moved `ReadF32Scalar` down 63 lines and `:101` is now an unrelated `try`/`catch` probe), the scalar-only `Fp8Weight` [qwen3_5_weights.h:563](../include/vllm/model_executor/models/qwen3_5_weights.h#L563), the static activation quant [qwen3_5.cpp:3629](../src/vllm/model_executor/models/qwen3_5.cpp#L3629), and the stale predicate [qwen3_5_dense_weights.cpp:698](../src/vllm/model_executor/models/qwen3_5_dense_weights.cpp#L698) (the `.linear_attn.in_proj_` early-false at `:702`) that declares the GDN input projections never quantized — true for the 3.6 unsloth artifact ([hf_snapshot.h:287](../tests/parity/hf_snapshot.h#L287)) and false for this one | none yet — this row is spec-only. Manifest-capture precedent [minimax_h3_nvfp4_manifest.inc](../tests/vllm/models/minimax_h3_nvfp4_manifest.inc), captured the same way this row's numbers were: an HTTP range read of the file's own header. The token gate is PENDING on [#1185](https://github.com/mudler/vllm.cpp/issues/1185), which records the pinned oracle BUILDING inside an `rc` lease and RUNNING A MODEL as untested | [quantized arms of Qwen3.8-27B](specs/qwen38-27b-quant-arms.md) | `READY` | - | | `QUANT-GGUF-PRESETS` | Representative mixed-file gates for every llama.cpp output preset family | llama.cpp `tools/quantize/quantize.cpp:34-74` | only custom APEX mixed files are executable; no general preset dispatch | [APEX gates](../tests/parity/test_qwen36_gguf_engine.cpp#L143) do not prove llama.cpp preset breadth | [coverage spike](specs/quantization-coverage.md); split exact preset IDs before `READY` | `INVENTORIED` | - | ## 1. llama.cpp / GGUF encodings diff --git a/.agents/specs/expert-stream-device-slots.md b/.agents/specs/expert-stream-device-slots.md index 7431f7f8b..ce0055770 100644 --- a/.agents/specs/expert-stream-device-slots.md +++ b/.agents/specs/expert-stream-device-slots.md @@ -10,20 +10,119 @@ platform may read it. ## Now -`ACTIVE`. **W0a through W0e have all run. W0 does what it was built to do — -`--device cuda` now LOADS this checkpoint instead of refusing — and the CUDA arm -still produces no token, because the first forward exhausts the box for a reason -that is not the expert lane -([#1299](https://github.com/mudler/vllm.cpp/issues/1299)). The CPU arm is -measured, reproduced, and replaces the VOID #912 F1 figure.** See `## Evidence`. - -The three gates, reported one result each: +`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. **Is it W0f? The two grounds first offered here could not answer that, and + they are withdrawn.** The CPU arm was re-run on the SAME binary and the + SAME lease and reproduced the recorded answer byte for byte, and the + instrument counted `w0f-alias` calls **0** on that arm. Neither + discriminates. `ResidentWeight` takes an `is_cpu()` early return + (the `is_cpu()` branch of `qwen3_5.cpp`'s `ResidentWeight`) roughly ninety lines above the alias branch (the `host_memory_is_device_addressable()` branch of the same function; no line number, because this change moves it), + so a zero count on the CPU arm is true by construction for every possible + state of W0f, correct or corrupt; and "the CPU arm reproduces its own + reference" constrains only the arm W0f cannot reach. Both show the branch + is platform-gated. Neither separates "the arms' GEMM arithmetic differs" + from "W0f moved a logit". + 1b. **The discriminating experiment, RUN, and it clears W0f.** The only thing + a consumer can notice about the substitution is the pointer, so the + question is whether cuBLASLt answers differently for a 256-aligned HOST + block than for a `cudaMalloc` one. Measured on `thor:gpu0` (NVIDIA Thor + `sm_110`, driver 13020, cuBLASLt 130101), which answers this branch's own + predicate TRUE (`PageableMemoryAccess = 1`, `Integrated = 1`) and is + therefore in the population W0f serves. Six checkpoint shapes + (`embedding_length = 8192`, M = 1, 5, 32) crossed with BOTH formulations + the dense path issues (row-major NN, weight as B; column-major TN, weight + as A) — **12 measurements, `PROBE_FAILURES=0`**. The heuristic returned an + **identical selection 12/12** on a repeated call, 12/12 with the 256 + promise stated explicitly, and 12/12 with it weakened to 16; and + `cublasLtMatmul` over the same bytes gave **bit-exact output 12/12**, zero + differing elements, every status `SUCCESS`. Five DIFFERENT algo + configurations appear across the six shapes, so the instrument + discriminates. The structural reason needs no lease: + `cublasLtMatmulAlgoGetHeuristic` takes no operand pointers, so alignment + reaches it only through a preference this tree never sets. **Identical algo + and bit-exact output, so W0f cannot move a logit.** + 1c. **The GB10 leg RAN too, on the target silicon.** `rc` job + `7c7a05e9-be87-48f4-94ae-1bbe0340f063` on `dgx:gpu0`, 2026-08-19 17:47 UTC, + `NVIDIA GB10 sm_121`, driver 580.173.02, cuBLASLt 130101, the predicate + re-derived in the job's own output (`pageableMemoryAccess=1 integrated=1`). + Same six shapes, same two formulations, **12 measurements, + `PROBE_EXIT=0`, `PROBE_FAILURES=0`**: repeated heuristic call identical + 12/12, unset preference equal to the 256 default 12/12, the promise weakened + to 16 moving nothing 12/12, and `cublasLtMatmul` output **bit-exact 12/12** + (`differing=0`) between a `cudaMalloc` operand and a 256-aligned host block. + At least five distinct algo configurations appear across the twelve, and + they differ from Thor's, so the heuristic was re-resolved rather than + replayed. **The attribution is therefore measured on the silicon the token + gate ran on: the alias does not cause the divergence.** It does NOT say what + does; naming the first operation that differs is carried under `## Owed`. + 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. **No ratio of the two medians is written here or in + `.agents/benchmark-record.md`, deliberately**: it would rest on a token + comparison that FAILED, and a disowned figure written out in digits is how a + number this repository never measured becomes one it is quoted as having + measured. Both medians are above for anyone entitled to the quotient. Second + caution: 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, and read at a stated point.** 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. **The qualifier is part of the +number**: those are the FIRST-FORWARD totals, taken at the point re-homing +plateaus, call **1361**. The counters are per CALL and there is no memo on the +alias branch, so they keep growing at roughly 70 GiB per decode step; quoted +without the qualifier the same figure is a traffic count and not a residency +measurement. The residency claim it supports is corroborated independently by +peak RSS **97.75 GiB** against a load that previously reached 61.20 GiB and then +exhausted the box. That is the whole difference between zero decode steps and 32. + +W0a HAS RUN as a standalone probe on `dgx:gpu0` and answered +`W0A_VERDICT=PAGEABLE_OK`, so the stop condition that would have returned this +row `NEEDS_DECISION` never fired. The completed decode above corroborates it from +the other end: the load succeeds only when +`host_memory_is_device_addressable()` answers true, because W0d's conditional +refusal is keyed on it. + +The three gates, reported one result each. Every row states what it read at +`95883dcae` where W0f moved it, so the two runs in `## Evidence` are not confused +for one: | Gate | Result | |---|---| -| **G0-CORRECT** | **NO CUDA SIDE.** The CUDA arm emits zero tokens, so there is nothing to compare. The CPU side is byte-identical across four runs and two slot counts (32 ids, listed in `## Evidence`), which is the strongest half of the comparison that this hardware allows today. | -| **G0-LIVE** | **PASS on CPU, NOT REACHED on CUDA.** CPU: `steps=32`, `forced=0`, decode-phase `exhausted` delta **0** at both 4000 and 8000 slots. CUDA: the store BUILDS and prints its banner, and no step boundary is ever reached, so there is no snapshot pair to difference. | -| **G0-SPEED** | **CPU only, no ratio.** Steady decode **11.05 s/token**, which is rep 2's median over 29 samples (min 9.43, max 13.25) at 4000 slots; rep 1's median is 11.22, so the two reps agree within 1.5%. No CUDA number exists, so no ratio is reported and none may be inferred. | +| **G0-CORRECT** | **FAIL, and it now has a CUDA side to fail on.** At source `95883dcae` the entry read "NO CUDA SIDE", because that arm emitted zero tokens. With W0f it emits 32, and they diverge from the CPU arm at step 7 on a MEASURED near-tie: the CPU arm's own runner-up is the token CUDA emitted, 1.4% behind, and one step later the margin is 0.1%. The CPU side remains byte-identical across four runs and two slot counts (32 ids, listed in `## Evidence`). The alias is measured ON GB10 not to be the cause; WHAT is remains open under `## Owed`. | +| **G0-LIVE** | **PASS on both arms.** CPU: `steps=32`, `forced=0`, decode-phase `exhausted` delta **0** at both 4000 and 8000 slots. CUDA with W0f: `steps=32`, decode-phase `exhausted` delta **0**, peak RSS 97.75 GiB of a 119.631 GiB box, swap untouched, container exit 0. At `95883dcae` this read "NOT REACHED on CUDA", because no step boundary was ever crossed. | +| **G0-SPEED** | **VOID, and no ratio is published.** The CPU denominator is measured: steady decode **11.05 s/token**, rep 2's median over 29 samples (min 9.43, max 13.25) at 4000 slots, rep 1's median 11.22, the two reps agreeing within 1.5%. A CUDA number exists now and is recorded in `../benchmark-record.md` for the record only, because this row's own stop condition VOIDS a speed result behind a failing correctness gate. No ratio may be inferred from the two. | What that means precisely, because both "W0 landed" and "W0 failed" would misstate it: @@ -55,9 +154,10 @@ misstate it: branch is the 1.1875 GiB allocation #1123 died on), reads the tower in place when the cache cannot serve a slice on a non-CPU platform, and refuses by name if a claimed tower ever reaches device staging. Gated in the new - `test_expert_stream_device_slot` (5 cases / 38 assertions) over a fake - staging, host-addressable platform, because no CPU tier can register a real - one. + `test_expert_stream_device_slot` (5 cases / **45** assertions after W0f + re-stated its "served normally" case; the count is read off the binary's own + last `test cases:` line, not off what this row expected) over a fake staging, + host-addressable platform, because no CPU tier can register a real one. * **W0d — the conditional refusal.** The fit bound gained a `StreamedExpertLane` input; the loader fills it when the platform stages, can read host slots, the resolved model's factory declares @@ -114,6 +214,22 @@ misstate it: assertions and exit status 0, RED against the current file at 17 cases with 1 failed and 130 assertions with 1 failed and exit status 1, and both files restored byte-identical by sha256. +* **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](https://github.com/mudler/vllm.cpp/issues/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` + (**12 cases / 71 assertions**) over the same fake staging, host-addressable + platform, plus one defect the work uncovered and fixed in flow + ([#1320](https://github.com/mudler/vllm.cpp/issues/1320)) and three a fresh + review found: the `d_dev_f32` disjunct in the dense host-mirror release, the + refusal that fired above the device-copy memo, and the source-page release that + repeated once per forward step. The count is quoted from the binary's own last + `test cases:` line, not from the number this row expected. * **W0a — the probe. RUN, and it answered the question W0b rests on.** On `dgx:gpu0` inside an `rc` lease: `cudaDevAttrPageableMemoryAccess = 1` and `cudaDevAttrIntegrated = 1`, which is exactly the pair `CudaPlatform` @@ -132,18 +248,29 @@ misstate it: `ConcurrentManagedAccess` — are not carried here, because the verdict turns on the two that are and inventing the other two would be worse than omitting them. -* **W0e — the measurement.** RAN, on one `rc hold` on `dgx:gpu0`. It produced a - reproduced CPU figure, a CUDA load that works, and a CUDA arm that generates - nothing. `## Evidence` has all of it. - -Today `--device cuda` on `Qwen3.8-2.4T-A95B UD-Q1_0` **loads**, which is what -W0 was for and is new. It then dies in the first forward, and the cause is the -DENSE half of the model rather than the expert lane: the non-expert weights are -resident twice on a unified part, and a 0.15 GiB slot arena fails in exactly the -place an 18.55 GiB one does. That is -[#1299](https://github.com/mudler/vllm.cpp/issues/1299), listed under `## Owed`. -The developer's target remains a GPU figure, and this row cannot produce one -until #1299 moves. +* **W0e — the measurement. RAN TWICE, on two trees, and the pair is the + result.** The first run, on one `rc hold` on `dgx:gpu0` at source `95883dcae` + (W0f's parent), produced a reproduced CPU figure of **11.05 s/token at 4000 + slots**, a CUDA load that works, and a CUDA arm that generates nothing. The + second, at source `9c783a8be` with W0f in the tree, produced **32/32 CUDA + decode steps**: G0-LIVE **PASS**, G0-CORRECT **FAIL** on a measured near-tie, + G0-SPEED **VOID** by this row's own stop condition. `## Evidence` and + [`../benchmark-record.md`](../benchmark-record.md) carry both, in two sections + that must not be mixed: 11.05 s/token is the standing CPU number and the second + run's CPU column is a same-lease control rather than a second attempt at it. + +Today `--device cuda` on `Qwen3.8-2.4T-A95B UD-Q1_0` **loads and decodes**. The +load is what W0b-W0d were for; at `95883dcae` it then died in the first forward, +and the cause was the DENSE half of the model rather than the expert lane -- the +non-expert weights were resident twice on a unified part, and a 0.15 GiB slot +arena failed in exactly the place an 18.55 GiB one did. That was +[#1299](https://github.com/mudler/vllm.cpp/issues/1299), and W0f is the fix: +32/32 steps at peak RSS 97.75 GiB. + +**The developer's target is a GPU FIGURE, and this row still cannot publish +one.** G0-CORRECT fails on the step-7 near-tie, so G0-SPEED is VOID by this row's +own stop condition. A decode number exists in `../benchmark-record.md` for the +record; it is not a result and no ratio may be inferred from it. ## Scope @@ -328,6 +455,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 | @@ -349,6 +477,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 @@ -520,7 +649,12 @@ which detokenize to " Paris. Paris is a city located in the northern part of France, on the Seine River. It is the largest city in France and is known for its iconic", `finish_reason=length`, `completion_tokens=32`. -### The CUDA arm: it loads, and it does not generate +### The CUDA arm at `95883dcae`: it loads, and it does not generate + +**Everything in this subsection was read on W0f's PARENT tree and is kept as +measured.** It is the diagnosis W0f was built from, not a claim about the tree +this spec describes; the subsection after it is the same harness re-run with W0f +and it reaches 32 decode steps. **The load is the new thing and it works.** `--device cuda` on this checkpoint used to refuse ([#1123](https://github.com/mudler/vllm.cpp/issues/1123)); W0d's @@ -589,8 +723,66 @@ this: the GDN V-head reorder makes `attn_qkv` and `ssm_out` pays it twice and cannot. That spec's own sentence — "Whoever takes this needs BOTH: the streaming lane for the ~330 GiB of experts, and a transformed-weight path that does not expand" — is exactly this result, and W0 delivered the first -half. Filed as [#1299](https://github.com/mudler/vllm.cpp/issues/1299) and -listed under `## Owed`. +half. Filed as [#1299](https://github.com/mudler/vllm.cpp/issues/1299), and +**FIXED as W0f**, which is the next subsection. The prediction that the fix would +need a transformed-weight path turned out to be one option rather than the only +one: not expanding, and not paying for the expansion twice, are different repairs +and W0f is the second. + +### The CUDA arm with W0f at `9c783a8be`: it decodes + +Same harness, same lease shape, same box, same prompt ids `760,6511,314,9338,369`, +streaming ON at 4000 slots, greedy, 32 tokens, page cache dropped between arms. +`../benchmark-record.md` carries the full entry including the two VOID attempts +that preceded it, which were void because the build target relinked nothing. + +| Observable | CUDA | CPU (same-lease control) | +|---|---|---| +| load | 266.330 s | 253.504 s | +| RSS after load | 61.20 GiB | 62.45 GiB | +| decode steps | **32** | 32 | +| decode-phase `exhausted` delta | **0** | **0** | +| peak RSS | **97.75 GiB** | 92.19 GiB | +| swap used at peak | 0 | 0 | +| container exit | 0 | 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. Read at the point re-homing plateaus, call **1361** of +the first forward: **60.793 GiB** re-homed into an aligned host block and then +aliased, **~9.2 GiB** declined as misaligned GGUF borrows and still staged, and +0.02 GiB aliased in place. The qualifier is part of the number, because the +counters are per CALL with no memo on the alias branch and keep growing at +roughly 70 GiB per decode step. On the CPU arm the same counter reads **0 +calls**, which is a live control that the branch is platform-gated and not an +argument that it is. + +**And the ids diverge at step 7.** `...,264,3177,7172,...` on CPU against +`...,264,3177,303,...` on CUDA. An instrumented CPU run printing the top-2 logits +per step shows the CPU arm's own runner-up at that step is `303` — exactly what +CUDA emitted — 0.264709 logits behind, 1.4 % of the winner; two steps later the +margin is 0.022802, about 0.1 %. The declared gate fails and the wave stops. + +**The alias is excluded as the cause, on the target silicon.** `rc` job +`7c7a05e9-be87-48f4-94ae-1bbe0340f063` on `dgx:gpu0` (`NVIDIA GB10 sm_121`, +driver 580.173.02, cuBLASLt 130101, the predicate re-derived in the job's own +output as `pageableMemoryAccess=1 integrated=1`) ran six checkpoint shapes +crossed with both cuBLASLt formulations the dense path issues: 12 measurements, +`PROBE_EXIT=0`, `PROBE_FAILURES=0`. A repeated heuristic call is identical 12/12; +the tree's unset preference equals the documented 256 default 12/12; weakening +the promise to 16 moves nothing 12/12; and `cublasLtMatmul` output is bit-exact +between a `cudaMalloc` operand and a 256-aligned host block 12/12, +`differing=0`, every status `SUCCESS`. At least five distinct algorithm +configurations appear across the twelve and they differ from the earlier +`thor:gpu0` leg's, so the heuristic was re-resolved rather than replayed and the +instrument discriminates. The structural reason needs no lease: +`cublasLtMatmulAlgoGetHeuristic` takes no operand pointers, so alignment reaches +it only through a preference this tree never sets. + +**Excluding one cause is not identifying another.** What the divergence IS +remains unmeasured and is carried under `## Owed` with its next traceable step +named. ### What was running beside the measurement @@ -676,6 +868,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. @@ -750,7 +969,13 @@ 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 mutation of W0b's CUDA leg.** `CudaPlatform::host_memory_is_device_addressable` compiles only in a CUDA build, so no CPU-tier gate can invert it. The bullet in `## Now` promised this line and the table did not carry it, which is fixed here. | **Half discharged by W0e and stated as half.** The lane engaged on a real `--device cuda` run — the `[expert-stream] ON` banner printed and the #1123 refusal did not fire — and neither happens unless the probed predicate returned true on the actual CUDA platform, so the leg is now proven REACHED and proven to answer true on a GB10. What is still owed is the negative: a mutation that makes it answer false and shows a gate go red. That needs a CUDA build with a test target, and W0e built with `-DVLLM_CPP_BUILD_TESTS=OFF` because the lease was for the measurement. | | **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 CUDA arm loads and then exhausts the box in its first forward, so this row still has no GPU number.** [#1299](https://github.com/mudler/vllm.cpp/issues/1299). The non-expert weights are resident twice on a unified part, once as the host-side `OwnedTensor` and once as the `ResidentWeight` device staging copy, and about 50 GiB of that is the bf16 expansion the GDN V-head reorder forces on `attn_qkv` and `ssm_out`. | Not fixable inside this row's scope, and measured rather than inferred: a 0.15 GiB arena fails where an 18.55 GiB one does, and the growth is `RssAnon` while `RssFile` stays flat. The fix is a transformed-weight path that does not expand, or a staging path that releases the host copy — either is its own row with its own spec. W1 and W2 are unaffected: they are about WHERE a slice lives, and this is about the dense remainder beside it. | +| ~~**The CUDA arm loads and then exhausts the box in its first forward, so this row still has no GPU number.**~~ **CLOSED by W0f**, 2026-08-19 ([#1299](https://github.com/mudler/vllm.cpp/issues/1299)): the non-expert weights were resident twice on a unified part, and `ResidentWeight` now aliases the host bytes where `host_memory_is_device_addressable()`. The same checkpoint reaches **32/32 decode steps** at peak RSS 97.75 GiB. | Kept as a line rather than deleted, because the entry recorded a diagnosis as well as a debt and the diagnosis held: a 0.15 GiB arena failed where an 18.55 GiB one did, and the growth was `RssAnon` while `RssFile` stayed flat, which is what pointed at the dense remainder rather than at the lane. What it got wrong was the scope call -- "not fixable inside this row's scope" -- and W0f fixing it in one branch is the correction. What is NOT closed is the GPU NUMBER: G0-CORRECT fails, so G0-SPEED stays VOID and no rate is published. | | ~~**The CPU arm's streaming decode figure is still VOID.**~~ **CLOSED by W0e**, 2026-08-18: streaming-ON decode on a live cache is **11.05 s/token** steady at 4000 slots, rep 2's median with rep 1 at 11.22, and the decode-phase `exhausted` delta is 0 in the same run. See `## Evidence`. | Kept as a line rather than deleted because `docs/BENCHMARKS.md:8` still carries the parent row's VOID (#912 F1) text for `ENG-EXPERT-STREAM`, which owns that row's own re-measure. This row measured its own denominator and is no longer waiting on one. | +| **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. | +| **WHAT the step-7 divergence IS.** The alias is EXCLUDED as its cause, on the target silicon: the algo-identity probe ran on `dgx:gpu0` as well as on `thor:gpu0` (`rc` job `7c7a05e9-be87-48f4-94ae-1bbe0340f063`, `NVIDIA GB10 sm_121`, cuBLASLt 130101, 12/12 identical selection, 12/12 bit-exact output, `PROBE_FAILURES=0`), so this entry is no longer about the probe. | Excluding one cause is not identifying another, and nothing here may present it as one. The standing hypothesis is that the two arms run genuinely different GEMM kernels and the greedy path is a coin flip, and that is not measured. The next traceable step needs no new instrument beyond a two-arm dump of the step-7 forward: name the FIRST tensor whose values differ between the CPU and CUDA arms at that step, and the operation that produced it. It needs the same `dgx:gpu0` lease as W0e. Until it runs, G0-CORRECT stays FAILING and G0-SPEED stays VOID. | +| **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`). **WHY IT FAULTS IS UNMEASURED.** An earlier draft of this row wrote "almost certainly because the pointer it is handed there is not host memory on that arm", and that is a hypothesis, not a reading: nothing printed the pointer, nothing asked `cudaPointerGetAttributes` about it, and no fault address was recorded. In a change whose central risk is handing device kernels host pointers, a segfault whose cause was guessed at is exactly the finding that must not be dismissed — so it is recorded as unmeasured rather than as explained. The CPU arm's margin is enough to establish the near-tie (`303` is its own runner-up), and the next lease should print `cudaPointerGetAttributes(logits)` in that callback before anything else. | +| **No CI gate reaches the alias branch through a production entry point.** `test_expert_stream_wiring` enters `Qwen3_5Model::Forward` and the reachability mutation reds it, but it runs on the **CPU** device, where `ResidentWeight` returns at the `is_cpu()` early return roughly ninety lines above the alias branch. In CI the branch is reached only through `detail::StageWeightForTest`, a test-only seam. | Deliberate, and this is the entry `## Nothing lands dead` requires for it. The branch is selected by `needs_weight_staging() && host_memory_is_device_addressable()`, and no CPU tier can register a platform that answers both — a real one exists on exactly one machine this project can reach. The device evidence is real and is the stronger of the two (the W0e run entered the branch **43,501 times** through `Qwen3_5Model::Forward` on `dgx:gpu0`); it is simply not repeatable in CI. Closing this means either a GPU CI lane on a probed-capable part, or a production entry point that a fake staging platform can drive end to end. It is owned by `ENG-EXPERT-STREAM-DEVICE` and tracked by [#1299](https://github.com/mudler/vllm.cpp/issues/1299) until either lands, and that pair is named in the landing commit body and the pull request body as well as here, because `## Nothing lands dead` requires all three and the spec alone is not the disclosure. | +| **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/.agents/specs/model-fp8-block-weight.md b/.agents/specs/model-fp8-block-weight.md index f84a91c97..de330be81 100644 --- a/.agents/specs/model-fp8-block-weight.md +++ b/.agents/specs/model-fp8-block-weight.md @@ -117,7 +117,7 @@ struct Fp8BlockWeight { }; ``` -`Fp8Weight` (`qwen3_5_weights.h:342-355`) is three host floats — `weight_scale`, +`Fp8Weight` (`qwen3_5_weights.h:563-576`) is three host floats — `weight_scale`, `input_scale`, and the `alpha = input_scale * weight_scale` folded at load. A block scheme has **no `input_scale` at all** (the activation scheme is dynamic; the target checkpoint ships zero such tensors) and its weight scale is a 2-D diff --git a/.agents/specs/qwen38-27b-quant-arms.md b/.agents/specs/qwen38-27b-quant-arms.md index bbfe6b01e..16ee699eb 100644 --- a/.agents/specs/qwen38-27b-quant-arms.md +++ b/.agents/specs/qwen38-27b-quant-arms.md @@ -351,7 +351,7 @@ unquantized. (`:147-163`) names this exact case: "A per-output-channel `[out] BF16` scale passed at two bytes an element and was read as one float built from the first two entries." So even after (1) is fixed, `weight_scale` BF16 `[10240,1]` fails the count check first. `Fp8Weight` - (`include/vllm/model_executor/models/qwen3_5_weights.h:342-355`) is three host + (`include/vllm/model_executor/models/qwen3_5_weights.h:563-576`) is three host floats with **no tensor-valued scale slot**, so this is a type change, not a read fix. diff --git a/docs/BENCHMARKS.md b/docs/BENCHMARKS.md index deed595fa..7f9b0bc2e 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` 370 GiB, `--device cuda` vs `cpu`, one GB10, greedy 32 tokens, **4000 expert-stream slots**, 2026-08-18, [record](../.agents/benchmark-record.md) | CPU **MEASURED**, replaces #912 F1's VOID: **11.05 s/token** steady, rep 2's median, rep 1 11.22 (1.5%), live cache. 8000 slots is 3.6-4.1x slower. CUDA: loads, dies in forward 1, 0 steps (#1299). W0a `PAGEABLE_OK` | correct answer; no oracle runs this | +| **Device expert slices (`ENG-EXPERT-STREAM-DEVICE`, [#1124](https://github.com/mudler/vllm.cpp/issues/1124), [#1299](https://github.com/mudler/vllm.cpp/issues/1299))** | Qwen3.8-2.4T-A95B `UD-Q1_0` 370 GiB on `cuda` vs `cpu`, one GB10, greedy 32 tokens, **4000 expert-stream slots** (8000 is 3.6-4.1x slower), [record](../.agents/benchmark-record.md) | CPU **11.05 s/token** steady, replacing #912 F1's VOID. CUDA **G0-LIVE PASS**: 32/32 steps, `exhausted` delta 0, peak RSS 97.75 GiB. **G0-CORRECT FAIL**, cause NOT identified (not the alias: GB10). **G0-SPEED VOID** | 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/docs/ENVIRONMENT.md b/docs/ENVIRONMENT.md index 8880dff65..70d652a4e 100644 --- a/docs/ENVIRONMENT.md +++ b/docs/ENVIRONMENT.md @@ -210,6 +210,7 @@ portable/reference path. In normal operation leave them unset. | `VT_DFLASH_ATTN_KEYLANE` | off (CUDA) | `=1` selects the one-key-per-lane block-attention form. **MEASURED NEGATIVE and not a tuning knob:** 28.90 s/step against the per-key warp kernel's 18.73 on the same binary (sm_110, MiniMax-H3 512x512/33f, seq 3224), 54% slower, because giving each lane a whole K row makes every K load 32-way scattered. Kept only because it is the experiment that located the real constraint | | `VT_ADOPT_DEVICE_BYTES` | on (only acts where the backend advertises host-addressable device memory — Vulkan today) | After a weight is uploaded, re-point its host buffer AT the device allocation instead of keeping a second copy. On a unified box the two copies come out of the same RAM. MEASURED on GB10, Qwen3.6-27B bf16 (50.89 GiB): with the mirror the process reaches **VmRSS 100.759 GiB** and drives the machine to MemAvailable 13.85 / MemFree 1.13 GiB of 119.6 GiB before it has even finished allocating; without it the same load completes at **VmHWM 53.413 GiB**. Qwen3-4B: **16.392 -> 9.607 GiB**. Vulkan allocation is byte-identical either way. `0` is the same-binary A/B back to the two-copy behaviour, and it moves ONLY this lever: the `VT_LOAD_DIRECT_UPLOAD` post-upload source-page release still runs at `0` and on a backend without host-addressable device memory, so the two knobs stay independent. It is an adoption, not a release — the bytes survive at the device address and every reader sees them — so tokens are unchanged either way (`test_opt_paged_engine` on Vulkan is 6/6 token-exact, 96/96, both arms). No effect on CUDA/CPU/Metal, whose backends do not advertise the property | +| `VT_QWEN35_ALIAS_HOST_WEIGHTS` | on (Qwen3.5/3.6 family; acts only where the platform advertises `host_memory_is_device_addressable()` — a GB10-class integrated part today) | Hand a dense weight's HOST bytes straight to the device kernel instead of allocating a device copy and uploading into it. On a part whose kernels can dereference host storage that copy buys nothing and costs a second full resident copy of every dense weight out of the same RAM. MEASURED on GB10, `Qwen3.8-2.4T-A95B UD-Q1_0` (369.97 GiB) on `--device cuda`: with the second copy the load completes at 61.20 GiB resident and then exhausts the 119.631 GiB box inside its FIRST forward, **zero decode steps over seven attempts** (issue #1299); with the alias the same run reaches **32/32 decode steps at peak RSS 97.75 GiB**. The instrument counts 60.793 GiB of dense weight aliased rather than duplicated (first-forward totals, at call 1361, where re-homing plateaus) against ~9.2 GiB of misaligned GGUF borrows that decline and still stage. `0` is the same-binary A/B back to the staging behaviour, and it is not only a bisect lever: `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, so a decode regression has to be separable from the workload. The substitution is otherwise indistinguishable — the aliased pointer is re-homed to 256 bytes, cuBLASLt's own `CUBLASLT_MATMUL_PREF_MIN_ALIGNMENT_A_BYTES` default, which dominates every explicit pointer gate in the CUDA kernels (the strictest asks 32). A DISCRETE device answers the predicate false and is byte-for-byte unchanged. No effect on CPU, Vulkan, Metal or XPU, whose platforms do not advertise the property | | `VT_LOAD_DIRECT_UPLOAD` | on | Load a weight the device consumes VERBATIM by VIEWING the safetensors mmap (`OwnedBytes::Borrow`, keep-alive on the mapping) instead of copying it into an owned host buffer first, so the device upload reads the file mapping and the load moves those bytes ONCE rather than twice. Only whole-range same-size copies qualify — a transpose, a dtype conversion, a dequant, a concatenation or a load-time repack always takes the copy path, and the helper re-checks `numel * sizeof(dtype) == span` and fails closed to the copy on any mismatch. `0` is the same-binary A/B back to copy-then-upload. Bytes are identical either way, so tokens are unchanged. MEASURED on GB10, Qwen3.6-27B bf16 (50.098 GiB), Vulkan, same binary both arms: the weight-load phase goes **19.27 -> 12.48 s warm** (1.54x) and **52.62 -> 32.75 s cold** (1.61x), load-and-one-token **30.39 -> 22.47 s** warm and **62.98 -> 55.60 s** cold, with every ON leg beating every OFF leg. Total bytes MOVED **100.196 -> 81.260 GiB**: the host materialization pass drops **50.098 -> 31.162 GiB** while the 50.098 GiB device upload is unchanged (the model still has to be uploaded once). 37.8% of this checkpoint qualifies; the rest is merged (qkv, gate_up) or transposed (lm_head) at load and correctly still copies | | `VT_LOAD_STATS` | off | `=1` prints one line per load phase (mmap+header, weights) with its wall time, plus the bytes the load MOVED: `host_copy` (source bytes materialized into an owned host buffer), `borrowed` (source bytes viewed in place by the direct-upload path) and `device_upload` (bytes copied host to device). Diagnostic only; it changes no numerics. Issue #150 | | `VT_VULKAN_ALLOC_STATS` | off | `=1` prints a device-memory line on every 1 GiB high-water crossing and a summary at exit: live buffer count, bytes REQUESTED by the caller, bytes COMMITTED by the driver (`VkMemoryRequirements::size`), peak live bytes, and the process/system context (`VmRSS`, `VmHWM`, `MemAvailable`, `Cached`) read from `/proc`. On a unified-memory device the Vulkan heap IS system RAM, so separating "the backend allocated it", "the process allocated it some other way" and "it is page cache" is the whole of a memory attribution. Diagnostic only; it changes no numerics. The counters themselves are always maintained (one relaxed atomic per allocation) and are readable from a test through `vt::vulkan::DeviceAllocStatsSnapshot()`. Vulkan-only | diff --git a/docs/FEATURES.md b/docs/FEATURES.md index 7e50aa603..c4cd02266 100644 --- a/docs/FEATURES.md +++ b/docs/FEATURES.md @@ -61,7 +61,7 @@ are our reading of their documented behavior, not measurements. | Scratch allocator keyed by device (two backends, one process) | ✅ since [#516](https://github.com/mudler/vllm.cpp/issues/516); a pool is bound to one backend and refuses any other, and a backend with no registered platform is refused rather than given another's residency cap | ✅ device is field 0 of the allocation handle | ✅ | ✅ | | Automatic memory sizing (no hand-tuned budget) | ☐ hand-typed block count | ☐ percent, hand-tuned | ☐ | ◐ | | Memory cap with a pre-flight error instead of an OOM | ☐ | ◐ KV pool only | ◐ | ☐ | -| Routed-expert weight streaming from disk | ◐ default OFF (`VT_MOE_EXPERT_STREAM=1`), keep-quant/keep-f16 towers (#1378); bounded slot cache; refuses an unfittable slice by name. c1-c4 capacity, not throughput. CPU, plus a host-readable staging device (#1124) | ☐ blanket `cpu_offload_gb`, not expert-granular | ☐ | ◐ mmap only | +| Routed-expert weight streaming from disk | ◐ default OFF (`VT_MOE_EXPERT_STREAM=1`), keep-quant/keep-f16 towers (#1378); bounded slot cache; refuses unfittable slices by name. c1-c4 capacity, not throughput. CPU; staging device DECODES, token gate FAILS (#1299) | ☐ blanket `cpu_offload_gb`, not expert-granular | ☐ | ◐ mmap only | ## Quantization and weight formats diff --git a/docs/STATUS.md b/docs/STATUS.md index 3492f4de5..85aa52503 100644 --- a/docs/STATUS.md +++ b/docs/STATUS.md @@ -164,7 +164,7 @@ token-for-token correctness against the pinned oracle. | LoRA / multi-LoRA adapters | In progress (W1+W2 CPU runtime), not yet usable end-to-end | Highest-demand missing feature. W1+W2 landed the punica CPU brick (shrink/expand, `-1`-slot skip) plus packed adapters and the wrapped layer family (merged qkv/gate_up slices, TP slicing, embedding and logits LoRA). CPU-gated vs double references, RED-first: `test_punica_cpu` 8/8 + `test_lora_layers` 16/16 (4,647 assertions); the fully-sharded (S-LoRA) apply REFUSES at tp_size>1 instead of returning a partial delta. W3-W7 (mapping, adapter load, manager, endpoints, GPU kernels + model gate) are in `.agents/specs/lora-adapter.md`. No model can be served with an adapter yet | | Safetensors loading | Supported | Both gate models plus every registered dense/MoE family | | GGUF loading (F32/F16/BF16/Q4_0/Q8_0/Q2_K/Q3_K/Q4_K/Q5_K/Q6_K/IQ2_XXS/IQ3_XXS/IQ2_S/MXFP4/NVFP4) | Supported; compute-in-quant (keep-quant) on CPU AND now CUDA for the six K-block encodings PLUS Q2_K/IQ2_XXS/IQ3_XXS (DeepSeek-V4 W8, 2026-07-29 - the FIRST CUDA keep-quant GGUF k-quant GEMM `KERNEL-QUANT-CIQ-GEMM-CUDA`, MMVQ-style dequant-in-kernel, GB10-gated 92401/92401 vs the CPU oracle, so a CUDA runner keeps blocks compressed and dots them on the GPU instead of the ARM cores); **NVFP4 now COMPUTES IN FP4 on CUDA for the dense-MLP and full-attention projections (2026-07-29, `CLAIM-GGUF-NVFP4-COMPUTE`), no longer materialize-only** | Weights in six block encodings stay compressed from file to matmul on CPU (no BF16 expansion). NVFP4 (ggml type 40) DEQUANTIZES, including the per-tensor (per-expert) `.scale` sidecar the container keeps outside the blocks; gated BIT-EXACT against the compressed-tensors NVFP4 path on real Qwen3.6-27B bytes from both containers. **It no longer expands to bf16 on CUDA:** an NVFP4 matmul/expert weight is REPACKED at load into the same (`weight_packed [N,K/2]`, `weight_scale [N,K/16]`) operand pair the compressed-tensors path produces - a pure byte permutation, gated BYTE-IDENTICAL against that container - and the existing `vt::MatmulNvfp4*` kernels run on it, so no new kernel exists and no numerics are re-derived. Covers the dense MLP + full-attention q/k/v/o and the MoE shared/routed experts; the GDN `in_proj_*` family and `ssm_out` still expand (the V-head reorder rewrites their layout) and a CPU build still expands everything - the documented `part` subset. **MEASURED GB10 (2026-07-29), same-binary A/B, 2 reps/arm:** peak RSS **50.8 -> 25.7 GiB**, load-and-generate **1:58 -> 0:41**; the 256 moved projections cost 35 840 MiB expanded vs 10 080 fp4-resident (3.56x). **The safetensors-sibling divergence CLOSES:** the fp4 arm is token-IDENTICAL over 24 greedy tokens where the same binary's bf16 arm diverges at index 4. REPORTED, not gated: the containers are not the same model - the GGUF NVFP4-quantizes 192 GDN `in_proj` tensors the safetensors keeps BF16 (mean rel. weight error ~0.18) and their activation global scales differ, so identity is not guaranteed and a cross-container throughput arm is invalid. SACRED gates unmoved: `test_qwen27_paged_engine` 235/235, `test_qwen36_paged_engine` 315/315. **The MoE (35B) stacked-expert arm is HARDWARE-GATED too (2026-07-29)**: the real 35B A3B NVFP4 GGUF loads and generates through the fp4 path, its 120 routed-expert stacks x 256 experts repack to the modelopt safetensors' operands with ZERO differing bytes over 840 sampled (tensor, expert) slabs, and all 840 per-expert `.scale[e]` are bit-identical to that expert's `weight_scale_2` - the scale INDEXING, mutation-proved against a `scales[0]`-for-all and an expert-0-slab-for-all mutant. Same-binary A/B: peak RSS 68.5 -> 22.7 GiB (3.01x), load-and-generate 1:51.9 -> 0:28.8, tokens IDENTICAL (the 35B routed experts run the W4A16 grouped GEMM in both arms). Recorded OPEN: this case's 24-token greedy stream is NOT run-to-run stable (1 of 3 `use_a16` and 1 of 4 reference runs differed), so the binding results are the weight-level byte identity and the residency audit, not token-exactness; `test_qwen36_paged_engine` is token-exact at ITS engine params, so the instability belongs to this case's configuration and attributing it is owed work. It also FIXED a latent defect the MoE arm made reachable: the two fp4 fused MoE blocks issued the router GEMM assuming the safetensors `[K,N]` gate layout and threw `matmul: inner dims mismatch` on the GGUF's `[N,K]`; `MoeRouterLogits` now branches on `nk` (inert for safetensors, SACRED unmoved). **Q2_K (id 10) + IQ2_XXS (id 16) DEQUANTIZE (2026-07-29, `CLAIM-DSV4-GGUF-LOADER`):** the ~2-bit types the single-Spark `DeepSeek-V4-Flash-GGUF UD-IQ2_XXS`/`UD-Q2_K_XL` vehicles use, ported 1:1 from llama.cpp `ggml-quants.c` (`iq2xxs_grid` codebook + signs; Q2_K nibble sub-scale/min), unit-gated on hand-derived bytes (`test_gguf_dequant` 15/15). Dequant-only (no vec_dot -> expand-bf16). A V4-GGUF model cannot RUN yet: the name map (tensor-manifest-blocked) + the V4 forward (W3-W8) remain. **Multi-shard split GGUF READING (2026-08-03, `CLAIM-GGUF-SPLIT-SHARDS`):** `GgufFile::Open` now transparently stitches llama.cpp `gguf-split` shards (`...-00001-of-00003.gguf`) — every shard mmap'd, tensor tables merged, KV metadata from shard `00001`, sibling mappings kept alive by the primary so keep-quant mmap-borrows stay valid across shards (`OwnsSpan` is shard-aware); `VT_GGUF_NO_SPLIT=1` opts out; unit-gated (`test_gguf` split-merge / no-split / count-mismatch cases, 33/33 local). This unblocks the real 3-shard `unsloth/DeepSeek-V4-Flash-0731 UD-IQ2_M` (~91 GiB), whose layout is the NATIVE `deepseek4` arch — per-block `ffn_gate_tid2eid` hash tables (hash layers 0/1/2) + `hc_*` MHC + DSA compressor/indexer are all PRESENT (name-map 1328/1328), `vocab_size` derives from `token_embd` — NOT a standard llama.cpp conversion, so no loader-layout change is owed. It now loads THROUGH 1324/1328 tensors; the sole remaining gap is 4 routed-expert slabs quantized with IQ2_S (id 22, ×2) + MXFP4 (id 39, ×2) — encodings we have GGUF block traits for but no keep-quant vec_dot, so they hit the expand→dequant path which lacks them. Expanding those 4 expert tensors to bf16 would add ~17 GiB (~106 GiB total → OOM-reboot risk), so the memory-safe fix is an IQ2_S+MXFP4 keep-quant kernel (CPU dequant dispatch + `iq2s_grid` + CUDA `DotSuperblock`), spec'd as the next brick **IQ2_S (id 22) + MXFP4 (id 39) DEQUANTIZE + KEEP-QUANT on CPU (2026-08-03, `CLAIM-DSV4-UDIQ2M-QUANT`, off-GPU):** the extra per-tensor "dynamic" encodings the `unsloth/DeepSeek-V4-Flash-GGUF UD-IQ2_M` checkpoint mixes into its last routed-expert slabs (IQ2_S `ffn_gate/up` dotting Q8_K, MXFP4 `ffn_down` dotting Q8_0) — ported 1:1 from llama.cpp `ggml-quants.c` @ 237ad9b96 (`iq2s_grid` 1024-entry codebook + DIRECT sign bytes; MXFP4 `kvalues_mxfp4` + `e8m0_to_fp32_half` micro-scaling, distinct from the compressed-tensors `E8M0ToF32` NVFP4 path). CPU dequant + keep-quant `vec_dot`, unit-gated on hand-derived golden bytes (`test_gguf_dequant` 17/17), an INDEPENDENT f64 dequant-then-dot + GEMM NMSE (`test_ops_quant_dot` 19/19), and keep-quant routing (`test_gguf_keep_quant` 37/37) — all CPU-green, so UD-IQ2_M's four previously-`unsupported ggml type 22/39` slabs now load COMPRESSED (no ~17 GiB bf16 expansion that OOM-reboots the box). CUDA: the IQ2_S device `DotSuperblock` is wired into the Q8_K grouped-MoE GEMM and now **CUDA-BUILT + LINKED on GB10 (sm_121a, CUDA 13.0, `-Werror`, 2026-08-03 integration)** — it compiles clean and the merged binary links; MXFP4's device dot (`DotMXFP4`) is written but NOT wired (Q8_0-activation needs a separate 32-block GEMM) so it is marked `[[maybe_unused]]` to keep the ready math without tripping nvcc #177-D, and on GPU MXFP4 CPU-fallbacks like Q4_0/Q8_0. The V4-GGUF forward + a real UD-IQ2_M GPU load/coherence run are owed | -| Larger-than-pool GGUF on an accelerator (`ENG-EXPERT-STREAM-DEVICE`, [#1124](https://github.com/mudler/vllm.cpp/issues/1124)) | W0a-W0e all RUN. It now **loads** on `--device cuda` and does **not** generate: forward 1 exhausts the box, 0 steps ([#1299](https://github.com/mudler/vllm.cpp/issues/1299)). CPU serves, 11.05 s/token at 4000 slots | W0 removed the #1123 refusal: `*_exps` from the HOST slot store on a PROBED `host_memory_is_device_addressable()`, keep-quant/keep-f16 only (#1378). The DENSE half blocks it now: 0.15 GiB dies where 18.55 GiB does | +| Larger-than-pool GGUF on an accelerator (`ENG-EXPERT-STREAM-DEVICE`, [#1124](https://github.com/mudler/vllm.cpp/issues/1124), [#1299](https://github.com/mudler/vllm.cpp/issues/1299)) | **W0b-W0f in; `--device cuda` DECODES it on a GB10.** G0-LIVE PASS (32/32 steps, `exhausted` delta 0, peak RSS 97.75 GiB); G0-CORRECT FAIL, cause NOT identified, not the alias (GB10); G0-SPEED VOID. CPU 11.05 s/token | A 369.96 GiB GGUF refused against a 119.631 GiB pool (#1123), then exhausted the box in its first forward (#1299). W0 serves `*_exps` from host slots, keep-quant/keep-f16 only (#1378); W0f ends the DENSE double residency | | AWQ / GPTQ quantization | W0 spike + W1 CPU INT4 dequant primitive; not yet loadable end to end | INT4 unpack+dequant-to-bf16 for BOTH community formats, mirroring vLLM 1:1 (AWQ reverse-order `awq_triton.py`; GPTQ `qdq_4.cuh` with zero_offset v1/v2 + act-order g_idx). Unit-gated RED-first (hand-computed known bytes + double-precision roundtrip). NOT wired to a loader, no GPU Marlin compute, no model run yet: config recognizer (W2), Marlin GPU GEMM riding the vendored NVFP4 Marlin (W4), CPU e2e (W3), GPTQ 8/2/3-bit (W5) and MoE (W6) are named next bricks. See [.agents/specs/awq-gptq-quant.md](../.agents/specs/awq-gptq-quant.md) | | MXFP4 (compressed-tensors `mxfp4-pack-quantized`) | Compute PROVEN (#38); GQA-swap ON (#49); decode-graph+gate_up FUSION default-ON. `VT_MARLIN_DENSE` DEFAULT-ON (`KERNEL-MARLIN-DENSE-EXEC`): dense marlin 48-CTA byte-faithful (32B 0.000, 263/263), binding beats #51 every axis (c1 1.020, c8 0.969, mem 2.63x). **`QUANT-CT-MXFP4-FINAL-STACK` TERMINAL — both last levers exhausted: num_splits cap `VT_FA2_NSPLITS_CAP` gated-OFF (c1-only, self-corrects@c8; 32B strict char-identical); glue folds via `vt::FusedChain`; `FLASH-AUDIT` #68: c8 flash gap +12.5us/call is occupancy/L2-bound; `-use_fast_math` TRIED, REGRESSES flash (168.8→189.8), rejected. c1 1.020x PASS, c2-c8 0.962-0.969.** state.md | Shared with DeepSeek-V4-Flash + Kimi-K3 MXFP4 paths. CPU E8M0 dequant 5/5·1142. GPU W4A4 + MoE-expert e2e later | | CPU backend vs llama.cpp | 20-core Arm at floor, **denominator SUPERSEDED** (fork `237ad9b96`, owed a re-take vs stock `b10451`, #1003); RPi5/A76 below floor `GATING`; **x86_64 open on every axis** | Pi: AAPCS64 beats SDOT 3.66-5.08%; llama.cpp 2.17x pf / 1.53x dec faster (0.461x/0.653x); RSS -24.2% vs stock `b9892`; 64-tok byte-exact. BF16 GEMM open. x86_64 first measured 2026-08-11 (#433): peak RSS 1.0022x = hairline OPEN GAP (6.33 MB against us), throughput pending a quiet host, `G5` load-discipline gate failing, quant path portable-tier only (CIQ `G5`) | diff --git a/docs/USAGE.md b/docs/USAGE.md index 527dcacc2..d67115403 100644 --- a/docs/USAGE.md +++ b/docs/USAGE.md @@ -4546,7 +4546,8 @@ VT_MOE_EXPERT_STREAM_SLOTS=4000 \ ### Which device can serve it `--device cpu` serves this today, and that is the arm every published number for -this checkpoint was measured on. +this checkpoint was measured on. `--device cuda` now decodes it on a probed +integrated part; see the six limits below before you rely on that. `--device cuda` refuses at load, by design, when the weights cannot be staged into device memory (issue @@ -4560,7 +4561,19 @@ staged at all — their slices are read from the host slot store in place — so what has to fit is the NON-expert remainder plus the slot arena rather than the whole file. -Four limits, stated plainly rather than left to be discovered. +The lane alone was not enough to produce a token. With it on, the checkpoint +loaded on `--device cuda` and then exhausted the machine inside its first +forward — zero decode steps, seven attempts, every one identical (issue +[#1299](https://github.com/mudler/vllm.cpp/issues/1299)) — because the DENSE +weights were resident twice: once as the host buffer and once as the device +staging copy, which on a part where device memory IS host memory comes out of the +same RAM. `VT_QWEN35_ALIAS_HOST_WEIGHTS` (default **on**, `docs/ENVIRONMENT.md`) +removes the second copy by handing the kernels the host bytes directly, and it is +what makes the CUDA arm decode at all. Set it to `0` for the same-binary A/B back +to the staging behaviour. + +**It now decodes: 32/32 steps, at peak RSS 97.75 GiB of a 119.631 GiB box.** +Six limits, stated plainly rather than left to be discovered. * **The device has to be probed capable, and most are not.** The condition is `cudaDevAttrPageableMemoryAccess AND cudaDevAttrIntegrated` — an integrated, @@ -4597,23 +4610,34 @@ Four limits, stated plainly rather than left to be discovered. file, against the residency this process resolved, and a file that mixes a kept tower with a staged one keeps the whole bound as well (issue [#1378](https://github.com/mudler/vllm.cpp/issues/1378)). -* **The load now succeeds and the generation does not, so there is still no - speed claim.** The measurement ran on the one machine that answers true - (GB10, 2026-08-18) and it split: `--device cuda` loads this checkpoint in - 255-272 s, which it could not do before, and then exhausts the machine inside - its first forward without emitting a token - ([#1299](https://github.com/mudler/vllm.cpp/issues/1299)). The slot arena is - measurably not the cause — a 64-slot 0.15 GiB arena fails exactly where an - 8000-slot 18.55 GiB one does — so raising or lowering - `VT_MOE_EXPERT_STREAM_SLOTS` will not get you a token. **Use `--device cpu` - for this checkpoint today.** That arm serves it at a steady **11.05 s/token - at 4000 slots**, which is the count both recipes in this section set and the - only count that figure holds for. The same binary at 8000 slots measured a - 39.98-45.40 s/token median over two runs, and the second of them consumed all - 30,625 MiB of the box's swap, so **more slots is not a free knob here**: the - extra 9.27 GiB of arena takes the free memory the borrowed 370 GiB expert - mapping is served out of. Read `docs/BENCHMARKS.md` before assuming the GPU is - the faster arm here. +* **The correctness gate does NOT pass.** The 32 ids match the CPU arm for six + tokens and diverge at the seventh. Both continuations are coherent, and the + margins around it are measured and small: at that step the CPU arm's own + second-ranked token is exactly the one the CUDA arm emitted, behind by 1.4% of + the winning logit, and one step later the margin is 0.1%. **What CAUSES the + divergence is NOT identified.** The host-weight alias is EXCLUDED, measured ON + GB10 — same shapes, same algorithm, bit-identical output from a `cudaMalloc` + operand and from a 256-aligned host one — but excluding one cause is not + identifying another, and that the two arms simply run different GEMM kernels + over a near-tie is a standing hypothesis rather than a reading. Treat the CUDA + arm as unverified against the CPU arm until that gate is settled, and **use + `--device cpu` for this checkpoint today**: it is the arm every published + number here was measured on. +* **No speed claim is attached.** `docs/BENCHMARKS.md` carries G0-SPEED as + `VOID`, because a speed number behind a failing correctness gate is not a + result. The CPU arm serves this checkpoint at a steady **11.05 s/token at 4000 + slots**, which is the count both recipes in this section set and the only count + that figure holds for. Device access to host-resident weights on that part also has + a recorded penalty, and this lane reads ~6.95 GB of expert bytes per token that + way, so a CUDA arm slower than the CPU arm remains a real possible outcome. +* **More slots is not a free knob, and the reason is the page cache rather than + the arena.** The same binary at 8000 slots measured a 39.98-45.40 s/token + median over two runs, and the second consumed all 30,625 MiB of the box's swap: + the extra 9.27 GiB of arena takes the free memory the borrowed 370 GiB expert + mapping is served out of. The arena is also measurably not what exhausted the + box in [#1299](https://github.com/mudler/vllm.cpp/issues/1299) — a 64-slot + 0.15 GiB arena failed exactly where an 8000-slot 18.55 GiB one did — so this + knob was never the lever there either. ### The same thing as config, and which one wins diff --git a/include/vllm/model_executor/models/qwen3_5_weights.h b/include/vllm/model_executor/models/qwen3_5_weights.h index 6055fe830..b63edfbee 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,226 @@ 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. +// +// THE POINTER GATES, ENUMERATED. Two earlier revisions of this paragraph said +// there was exactly one, in `cuda_matmul_nvfp4.cu`, asking for 16. That was +// wrong both times. `grep -rn 'reinterpret_cast' src/vt/cuda/` plus +// `grep -rn PointerAligned src/vt/` finds at least seven, across four files: +// +// | site | asks | operand | +// |---|---|---| +// | `cuda_nvfp4_sm12x.cu:401` `PointerAligned(gate_up, 32)` | **32** | activation | +// | `cuda_nvfp4_sm12x.cu:401` `PointerAligned(packed, 8)` | 8 | weight | +// | `cuda_matmul_nvfp4.cu:204` `(prow) & 0xf` | 16 | weight row | +// | `cuda_matmul_nvfp4.cu:1757,1778` `(p) & 0xF` | 16 | scratch | +// | `cuda_matmul_nvfp4.cu:1841` `(out) & 0x7` | 8 | output | +// | `cuda_laguna.cu:67,70` `LagFastNormAligned{16,8}` | 16 | weight/act | +// | `cuda_ops.cu:370,395` `aligned16(w.data)` | 16 | NORM WEIGHT | +// +// The strictest is 32, and `cuda_ops.cu` is the one that binds a WEIGHT pointer +// on the ordinary decode path rather than a packed-arm buffer. The `% 32` and +// `% 64` tests near `cuda_matmul_nvfp4.cu` still read as alignment gates and +// still are not: they check a DIMENSION (`d`, `dv`), not an address. +// +// The conclusion is unchanged and is now correct at the widest of them: 256 +// dominates 32 as comfortably as it dominated 16. 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), while 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`). 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. +// +// CAN THE SUBSTITUTION MOVE A LOGIT? MEASURED, AND THE ANSWER IS NO. +// An earlier revision of this comment worried that "the heuristic may pick an +// algorithm on the strength of a promise a 16-aligned pointer breaks", and a +// fresh review was right that the worry was recorded here and never measured. +// It is measured now, on `thor:gpu0` (NVIDIA Thor `sm_110`, driver 13020, +// cuBLASLt 130101), which answers this file's own predicate TRUE +// (`cudaDevAttrPageableMemoryAccess = 1`, `cudaDevAttrIntegrated = 1`) and is +// therefore a member of the population this branch serves. Six shapes off the +// target checkpoint (`embedding_length = 8192`; M = 1, 5 and 32) crossed with +// BOTH formulations this tree issues — the row-major NN `MatmulKernelCuda`, +// where the weight is operand B, and the column-major TN `MatmulBTKernelCuda`, +// where it is operand A — for 12 measurements, `PROBE_FAILURES=0`: +// +// * The heuristic CANNOT see a pointer. `cublasLtMatmulAlgoGetHeuristic` takes +// (handle, desc, four layouts, preference) and no operands, so alignment +// reaches it only through the preference. 12/12 identical on a repeated +// call, 12/12 identical with the 256 promise stated EXPLICITLY, and 12/12 +// identical when the promise is weakened to 16 — the selection does not move +// on alignment at all. Five DIFFERENT algo configurations appear across the +// six shapes (tiles 393/537/573/576, workspaces 0 to 5,242,896), so the +// instrument does discriminate; it simply does not discriminate on this. +// * The OUTPUT is bit-exact. Running `cublasLtMatmul` with the same algo over +// the same bytes, the weight operand once from `cudaMalloc` and once from a +// 256-aligned host block, gives byte-identical results: 12/12 with zero +// differing elements, every status `SUCCESS`. +// +// So a 256-aligned host pointer cannot change a logit. +// +// AND THE SAME PROBE RAN ON THE GB10 ITSELF — the silicon the token gate ran on, +// so this is no longer a Thor result read across to another part. `rc` job +// `7c7a05e9-be87-48f4-94ae-1bbe0340f063` on `dgx:gpu0`, 2026-08-19 17:47 UTC, +// `NVIDIA GB10 sm_121`, driver 580.173.02, cuBLASLt 130101, the predicate +// re-derived in the job's own output (`pageableMemoryAccess=1 integrated=1`). +// Same six shapes, same two formulations, 12 measurements, `PROBE_EXIT=0`, +// `PROBE_FAILURES=0`: repeated heuristic call identical 12/12, the promise +// weakened to 16 moving nothing 12/12, and `cublasLtMatmul` output bit-exact +// 12/12 (`differing=0`). At least five distinct algo configurations appear +// across the twelve and they DIFFER from Thor's, so the heuristic was +// re-resolved rather than replayed. +// +// WHAT THAT ESTABLISHES, AND WHAT IT DOES NOT. It EXCLUDES this branch as the +// cause of the row's CUDA-versus-CPU token divergence, measured on the target +// silicon. It does not IDENTIFY the cause, and no reader of this block may take +// it as if it did: excluding one cause is not identifying another. An earlier +// revision of this comment named the two arms' GEMM arithmetic as the cause, +// which the row's own spec forbids asserting — that is the STANDING HYPOTHESIS, +// together with a greedy path whose top-2 margin at the divergent step is +// 0.264709 logits, and it is NOT MEASURED. Naming the first tensor whose values +// differ between the arms at that step, and the operation that produced it, is +// carried under `## Owed` in `.agents/specs/expert-stream-device-slots.md`. +// +// One more thing the probe does not license. The 16-aligned arm also came back +// bit-exact at these shapes, which is NOT a reason to lower this constant: the +// enumeration above still does not close, and a promise kept by luck at twelve +// shapes is not a promise. The literal below is pinned by a case in +// `tests/vllm/model_executor/test_resident_weight_host_addressable.cpp`, so +// lowering it reds a gate instead of passing every one of them silently. +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. +// +// PER CALL, NOT PER WEIGHT, AND THE DIFFERENCE IS THE WHOLE READING. There is no +// memo on the alias branch: `ResidentWeight` re-enters it for every weight on +// every forward step, so a weight aliased 32 times is counted 32 times. These are +// therefore BYTES SEEN — traffic — and they become a residency statement only +// when read at a stated point. The recorded 60.793 GiB is one such reading: the +// first-forward totals at call 1361, where re-homing plateaus and every dense +// weight has been seen exactly once. Quoted without that qualifier the same +// number is a traffic count, and after two decode steps the counter has passed +// 200 GiB on a checkpoint whose resident dense half is 60.8. +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 d8ec23dfc..9196a2f6c 100644 --- a/src/vllm/model_executor/models/qwen3_5.cpp +++ b/src/vllm/model_executor/models/qwen3_5.cpp @@ -1069,6 +1069,61 @@ 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; + // BOUNDED, because the counter this trips on is CUMULATIVE OVER CALLS and + // never stops growing. `ResidentWeight` re-enters the alias branch about 1,361 + // times per decode step, roughly 70 GiB of counted bytes, so a 4 GiB step + // prints about 17 lines EVERY step for the life of the process. The first + // forward is what this instrument exists to explain — it is where the aliasing + // set is established and where the recorded 60.793 GiB was read — and that + // fits inside the cap with room to spare. Everything after it is the same + // weights being counted again. + static int lines = 0; + constexpr int kMaxLines = 24; + if (lines >= kMaxLines) return; + ++lines; + const double gib = 1024.0 * 1024.0 * 1024.0; + // "per call", spelled out in the line itself. These are BYTES SEEN, not bytes + // resident: a weight aliased on every step is counted on every step, so the + // figures are traffic and become a residency measurement only when read at a + // stated point (see HostAliasStats in qwen3_5_weights.h). + std::fprintf(stderr, + "[vt load] w0f-alias per-call totals: 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); + if (lines == kMaxLines) + std::fprintf(stderr, + "[vt load] w0f-alias: %d lines printed; further lines are " + "suppressed (the counters keep running)\n", + kMaxLines); +} + 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 @@ -1135,6 +1190,93 @@ 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 with NEITHER host bytes NOR a device copy cannot be served at + // all, and the staging branch below would not notice: it would `Alloc(0)`, + // copy nothing, and hand out a pointer to nothing. That is the precondition + // this states. + // + // THE `w.d_dev` HALF IS NOT DEFENSIVE, AND THIS CHANGE IS WHAT CREATED THE + // POPULATION IT SERVES (found by a fresh review of #1299). A weight whose + // host mirror is gone but whose `d_dev` is populated has ALWAYS been served, + // by the memo below, and it returned the device tensor without complaint. + // W0f put this check ABOVE that memo, so the same weight began to throw. The + // justification written here first — "`ReleaseHost()` is not reachable for + // the dense weights this branch serves" — is true of the dense weights and + // FALSE of the expert weights the same function serves at the `gp/up/dp` + // capture below, whose misaligned GGUF borrows decline the alias, stage, get + // a `d_dev`, and are then released by the guarded loop beside that capture. + // So the condition is "nothing to serve", not "no host bytes". + VT_CHECK(!w.bytes.empty() || w.d_dev != nullptr, + "qwen3_5: a weight reaching device residency has no host bytes and " + "no device copy; its host mirror was released and there is nothing " + "to alias or upload"); + // ...and with no host bytes there is nothing to alias, so skip the attempt + // rather than charging a `kDeclinedEmpty` to the residency instrument for a + // weight that is already device-resident. + if (!w.bytes.empty()) { + 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); @@ -6853,11 +6995,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; @@ -7786,9 +7949,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_dense_weights.cpp b/src/vllm/model_executor/models/qwen3_5_dense_weights.cpp index 8c2c07709..f6dd8c3f4 100644 --- a/src/vllm/model_executor/models/qwen3_5_dense_weights.cpp +++ b/src/vllm/model_executor/models/qwen3_5_dense_weights.cpp @@ -969,7 +969,31 @@ size_t ReleaseResidentQwen3_5DenseHostWeights( Qwen3_5DenseWeights& weights) { size_t released = 0; const auto release = [&released](OwnedTensor& tensor) { - if (tensor.HasHostBytes() && (tensor.d_dev || tensor.d_dev_f32)) { + // `HostMirrorIsRedundant`, NOT `d_dev || d_dev_f32` (the second instance of + // the use-after-free W0f introduced; #1299). + // + // `d_dev_f32` is a bf16->f32 UPCAST into a separate device allocation. It is + // not a copy of these bytes and it can never stand in for them, so it never + // made a host mirror redundant — it only looked like it did while every + // weight that had one also had a `d_dev`. On the aliasing arm that stopped + // being true: `PrepareBf16Resident` passes exactly four weights to BOTH + // `raw()` and `f32()` — `gdn.conv1d_weight`, `gdn.norm_weight`, + // `attn.q_norm`, `attn.k_norm` — and there `raw()` ALIASES (leaving `d_dev` + // null) while `f32()` allocates (setting `d_dev_f32`). The disjunction then + // passes and frees the very bytes the aliased raw tensor points at. + // + // Nothing reaches that combination today, and the reason is worth writing + // down because it is an accident rather than a design: + // `DirectDeviceLoadEligible` above requires `!platform.is_unified_memory()`, + // and on CUDA `is_unified_memory()` and `host_memory_is_device_addressable()` + // are the SAME `pageable && integrated` conjunction, computed independently + // in `src/vt/cuda/cuda_backend.cu` and `src/vllm/platforms/cuda.cpp`. No + // rule states that equality, no document records it and no gate holds it, so + // a platform that ever separates the two arrives here with a live + // use-after-free. Asking the invariant instead of a proxy for it costs + // nothing: the disjunct is redundant on the staging arm, where every weight + // with a `d_dev_f32` has a `d_dev` too. + if (tensor.HasHostBytes() && HostMirrorIsRedundant(tensor)) { released += tensor.bytes.size(); tensor.ReleaseHost(); } diff --git a/src/vllm/model_executor/models/qwen3_5_internal.h b/src/vllm/model_executor/models/qwen3_5_internal.h index 8172dd5fc..88a4bdc18 100644 --- a/src/vllm/model_executor/models/qwen3_5_internal.h +++ b/src/vllm/model_executor/models/qwen3_5_internal.h @@ -564,7 +564,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..65a8c9ad0 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,190 @@ 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 { + +// ATOMIC, although the function that feeds it is single-threaded by +// precondition. The re-homing branch below genuinely requires one thread — it +// destroys the source vector — and saying so is honest. These counters are a +// different thing: a DIAGNOSTIC that `HostAliasSnapshot` publishes to any +// caller, so a plain `uint64_t` incremented on the model path and read +// elsewhere is a data race in the standard's terms whatever the precondition +// says about the branch. Relaxed ordering costs nothing measurable here (five +// increments against a memcpy of a weight) and makes the snapshot well defined. +struct AtomicHostAliasStats { + std::atomic aliased_in_place_bytes{0}; + std::atomic rehomed_bytes{0}; + std::atomic declined_borrow_bytes{0}; + std::atomic declined_other_bytes{0}; + std::atomic calls{0}; +}; + +AtomicHostAliasStats& AliasStats() { + static AtomicHostAliasStats s; + return s; +} + +} // namespace + +HostAliasStats HostAliasSnapshot() { + const AtomicHostAliasStats& s = AliasStats(); + HostAliasStats out; + out.aliased_in_place_bytes = s.aliased_in_place_bytes.load(std::memory_order_relaxed); + out.rehomed_bytes = s.rehomed_bytes.load(std::memory_order_relaxed); + out.declined_borrow_bytes = s.declined_borrow_bytes.load(std::memory_order_relaxed); + out.declined_other_bytes = s.declined_other_bytes.load(std::memory_order_relaxed); + out.calls = s.calls.load(std::memory_order_relaxed); + return out; +} + +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) { + AtomicHostAliasStats& st = AliasStats(); + st.calls.fetch_add(1, std::memory_order_relaxed); + 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.fetch_add(nbytes, std::memory_order_relaxed); + break; + case HostAliasOutcome::kRehomed: + st.rehomed_bytes.fetch_add(nbytes, std::memory_order_relaxed); + break; + case HostAliasOutcome::kDeclinedBorrow: + st.declined_borrow_bytes.fetch_add(nbytes, std::memory_order_relaxed); + break; + default: + st.declined_other_bytes.fetch_add(nbytes, std::memory_order_relaxed); + 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. + // + // ONCE, AND THE MEMO IS THE RECORD ITSELF (a fresh review of #1299 caught + // the repeat). `AdoptDeviceBytesAsHost` calls this from behind + // `if (!w.d_dev)`, so it ran exactly once per weight. THIS branch has no + // memo: `ResidentWeight` re-tests the alignment on every call, about 1,361 + // times per forward step on the target checkpoint. Left uncleared, + // `mmap_src` would make every one of those steps `madvise(MADV_DONTNEED)` + // the very pages the GPU is about to read, and the kernel would fault them + // straight back in. Correctness survives that; throughput does not. Clearing + // the record IS the memo, and it says the true thing: these source pages are + // spent. It also matches what `ReleaseHost` and the adoption branch below + // already do after they consume the same record. The borrow itself is + // untouched and stays a valid, re-faultable PROT_READ MAP_PRIVATE view. + ReleaseDirectUploadSource(w); + auto& consumed = *const_cast(&w); + consumed.mmap_src = nullptr; + consumed.mmap_src_bytes = 0; + 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 3b47e837d..308bbaf94 100644 --- a/tests/CMakeLists.txt +++ b/tests/CMakeLists.txt @@ -1258,6 +1258,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 d000b7c55..2825f5e43 100644 --- a/tests/vllm/model_executor/test_expert_stream_device_slot.cpp +++ b/tests/vllm/model_executor/test_expert_stream_device_slot.cpp @@ -344,15 +344,45 @@ 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`. const HostAddressable 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. + { + const HostAddressable discrete_platform(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()); + } } 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..643627354 --- /dev/null +++ b/tests/vllm/model_executor/test_resident_weight_host_addressable.cpp @@ -0,0 +1,613 @@ +// 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_dense.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 + +// THE CONSTANT ITSELF, PINNED TO ITS LITERAL. +// +// Every OTHER assertion about `kDeviceAliasAlignment` in this tree is written +// `% vllm::kDeviceAliasAlignment == 0` -- three times in the cases below, and +// twice more in `test_load_direct_upload.cpp`. Those check the BUFFERS, which is +// a real and separate property, and they stay. What none of them checks is the +// CONSTANT: they are tautologies in it. A fresh review measured that directly -- +// lower the constant from 256 to 16 and every one of them still holds, all +// twelve cases in this file pass, the whole suite reports SUCCESS and exits 0. +// The promise the alias branch rests on would be gone with no gate saying so. +// +// WHY 256, AND WHY IT MAY NOT BE QUIETLY LOWERED. The header's argument +// (`qwen3_5_weights.h`, above this constant) is that the enumeration of what a +// kernel may dereference does not close, while cuBLASLt is PROMISED 256: +// `CUBLASLT_MATMUL_PREF_MIN_ALIGNMENT_A_BYTES` defaults to 256 and this tree +// never overrides it -- `grep -rn MIN_ALIGNMENT src/vt/` returns nothing, and a +// fresh review re-confirmed that with a positive control. Matching the +// allocator's promise rather than the consumers is what makes handing a host +// pointer to a device kernel CORRECT. The tempting change is to lower the +// constant so the re-home memcpy is skipped; a 16-aligned arm even came back +// bit-exact at twelve probe shapes. A promise kept by luck at twelve shapes is +// not a promise, and this case is what makes that argument executable. +TEST_CASE("the alias alignment constant IS 256, the promise cuBLASLt is given") { + CHECK(vllm::kDeviceAliasAlignment == 256u); +} + +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 with NOTHING to serve 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()); + REQUIRE(w.d_dev == nullptr); + Queue q = XpuQueue(); + + CHECK_THROWS_WITH_AS( + vllm::detail::StageWeightForTest(q, w), + doctest::Contains("no host bytes and no device copy"), std::runtime_error); +} + +TEST_CASE("a released host mirror with a DEVICE copy is still served, not refused") { + // THE OTHER HALF OF THAT PRECONDITION, and W0f is what created the population + // it protects (a fresh review of #1299 found the regression). A weight whose + // host mirror is gone but whose `d_dev` is populated has always been served — + // by the `if (!w.d_dev)` memo, which returns the device tensor. W0f put the + // "no host bytes" refusal ABOVE that memo, so the same weight began to throw, + // and the justification written beside it ("`ReleaseHost()` is not reachable + // for the dense weights this branch serves") is true of the dense weights and + // false of the expert weights the same function serves: a misaligned GGUF + // expert borrow declines the alias, stages, gets a `d_dev`, and is then + // released by the guarded loop beside the `MoeBlockBf16Cuda` pointer capture. + // + // Built on the arm that STAGES, then released, then asked again on the ALIASING + // arm — which is the ordering that reproduces it, because the refusal only + // lives on the aliasing side. + OwnedTensor w = MakeWeight(/*tag=*/13); + Queue q = XpuQueue(); + { + const PlatformArm discrete(false); + const Tensor staged = vllm::detail::StageWeightForTest(q, w); + REQUIRE(w.d_dev != nullptr); + REQUIRE(staged.data == w.d_dev.get()); + } + w.ReleaseHost(); + REQUIRE(w.bytes.empty()); + REQUIRE(w.d_dev != nullptr); + + const PlatformArm arm(true); + const int allocs_before = Fake().allocs; + const Tensor t = vllm::detail::StageWeightForTest(q, w); + + // Served from the copy that exists, with nothing allocated and nothing thrown. + CHECK(t.data == w.d_dev.get()); + CHECK(Fake().allocs == allocs_before); + CHECK(t.dtype == w.dtype); + CHECK(t.shape[0] == kN); + CHECK(t.shape[1] == kK); +} + +TEST_CASE("an F32 UPCAST does not make an aliased weight's host mirror redundant") { + // THE SECOND INSTANCE OF THE USE-AFTER-FREE (a fresh review of #1299 found it + // after the first was fixed). `ReleaseResidentQwen3_5DenseHostWeights` guarded + // on `d_dev || d_dev_f32`, and `d_dev_f32` is a bf16->f32 UPCAST into a + // separate device allocation — not a copy of these bytes, and never able to + // stand in for them. The disjunction was harmless only while every weight with + // an upcast also had a `d_dev`. + // + // `PrepareBf16Resident` passes exactly four weights to BOTH `raw()` and + // `f32()`: `gdn.conv1d_weight`, `gdn.norm_weight`, `attn.q_norm` and + // `attn.k_norm`. On the aliasing arm `raw()` leaves `d_dev` null while `f32()` + // sets `d_dev_f32`, so the guard passed and freed the bytes the aliased raw + // tensor points at. + // + // `d_dev_f32` is set directly here, which is the state that pairing produces: + // `ResidentWeightF32` is private to qwen3_5.cpp and has no gate seam, and + // inventing one would add test-only surface to observe a field the production + // release site reads directly. + const PlatformArm arm(true); + OwnedTensor w = MakeWeight(/*tag=*/14); + Queue q = XpuQueue(); + const Tensor t = vllm::detail::StageWeightForTest(q, w); + REQUIRE(w.d_dev == nullptr); + REQUIRE(t.data == static_cast(w.bytes.data())); + w.d_dev_f32 = std::shared_ptr(reinterpret_cast(1), [](void*) {}); + + // The DECISION, which is what the release site now asks and what a mutation of + // the invariant moves. + CHECK_FALSE(vllm::HostMirrorIsRedundant(w)); + + // ...and the PRODUCTION release, driven end to end over the four-weight + // pairing that creates the state. Nothing may be freed, and the aliased + // pointer the kernel holds must still be the weight's own bytes afterwards. + vllm::Qwen3_5DenseWeights weights; + weights.layers.resize(1); + vllm::Qwen3_5DenseLayerWeights& layer = weights.layers[0]; + layer.is_linear_attention = true; + layer.gdn.conv1d_weight = std::move(w); + const void* aliased_at = static_cast(layer.gdn.conv1d_weight.bytes.data()); + REQUIRE(t.data == aliased_at); + + CHECK(vllm::ReleaseResidentQwen3_5DenseHostWeights(weights) == 0); + CHECK(layer.gdn.conv1d_weight.HasHostBytes()); + CHECK(static_cast(layer.gdn.conv1d_weight.bytes.data()) == aliased_at); + + // THE DISCRETE HALF, without which the invariant could be satisfied by + // refusing every release — silently undoing a measured host-memory lever. A + // STAGED weight has an authoritative device copy and is still freed. + vllm::Qwen3_5DenseWeights staged_weights; + staged_weights.layers.resize(1); + vllm::Qwen3_5DenseLayerWeights& staged_layer = staged_weights.layers[0]; + staged_layer.is_linear_attention = true; + { + const PlatformArm discrete(false); + staged_layer.gdn.conv1d_weight = MakeWeight(/*tag=*/15); + (void)vllm::detail::StageWeightForTest(q, staged_layer.gdn.conv1d_weight); + } + REQUIRE(staged_layer.gdn.conv1d_weight.d_dev != nullptr); + staged_layer.gdn.conv1d_weight.d_dev_f32 = + std::shared_ptr(reinterpret_cast(1), [](void*) {}); + const size_t expect_freed = staged_layer.gdn.conv1d_weight.bytes.size(); + CHECK(vllm::ReleaseResidentQwen3_5DenseHostWeights(staged_weights) == expect_freed); + CHECK_FALSE(staged_layer.gdn.conv1d_weight.HasHostBytes()); +} diff --git a/tests/vllm/test_load_direct_upload.cpp b/tests/vllm/test_load_direct_upload.cpp index 4bdd79a56..71046d1c6 100644 --- a/tests/vllm/test_load_direct_upload.cpp +++ b/tests/vllm/test_load_direct_upload.cpp @@ -522,6 +522,80 @@ 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_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); + // The record is SPENT, and clearing it is the memo. See the once-only case + // below for why this branch needs one at all. + CHECK(w.mmap_src == nullptr); + CHECK(w.mmap_src_bytes == 0); +} + +TEST_CASE("alias: the source pages are released ONCE, not once per forward step") { + // THE REPEAT THE PAGE-RELEASE FIX INTRODUCED (a fresh review of #1299 caught + // it). `AdoptDeviceBytesAsHost` calls `ReleaseDirectUploadSource` from behind + // `if (!w.d_dev)`, so it ran exactly once per weight for the life of the + // process. The alias branch has NO such memo: `ResidentWeight` re-tests the + // alignment on every call, about 1,361 times per forward step on the target + // checkpoint. Left uncleared, `mmap_src` would make every one of those calls + // `madvise(MADV_DONTNEED)` the pages the GPU is about to read, and the kernel + // would fault them straight back in — correct output, and a throughput + // regression invisible to every token gate. + // + // THE FIRST VERSION OF THIS CASE ASSERTED ONLY THAT THE RELEASE HAPPENED. That + // is satisfied by a release that happens every time, which is the defect. This + // asserts the COUNT, by restoring the pattern and looking for it again. + ForcedResidencyArm arm; + ObservableMapping m; + vllm::OwnedTensor w = BorrowedWeight(m, 2 * PageSize()); + REQUIRE(reinterpret_cast(w.bytes.data()) % + vllm::kDeviceAliasAlignment == 0); + + CHECK(vllm::MakeHostBytesDeviceAliasable(w)); + REQUIRE(m.addr[0] == 0); // the first release ran; the case is not vacuous + + // Re-arm the observation, then take the SAME branch again the way a second + // decode step would. + std::memset(m.addr, kSrcPattern, m.size); + CHECK(vllm::MakeHostBytesDeviceAliasable(w)); + CHECK(vllm::MakeHostBytesDeviceAliasable(w)); + + // The pattern SURVIVES: no second madvise. The alias itself still succeeded + // above, so this is a count and not a disabled branch. + CHECK(m.addr[0] == kSrcPattern); + CHECK(w.bytes.borrowed()); + CHECK(static_cast(w.bytes.data()) == static_cast(m.addr)); +} + 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");