diff --git a/.agents/benchmark-record.md b/.agents/benchmark-record.md
index 70feae65c..cb1a345dd 100644
--- a/.agents/benchmark-record.md
+++ b/.agents/benchmark-record.md
@@ -24800,3 +24800,114 @@ CUDA/ROCm/Vulkan is not covered. The counter sees every global `operator new`
 and nothing below it, so a raw `std::malloc` in the kernel is outside it (M-E
 above). No full-model render has been rerun, so nothing here says what such a
 run would now do — only that #1252 is not what stopped the last one.
+
+## ENG-EXPERT-STREAM-DEVICE W0e — the CPU arm is measured and the CUDA arm loads without generating (2026-08-18/19, `dgx:gpu0`, source `95883dcae`, #1124)
+
+One `rc hold` (`edb4b3d0-5d6e-422f-ade6-bff5339e3396`, 22:13:12Z to 00:58Z,
+released by interrupting its client). GB10, sm_121a, driver 580.173.02, CUDA
+13.0.88, 20 cores, 122,502 MiB RAM, 30,625 MiB swap. Checkpoint
+`Qwen3.8-2.4T-A95B-UD-Q1_0` (369.97 GiB, 10 shards) from LOCAL NVMe at
+`/home/mudler/ckpt/qwen3.8-q1_0`. Prompt fixed as token ids
+`760,6511,314,9338,369`, greedy, 32 tokens, `--max-num-seqs 1`,
+`VT_GGUF_PREFAULT=0 VT_MOE_EXPERT_STREAM=1 VT_MOE_EXPERT_STREAM_STATS_EVERY=1`.
+Page cache dropped before every arm. Same binary both arms, built with CUTLASS
+4.5.0 and Triton AOT: the configure log reports `fp4-mma`, `cutlass-nvfp4`,
+`cutlass-fp8`, `marlin-nvfp4` and `fa2` all `ENABLED for [121a]`, plus
+`FlashAttention-2 prefill/decode: ENABLED` and `Triton AOT W2: embedded trees
+[sm_80;sm_86;sm_89;sm_90a;sm_100a;sm_121a]`.
+
+Harness: `benchmarks/expert_stream_device_w0e.cpp`, a thin C-ABI client. It
+exists because one generation has to yield three things no shipped command gives
+together: the generated token ids (`vllm_complete_tokens`), a per-step arrival
+timestamp (a logits processor used as a pure observer, which edits nothing so
+the argmax is unchanged), and the expert-stream counters at both ends of the run.
+
+### CPU arm, four runs, two slot counts
+
+| slots | arena | load | TTFT | steady s/token, steps 4-32 | generate | peak RSS | min avail | peak swap |
+|---|---|---|---|---|---|---|---|---|
+| 4000 | 9.28 GiB | 271.1 s | 85.90 s | median 11.22, min 9.62, max 12.51 | 502.3 s | 86.5 GiB | 16,840 MiB | not sampled |
+| 4000 | 9.28 GiB | 255.7 s | 79.09 s | median **11.05**, min 9.43, max 13.25 | 460.7 s | 86.5 GiB | 16,347 MiB | 6,883 MiB |
+| 8000 | 18.55 GiB | 261.6 s | 94.25 s | median 45.40, min 20.73, max 82.68 | 1643.2 s | 86.6 GiB | 6,985 MiB | not sampled |
+| 8000 | 18.55 GiB | 266.5 s | 132.74 s | median 39.98, min 16.06, max 94.30 | 1581.6 s | 86.6 GiB | 6,941 MiB | 30,625 MiB (all) |
+
+The 8000-slot rows are a MEMORY result, not a cache result. A bigger cache came
+out 3.6x slower and 4x more variable, which is the wrong direction, and the last
+column says why: an 18.55 GiB arena does not fit beside 62 GiB of dense weights
+on a 119.63 GiB box, so those runs swapped and one exhausted swap entirely. Both
+reps of each slot count reproduce each other, so this is the box and not a
+fluke. The publishable figure is the 4000-slot one.
+
+Per-token deltas, 4000-slot rep 2, seconds: 79.09, 51.14, 12.52, then 11.69,
+11.51, 11.35, 11.59, 13.25, 11.58, 9.48, 12.16, 11.22, 11.42, 11.51, 11.05,
+10.70, 10.69, 11.93, 11.12, 11.26, 10.99, 9.85, 10.39, 9.58, 11.35, 10.17,
+11.03, 10.31, 9.43, 10.25, 10.38, 10.64. Step 1 is prefill; steps 2 and 3 are
+still filling a cold cache.
+
+G0-LIVE, gated on the DIFFERENCE and never the total. At 4000 slots the
+after-prefill snapshot reads `steps=1 hits=0 misses=10074 evictions=0 fills=4000
+bytes=9961472000 exhausted=6074` and the final one `steps=32 hits=37096
+misses=58538 evictions=48464 fills=52464 bytes=130654666752 exhausted=6074`, so
+the decode-phase delta is 0 over 31 decode steps. At 8000 slots the same pair
+reads 2074 and 2074, delta 0. Counters are byte-identical between the reps of
+each slot count.
+
+All four runs produced the same 32 ids:
+`11751,13,11751,369,264,3177,7172,303,279,17631,919,314,9338,11,383,279,181474,10629,13,1049,369,279,7526,3177,303,9338,321,369,3750,364,1141,25438`,
+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`.
+
+### CUDA arm, seven attempts, zero decode steps
+
+| slots | arena | prompt tokens | load | RSS after load | steps | peak system used | peak swap |
+|---|---|---|---|---|---|---|---|
+| 8000 | 18.55 GiB | 5 | 267.2 s | 61.20 GiB | 0 | 100,215 MiB | not sampled |
+| 4000 | 9.28 GiB | 5 | 271.6 s | 61.20 GiB | 0 | 94,737 MiB | not sampled |
+| 4000 | 9.28 GiB | 5 | 255.3 s | 61.20 GiB | 0 | 120,351 MiB | 30,569 MiB |
+| 3500 | 8.12 GiB | 1 | 262.8 s | 61.20 GiB | 0 | 120,296 MiB | 28,138 MiB |
+| 3500 | 8.12 GiB | 1 | 265.3 s | 61.20 GiB | 0 | 120,347 MiB | 30,461 MiB |
+| 3500 | 8.12 GiB | 1 | 272.3 s | 61.20 GiB | 0 | 120,306 MiB | 30,172 MiB |
+| 64 | 0.15 GiB | 1 | 268.9 s | 61.20 GiB | 0 | 118,257 MiB | 30,211 MiB |
+
+The LOAD is new and it works: this checkpoint used to refuse on `--device cuda`
+(#1123) and now loads, and the `[expert-stream] ON` banner prints on the device
+arm, which is the first production evidence that W0b's CUDA leg is reached and
+answers true on real hardware.
+
+The last row is the diagnosis. A 0.15 GiB arena dies where an 18.55 GiB one
+does, 124x smaller and no further along, so the slot arena is not the cost. The
+in-place tower fallback is not the cost either: a 1-token prompt has a protected
+set of 93 x 3 x 10 = 2,790 slices, fits 3500 slots with no fallback taken, and
+behaves exactly like the 5-token prompt whose 13,950-slice set fits no arena
+here.
+
+Growth is anonymous, not file-backed, so the GPU is not pinning the mapping's
+pages through its address translation. That was the first hypothesis and it is
+refuted. Sampling the container process every 5 s: `RssAnon` 8.1, 13.9, 25.9,
+38.3, 49.4, 61.4 GB through the load while `RssFile` stays at 0.1-0.2 GB, then
+`VmSwap` 0, 5.4, 11.4, 16.5, 27.1, 31.0 GB inside the forward. Host `RssAnon`
+plus `VmSwap` reaches about 65 GB while the system reports ~119 GiB used, and
+the ~42 GiB difference is device memory this unified part does not charge to the
+process RSS. The non-expert weights are therefore resident twice, once as the
+host-side `OwnedTensor` and once as the `ResidentWeight` device staging copy.
+About 50 GiB of the 61.20 is the bf16 expansion the GDN V-head reorder forces on
+`attn_qkv` and `ssm_out`, already measured in
+[expert-streaming.md](specs/expert-streaming.md). Filed as
+[#1299](https://github.com/mudler/vllm.cpp/issues/1299).
+
+Each CUDA run was stopped by a guard that kills the container when MemAvailable
+plus SwapFree falls under a floor, because an out-of-memory kill on GB10 takes
+the machine down rather than the process. The guard is not what makes them fail:
+the last rows were climbing at roughly 10 GB of swap per minute with under
+600 MiB of swap left. The guard's own first version counted MemAvailable ALONE
+and killed two runs while 29 GiB of swap sat untouched, which is worth recording
+because it reported a box about to die when the kernel had a whole tier left.
+
+Contention: the lease excluded every other `rc` job. Two things it did not
+exclude are recorded rather than assumed away. An orphaned `VLLM::EngineCore`
+from an earlier session held 3.32 GiB of host RSS throughout; it is not this
+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.
diff --git a/.agents/issue-index.md b/.agents/issue-index.md
index 01458865e..3918cf973 100644
--- a/.agents/issue-index.md
+++ b/.agents/issue-index.md
@@ -397,6 +397,7 @@ rather than merged. `scripts/check-agent-record.py` gates both.
 | [#1226](https://github.com/mudler/vllm.cpp/issues/1226) | `ENG-CUDAGRAPH-DEDUP` | The owed device A/B for #1162 RAN on `dgx:gpu0` 2026-08-18 (GB10, driver 580.173.02, nvcc 13.0.88, `rc` job f88d484b, gated commit `72de552c8`, whose four dedup sources are byte-identical to the merged `2a976eb9f`) and it SPLIT. CORRECTNESS PASSED and closes [#1184](https://github.com/mudler/vllm.cpp/issues/1184): 12/12 cells exit 0, zero `invalid device function` and zero `engine-fatal` where the pre-fix head `e4ce5571a` died after exactly one replay, ON replays as often as OFF (60=60, 33=33, 43=43), and `--output-token-ids` is IDENTICAL over 10/10 comparisons with the three OFF/OFF controls passing FIRST and three workloads hashing to three DIFFERENT values. THE BENEFIT IS REFUTED for the case the row was filed for: `N == M` in every `VT_CUDA_GRAPH_DEDUP=1` cell — 3 graphs to 3 execs on sizes [24 16 8], 2 to 2 on [16 8], 2 to 2 on [32 24] — with the registry count CLIMBING 1→1, 2→2, 3→3, so more than one capture reached it and the 1:1 is a measurement rather than the single-capture artefact of the first attempt. The cause is structural and was pre-registered before the run: `AppendKernelPayload` hashes (`func`, `gridDim.{x,y,z}`, `blockDim.{x,y,z}`, `sharedMemBytes`) at `src/vt/graph_dedup_runtime.h:121-128` and the memcpy payload hashes the copy extent, so the padded batch dimension is IN THE KEY, no candidate group forms and `cudaGraphExecUpdate` is never attempted. That contradicts the row's own premise, and SGLang keys the same fields (`cuda_graph_dedup_mixin.py:105-114`), so whatever folds upstream is not decode buckets either. NO throughput or memory number is recorded — clocks unpinned AND the ON arm allocated exactly as many executables as OFF. `VT_CUDA_GRAPH_DEDUP` therefore stays OFF and the default flip is NOT JUSTIFIED on this evidence. THE OPEN HYPOTHESIS this issue owns, deliberately not decided by the record that filed it: a COARSER key keeping the function addresses and topology but dropping the launch dimensions and memcpy extents would let two padded buckets form a candidate group at all, and the probe-before-fold design means a wrong grouping costs one wasted `cudaGraphExecUpdate` probe and a private executable rather than a wrong replay, so it is a cost question needing its own spec, a red-first discrimination test and a device probe-refusal rate. Also owed here: whether ANY other capture site or two models sharing the process-singleton registry can fold (only the Qwen3 dense decode driver was exercised), and a supporting `orin:gpu0` leg, BLOCKED because the Jetson 540.4.0 driver cannot run a CUDA 13 runtime (`cudaGetDeviceCount err=35`; a CUDA 12.x toolkit is the untried route). Honest gaps carried into the record: per-shape replay counts are unavailable, so B's ~30-per-shape is arithmetic over a printed TOTAL; the driver's "N captured size(s)" counts SLOTS not captures; the container's own cuBLASLt was never re-tested at CUDA 13.0. Full evidence [`benchmark-record.md`](benchmark-record.md) entry `ENG-CUDAGRAPH-DEDUP W4`; spec [`eng-cudagraph-dedup.md`](specs/eng-cudagraph-dedup.md) `## Outcome` | perf |
 | [#1229](https://github.com/mudler/vllm.cpp/issues/1229) | `ENG-CUDAGRAPH-DEDUP` | `origin/main` went RED at `fe24a3029` because `tests/scripts/test_check_gate_commands.py:670` pinned the literal string `Device byte-identity A/B (owed` inside `.agents/specs/eng-cudagraph-dedup.md`, and that commit legitimately rewrote the sentence from recording the device A/B as owed to recording that it RAN and what it found. The record edit is correct and the assertion described a state the record left, so every branch that merged `origin/main` inherited the failure. The protected property is stated by the assertion's own comment: the row is credited in `RUNNABLE_BASELINE` for `ctest -R test_graph_dedup` and `./scripts/agent-preflight.sh`, both CPU-tier, while gate item 6 of the spec contributes NO runnable command, so the credit covers the CPU tier and the record must not go silent about the leg it does not cover. That property is STATE-INDEPENDENT: `owed` was the leg's value on the day, not the property, and silence is the defect whether the leg was never run or the line was deleted after it ran. Asserting the NEW wording would only move the pin from `(owed` to `RAN 2026-08-18, PASS`, and the rerun this row already owes under [#1226](https://github.com/mudler/vllm.cpp/issues/1226) against a coarser signature key would red it again. FIXED by keying on structure: `gate_items`, `item_lead` and `gate_disposition` land in `scripts/check-gate-commands.py`, and the test now asserts that exactly one gate item names the device byte-identity A/B as its SUBJECT, that `runnable_commands` extracts nothing from it, and that its bold lead declares a disposition from a closed vocabulary. Lead-scoping is measured rather than assumed: a whole-item search credits item 5 with a disposition it does not declare, because its body contains `fail` in the phrase "proving the instrument can fail". Both historical wordings of item 6 are pinned as fixtures, which is what proves the new assertion is state-independent rather than re-pinned to today's sentence. The rule is NOT swept over every gate item and the reason is measured, not assumed: items 1 and 5 yield no command and declare no disposition, so a sweep is red on arrival, and widening the vocabulary until they pass would make it match ordinary English. That general rule is carried under `## Owed`. Also established and NOT a defect: `ERROR: these baseline rows left the gated population: ROW-THAT-IS-NOT-THERE` on the same run is the expected negative-path output of `test_check_mode_is_never_silently_swallowed_by_json`, confirmed by running that fixture alone. Spec [`fix-gate-commands-prose-pin.md`](specs/fix-gate-commands-prose-pin.md) | bug |
 | [#1240](https://github.com/mudler/vllm.cpp/issues/1240) | `MODEL-FP8-BLOCK-WEIGHT` | `test_qwen27_dense_forward` and `test_qwen27_dense_forward_glue_fuse_off` are RED on `origin/main` at `aba8d5ffb`. Both throw the block-wise-FP8 refusal `09597106e` (#1228, issue #1189) added at `src/vllm/model_executor/models/qwen3_5_dense_weights.cpp:428`: `'layer.linear_attn.in_proj_qkv.weight_scale_inv' is present ... but the checkpoint's quantization_config declares no weight_block_size`. ATTRIBUTED by measurement, not inference: `git log -1` on that file is `09597106e`, and `git show --stat 09597106e` changes it (+219) while touching neither test — the producer moved and its consumer did not. It hides behind the doctest thrown-case shape: the run reports 563 assertions with NONE failed beside 9 cases with ONE failed, because a case that THROWS records no failed assertion, so the assertions line reads clean and only `Status: FAILURE!` and the exit code tell the truth. NOT fixed in flow, deliberately: the refusal is correct in its own terms (guessing 128x128 is exactly what #1166 asked not to happen), so whether the synthetic fixture is wrong to emit `weight_scale_inv` without `weight_block_size`, or the refusal is too broad for a checkpoint that is not block-wise FP8 at all, is a semantics choice belonging to the owning row with its own red-first evidence — not a fixture edit that makes the message go away. Found by `ENG-EXPERT-STREAM-DEVICE` W0 running the full ctest suite after merging `origin/main` (523 of 525 passed; these two the only failures, and W0 touches neither file) | bug |
+| [#1299](https://github.com/mudler/vllm.cpp/issues/1299) | `ENG-EXPERT-STREAM-DEVICE` | `Qwen3.8-2.4T-A95B UD-Q1_0` (369.97 GiB) now LOADS on `--device cuda` on a 119.631 GiB GB10 — `ENG-EXPERT-STREAM-DEVICE` W0 works, the slot store builds and the #1123 refusal correctly does not fire — and then exhausts the machine inside the FIRST forward, producing zero decode steps. Measured on `dgx:gpu0` inside one `rc hold` at source `95883dcae`, production stack (CUTLASS 4.5.0 found, `cutlass-fp8`/`cutlass-nvfp4`/`fa2` ENABLED for `[121a]`, Triton AOT `sm_121a`). Seven attempts, all identical: load 255-272 s, RSS after load 61.20 GiB, 0 decode steps, peak system `used` 118-120 GiB with 28-31 GiB of swap consumed. The slot arena is NOT the cost, and that is measured rather than argued: a **64-slot 0.15 GiB** arena dies exactly where an 8000-slot 18.55 GiB one does, and a 1-token prompt (protected set 2,790 slices, which fits 3500 slots with no in-place tower fallback at all) behaves exactly like a 5-token one. The growth is ANONYMOUS, not file-backed, so the GPU is not pinning the mapping's pages: `RssAnon` climbs 8.1 -> 61.4 GB through the load while `RssFile` stays at 0.1-0.2 GB, then `VmSwap` climbs 0 -> 31.0 GB inside the forward. Host `RssAnon` + `VmSwap` reaches ~65 GB while system `used` sits at ~119 GB, and the ~42 GiB difference is device memory the process RSS does not account for on this unified part — so the non-expert weights are resident TWICE, once as the host-side `OwnedTensor` and once as the `ResidentWeight` device staging copy. `.agents/specs/expert-streaming.md` measured the host half already: the GDN V-head reorder makes `attn_qkv` and `ssm_out` `kTransformedWeight`, which expands them from ~5.5 bits to bf16, about 50 GiB of the 61.20. The CPU arm pays it once and serves (32 tokens, steady 11.05 s/token at 4000 slots, decode-phase `exhausted` delta 0); the CUDA arm pays it twice and cannot. NOT #1123, whose refusal fired at load and no longer applies, and NOT a defect in W0b/W0c/W0d, which do what their spec says. This is the device-side half of the sentence `expert-streaming.md` already carries: "Whoever takes this needs BOTH: the streaming lane for the ~330 GiB of experts, and a transformed-weight path that does not expand." Owed under `## Owed` in [`expert-stream-device-slots.md`](specs/expert-stream-device-slots.md) | bug |
 | [#1249](https://github.com/mudler/vllm.cpp/issues/1249) | `MODEL-FP8-BLOCK-LINEAR` | An f32 paged KV cache refuses every arm whose `v_proj` emits bf16. `vt::ReshapeAndCache`'s auto path requires `k`, `v`, `k_cache` and `v_cache` to share one float dtype (`src/vt/ops.cpp:3146`), and the CPU kernel behind it is a byte `memcpy` per token (`src/vt/cpu/cpu_cache.cpp:33-72`), so there is no conversion to widen. That mirrors upstream `reshape_and_cache_flash`, which torch types the same way, and it is correct for the production configuration, where the cache dtype follows the model dtype. It is wrong for `VT_KV_CACHE_F32=1`, which selects an f32 cache while every bf16 model's `v_proj` keeps emitting bf16 — so that lever is unusable on any bf16 arm, not merely on the block-wise FP8 one that found it. Found while wiring #1189 M4, whose G3 reachability case drives the bf16 cache and records in the source that it cannot drive the f32 one. Closing it is a `vt` semantic change (a converting store, or a refusal that names the lever) and needs its own spec and red-first test, so it is NOT fixed in M4's flow; owned by `MODEL-FP8-BLOCK-LINEAR` under its spec's `## Owed` | bug |
 | [#785](https://github.com/mudler/vllm.cpp/issues/785) | `BACKEND-ROCM` | ROCm: `PagedAttnPrefillSharedKWmma` host launches were behind `#if defined(VT_ROCWMMA_OK)`, which is never true on the host pass. This row repairs d=256 host dispatch only (`<2,8,16,32,false>`). d=512 stays scalar SharedK because shipping-f58b WMMA spills (52/212) and is a separate kernel repair. Spec [`rocm-sharedk-wmma-host-dispatch.md`](specs/rocm-sharedk-wmma-host-dispatch.md) | bug |
 | [#1217](https://github.com/mudler/vllm.cpp/issues/1217) | `MODEL-NEMOTRON-H-ABI-A2P` | `ModelForwardInput::device_token_ids` carries the async runner's device-combined ids and its contract is that `token_ids` is STALE for decode rows whenever the pointer is non-null (`model_registry.h:314-324`). A registered forward that embeds the host vector then embeds the same placeholder id on EVERY decode step. The field's own comment says a model that ignores it "is simply never given one", but `runner.cpp:1408` sets the pointer for whatever model the step routes to, with no per-model opt-in and no check \| two models have now been cut from the identical divergence: Kimi-Linear (`kimi_linear_device.cpp:2270-2280`, the GB10 9/128 case) and NemotronH's paged forward under [#1157](https://github.com/mudler/vllm.cpp/issues/1157), whose A3 gate read 4/24 on GB10 against 96/96 for the same binary on CPU where the pointer is always null \| invisible because the runner sets it only under `VLLM_CPP_CUDA` with a live device mirror, so no CPU gate reaches the branch, and the failure is fluent wrong tokens rather than an error \| two closes: give `ModelFactory` an explicit `honors_device_token_ids` and have the runner fall back to the synchronous host path for a forward that has not declared it, or add a checker over the registered `.forward` entry points (a file-level grep flags ~25 false positives because several models delegate through `detail::DeviceTokenIdsScope` or the shared dense block) \| NOT fixed in the #1157 flow because one close changes a shared seam and every model factory and the other changes checker semantics, which is the "needs its own spec" case rather than the in-flow case. Listed under `## Owed` in [`nemotron-h-a2p-paged-forward.md`](specs/nemotron-h-a2p-paged-forward.md) | bug |
diff --git a/.agents/specs/expert-stream-device-slots.md b/.agents/specs/expert-stream-device-slots.md
index 5faac4c32..819a8f1e3 100644
--- a/.agents/specs/expert-stream-device-slots.md
+++ b/.agents/specs/expert-stream-device-slots.md
@@ -10,11 +10,23 @@ platform may read it.
 
 ## Now
 
-`ACTIVE`. **W0a HAS RUN on `dgx:gpu0` and answered `PAGEABLE_OK`. W0b, W0c and
-W0d are implemented and unit-gated on the CPU tier. W0e is still QUEUED behind
-the lease, so no decode number exists yet and none is claimed.**
+`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`.
 
-What that means precisely, because "W0 landed" would overstate it:
+The three gates, reported one result each:
+
+| 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** median (min 9.43, max 13.25, 29 samples, 2 reps) at 4000 slots. No CUDA number exists, so no ratio is reported and none may be inferred. |
+
+What that means precisely, because both "W0 landed" and "W0 failed" would
+misstate it:
 
 * **W0b — the predicate.** `Platform::host_memory_is_device_addressable()`,
   base false, CUDA from `cudaDevAttrPageableMemoryAccess AND
@@ -68,14 +80,18 @@ What that means precisely, because "W0 landed" would overstate 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.** Not run. G0-CORRECT, G0-LIVE and G0-SPEED are all
-  `PENDING` on the lease.
-
-So `--device cuda` on `Qwen3.8-2.4T-A95B UD-Q1_0` is no longer refused by the
-predicate on this box: the probe has measured `PageableMemoryAccess == 1`, and
-the predicate was probed rather than assumed precisely so that this sentence
-could change on evidence. What has NOT run is the load itself. The developer's
-target remains that GPU figure, and W0e is what produces it.
+* **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.
 
 ## Scope
 
@@ -328,6 +344,193 @@ gate that stays green without it measured a class, not a capability.
 
 **G-DISCRETE (owed, cannot run here).** See `## Owed`.
 
+## Evidence
+
+W0a and W0e, `dgx:gpu0` (GB10, sm_121a, driver 580.173.02, CUDA 13.0.88, 20
+cores, 122,502 MiB RAM, 30,625 MiB swap), one `rc hold`
+`edb4b3d0-5d6e-422f-ade6-bff5339e3396`, 2026-08-18T22:13:12Z to
+2026-08-19T00:58Z, released by interrupting its client. Source `95883dcae`,
+the head of PR [#1242](https://github.com/mudler/vllm.cpp/pull/1242).
+
+### The build, because a degraded one would have voided every number
+
+`cmake -DCMAKE_BUILD_TYPE=Release -DVLLM_CPP_CUDA=ON
+-DVLLM_CPP_CUDA_ARCHITECTURES=121a -DVLLM_CPP_CUTLASS_DIR=/cutlass
+-DVLLM_CPP_TRITON=ON`, CUTLASS 4.5.0 staged host-side. The configure log
+reports `fp4-mma`, `cutlass-nvfp4`, `cutlass-fp8`, `marlin-nvfp4` and `fa2` all
+`ENABLED for [121a]`, `CUTLASS found at /cutlass; enabling sm120a NVFP4 cutlass
+GEMM`, `FlashAttention-2 prefill/decode: ENABLED for arch(es) [121a]`, and
+`Triton AOT W2: embedded trees [sm_80;sm_86;sm_89;sm_90a;sm_100a;sm_121a]`.
+Both arms are the SAME binary, and every arm ran with the page cache dropped
+first (`echo 3 > /proc/sys/vm/drop_caches`, ~90.9 GB available after each drop).
+
+### The harness, and why it is not `vllm-cli`
+
+`benchmarks/expert_stream_device_w0e.cpp`, a thin client of `include/vllm.h`
+linked against the packaged shared library, which exports the C ABI and nothing
+else. **The project builds it**, as the `expert-stream-device-w0e` target. It
+was first written unwired, beside `marlin_moe_standalone.cpp`, on the reading
+that a gate instrument is not a shipped capability. Wiring it found that the
+file did not compile under the project's own flags at all — three backslash
+continuations inside `//` comments, which `-Werror=comment` rejects — so the
+recorded recipe was the only thing that had ever built it. An instrument
+nothing compiles rots against the very ABI it measures, and a measurement whose
+harness no longer builds cannot be reproduced, so the target is the correct
+shape even though the file is not a capability. It needed to exist because the gate wants three things from ONE
+generation, and no shipped command produces all three: the generated TOKEN IDS
+(`vllm_complete_tokens`, ABI v13, which writes them into a caller buffer), a
+PER-STEP arrival timestamp (a custom logits processor, invoked once per decode
+step), and the expert-stream counters at both ends of the run
+(`VT_MOE_EXPERT_STREAM_STATS_EVERY=1`, so the periodic line prints at step 1 and
+every step after; at the default 16 a 32-token run prints two lines and a
+5-token one prints none, which is indistinguishable from a dead lane).
+
+The logits processor is a PURE OBSERVER: it records and returns without touching
+`logits`, so the argmax the sampler takes is the one it would have taken without
+it. Its own `token_ids` view is recorded but NOT used for the gate, because
+`tests/capi/test_capi.cpp` states that under the async scheduler that view can
+lag the emitted tokens.
+
+The prompt is fixed as TOKEN IDS rather than as text, `760,6511,314,9338,369`
+("The capital of France is", produced by `build/examples/tokenize` from the
+checkpoint's own GGUF metadata), so both arms are fed byte-identical input and
+no tokenizer step sits between the two.
+
+### The CPU arm: measured, reproduced, and it replaces the VOID #912 F1 figure
+
+`--device cpu`, `VT_GGUF_PREFAULT=0 VT_MOE_EXPERT_STREAM=1`, `--max-num-seqs 1`,
+greedy, 32 tokens. Four runs, two slot counts, two reps each:
+
+| slots | arena | load | TTFT (step 1) | steady decode s/token, steps 4-32 | total generate | peak RSS | min avail | peak swap |
+|---|---|---|---|---|---|---|---|---|
+| 4000 | 9.28 GiB | 271.1 s | 85.90 s | median **11.22**, min 9.62, max 12.51 | 502.3 s | 86.5 GiB | 16,840 MiB | not sampled |
+| 4000 | 9.28 GiB | 255.7 s | 79.09 s | median **11.05**, min 9.43, max 13.25 | 460.7 s | 86.5 GiB | 16,347 MiB | 6,883 MiB |
+| 8000 | 18.55 GiB | 261.6 s | 94.25 s | median 45.40, min 20.73, max 82.68 | 1643.2 s | 86.6 GiB | 6,985 MiB | not sampled |
+| 8000 | 18.55 GiB | 266.5 s | 132.74 s | median 39.98, min 16.06, max 94.30 | 1581.6 s | 86.6 GiB | 6,941 MiB | **30,625 MiB (all of it)** |
+
+**Read the 8000-slot rows as a memory result, not a cache result.** A bigger
+cache came out 3.6x SLOWER and 4x more variable, which is the wrong direction,
+and the last column says why: an 18.55 GiB arena does not fit beside this
+model's 62 GiB of dense weights on a 119.63 GiB box, so those runs swapped, one
+of them exhausting swap entirely. Both 8000-slot runs reproduce each other, so
+this is the box's behaviour and not a fluke. **The publishable CPU figure is the
+4000-slot one, 11.05 s/token**, and it is the first live-cache streaming-ON
+decode number this project has; `docs/BENCHMARKS.md:8` recorded streaming-ON as
+VOID (#912 F1, the step clock dead from token 3) with a re-measure owed.
+
+Steps 1-3 are excluded from the steady figure and reported separately because
+they are not steady state: step 1 is prefill, and steps 2-3 are still filling a
+cold cache (85.90, 80.38, 16.06 s, then 11.7 and below).
+
+**G0-LIVE, gated on the DIFFERENCE and never the total.** At 4000 slots the
+after-prefill snapshot is `steps=1 hits=0 misses=10074 evictions=0 fills=4000
+bytes=9961472000 exhausted=6074` and the final one is `steps=32 hits=37096
+misses=58538 evictions=48464 fills=52464 bytes=130654666752 exhausted=6074`, so
+the decode-phase `exhausted` delta is **0** over 31 decode steps. At 8000 slots
+the same pair reads 2074 and 2074, delta **0**. Both reps of each slot count
+produce byte-identical counters, so the lane is deterministic. `steps=32 > 0`.
+`forced` is 0 by construction and is deliberately absent from the stderr line:
+its only setter is `detail::ExpertStreamSetForceFallback`, which
+`src/vllm/model_executor/models/qwen3_5_internal.h:488` records as having no
+production caller, and `qwen3_5.cpp:5389` counts it separately from `exhausted`
+for exactly that reason.
+
+**The tokens.** All four CPU runs, across both slot counts, produced the same 32
+ids:
+
+```
+11751,13,11751,369,264,3177,7172,303,279,17631,919,314,9338,11,383,279,
+181474,10629,13,1049,369,279,7526,3177,303,9338,321,369,3750,364,1141,25438
+```
+
+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 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
+conditional bound removes that refusal when the lane serves the towers, and the
+load now completes in 255-272 s with 61.20 GiB resident. The lane then engages:
+the `[expert-stream] ON slots=... resident=... GiB` banner prints on the device
+arm, which is W0c's whole point.
+
+That banner is also the first production evidence that **W0b's CUDA leg is
+reached and answers true on real hardware**. Neither the banner nor the
+successful load is reachable unless `host_memory_is_device_addressable()`
+returned true from the CUDA platform, so the leg no longer rests on the fake
+platform in `test_expert_stream_device_slot` alone. The negative — a mutation
+that makes it answer false — is still owed and is in `## Owed`.
+
+**Then the first forward exhausts the machine. Seven attempts, zero decode
+steps, every one of them.**
+
+| slots | arena | prompt tokens | load | RSS after load | decode steps | peak system `used` | peak swap used |
+|---|---|---|---|---|---|---|---|
+| 8000 | 18.55 GiB | 5 | 267.2 s | 61.20 GiB | 0 | 100,215 MiB | not sampled |
+| 4000 | 9.28 GiB | 5 | 271.6 s | 61.20 GiB | 0 | 94,737 MiB | not sampled |
+| 4000 | 9.28 GiB | 5 | 255.3 s | 61.20 GiB | 0 | 120,351 MiB | 30,569 MiB |
+| 3500 | 8.12 GiB | 1 | 262.8 s | 61.20 GiB | 0 | 120,296 MiB | 28,138 MiB |
+| 3500 | 8.12 GiB | 1 | 265.3 s | 61.20 GiB | 0 | 120,347 MiB | 30,461 MiB |
+| 3500 | 8.12 GiB | 1 | 272.3 s | 61.20 GiB | 0 | 120,306 MiB | 30,172 MiB |
+| **64** | **0.15 GiB** | 1 | 268.9 s | 61.20 GiB | 0 | 118,257 MiB | 30,211 MiB |
+
+Each run was stopped by a guard that kills the container when MemAvailable plus
+SwapFree falls under a floor, because an out-of-memory kill on GB10 takes the
+whole machine down rather than the process. The guard is not what makes these
+runs fail: the last two rows above were still climbing at roughly 10 GB of swap
+per minute with under 600 MiB of swap left, so the kernel was seconds from the
+kill the guard exists to prevent.
+
+**The expert lane is not the cost, and that is measured rather than argued.**
+The 64-slot row is the control: a **0.15 GiB** arena dies in the same place an
+**18.55 GiB** one does, 124x smaller and no further along. The in-place tower
+fallback is not the cost either: a 1-token prompt has a protected set of
+`93 x 3 x 10 = 2,790` slices, which fits 3500 slots with no fallback taken at
+all, and it behaves exactly like the 5-token prompt whose 13,950-slice set
+cannot fit any arena on this box.
+
+**The growth is ANONYMOUS, not file-backed**, so the GPU is not pinning the
+mapping's pages through its address translation — which was the first
+hypothesis, and it is wrong. Sampling the container process every 5 s through
+the load and into the forward:
+
+```
+RssAnon:   8.1 -> 13.9 -> 25.9 -> 38.3 -> 49.4 -> 61.4 GB     (through the load)
+RssFile:   0.1-0.2 GB throughout, 4.6 GB at the very end
+VmSwap:    0 -> 5.4 -> 11.4 -> 16.5 -> 27.1 -> 31.0 GB        (inside the forward)
+```
+
+Host `RssAnon` plus `VmSwap` reaches about 65 GB while the system reports
+~119 GiB `used`, and the ~42 GiB difference is device memory that this unified
+part does not charge to the process RSS. **So the non-expert weights are
+resident twice** — once as the host-side `OwnedTensor` and once as the
+`ResidentWeight` device staging copy — and on a box where device memory IS host
+memory, that doubling is what runs it out.
+
+[`expert-streaming.md`](expert-streaming.md) already measured the host half of
+this: the GDN V-head reorder makes `attn_qkv` and `ssm_out`
+`kTransformedWeight`, which expands them from about 5.5 bits to bf16, roughly
+**50 GiB** of the 61.20. The CPU arm pays that once and serves. The CUDA arm
+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`.
+
+### What was running beside the measurement
+
+The lease excluded every other `rc` job for its duration. Two things were not
+excluded and are recorded rather than assumed away. An orphaned
+`VLLM::EngineCore` process from an earlier session held 3.32 GiB of host RSS for
+the entire window; it is not this row's process and it was left alone. And the
+FIRST arm started at 22:13, seconds after the previous lease holder's four-hour
+render released the device, with the one-minute load average still at 17.5 —
+which is one of the two reasons the first 8000-slot run is the noisiest of the
+four, the other being the swap it drove. Every later arm started from a quiet
+box.
+
 ## Dependencies
 
 | Dependency | Shape |
@@ -460,7 +663,9 @@ re-derived here.
 |---|---|
 | **`kQwen3MoeFactory.streams_routed_experts = true` is a CORRECT declaration that nothing READS today.** The flag's only reader is the loader's lane block, which is on the GGUF path, and `kGgufArchArms` (`model_loader.cpp`) maps no `general.architecture` onto `Qwen3MoeForCausalLM` (Qwen3-Coder), so no GGUF load can resolve to that factory. | It is set anyway because it is TRUE: `qwen3_moe.cpp` composes the same `RunMoeBlock` the Qwen3.5 MoE forward does, which is why it holds an `EndStepGuard` at all, so its experts do reach `KqExpertSlice`. Declaring it false to make every setting reachable would put a false statement in the registry, and the safe-direction default would then hide it. Named here per `## Nothing lands dead` rather than left for the next reader to find: `ENG-EXPERT-STREAM-DEVICE` owns the wiring under [#1124](https://github.com/mudler/vllm.cpp/issues/1124), and the flag becomes read the moment a `qwen3moe` GGUF arch arm exists. The `Qwen3_5Moe*` setting beside it IS read and IS gated (`test_gguf_device_fit_reach`, mutation M-A3). |
 | **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 CPU arm's streaming decode figure is still VOID.** `docs/BENCHMARKS.md:8` records it as VOID (#912 F1) with a re-measure owed. | Owned by `ENG-EXPERT-STREAM` and arranged separately by the operator. It is the DENOMINATOR for G0-SPEED, not a precondition for G0-CORRECT or G0-LIVE. |
+| **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 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, two reps, decode-phase `exhausted` delta 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. |
 | **`.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/benchmarks/expert_stream_device_w0e.cpp b/benchmarks/expert_stream_device_w0e.cpp
new file mode 100644
index 000000000..88b45bb3d
--- /dev/null
+++ b/benchmarks/expert_stream_device_w0e.cpp
@@ -0,0 +1,286 @@
+// w0e_gen — ENG-EXPERT-STREAM-DEVICE W0e measurement harness (issue #1124).
+//
+// A THIN CLIENT of the public C ABI (include/vllm.h) only: vllm_engine_load +
+// vllm_complete, with a custom logits processor used as a PURE OBSERVER. The
+// processor edits nothing, so the argmax the sampler takes is byte-identical to
+// a run without it; it exists because it is the only public surface that hands
+// back the request's GENERATED TOKEN IDS, which is what G0-CORRECT gates on,
+// and because it is invoked once per decode step, which is what gives the
+// per-token arrival deltas G0-SPEED needs.
+//
+// The AUTHORITATIVE id list comes from vllm_complete_tokens (ABI v13), which
+// fills a caller buffer with the generated ids. The processor's own token_ids
+// view is recorded too, but it is NOT trusted for the gate: the capi suite
+// records that under the async scheduler that view can LAG the emitted tokens,
+// because the bookkeeping is fed back by update_from_output. Pass --prompt-ids
+// to take that path; without it the harness falls back to vllm_complete, whose
+// only id evidence is the processor's lagging view.
+//
+// It builds as the `expert-stream-device-w0e` target, and links `vllm::shared`
+// rather than `vllm::vllm`, so it exercises the packaged C ABI and nothing else
+// and cannot reach an internal header even by accident. It was originally left
+// out of CMake beside marlin_moe_standalone.cpp, on the reading that a
+// gate instrument is not a shipped capability. That is the wrong trade for THIS
+// file: an instrument nothing compiles rots silently against the very ABI it
+// measures, and the numbers in `.agents/benchmark-record.md` cannot be
+// reproduced from a file that no longer builds. So the project builds it.
+//
+//   cmake -S . -B build -G Ninja -DCMAKE_BUILD_TYPE=Release
+//     -DVLLM_CPP_CUDA=ON -DVLLM_CPP_CUDA_ARCHITECTURES=121a
+//     -DVLLM_CPP_CUTLASS_DIR=$HOME/cutlass -DVLLM_CPP_TRITON=ON
+//   cmake --build build --target expert-stream-device-w0e -j 6
+//
+// One arm, with the lane on and its statistics line at every step so that the
+// after-prefill snapshot and the final one are both in the log. (Written on one
+// line: a backslash continuation inside a `//` comment is what -Werror=comment
+// rejects, and this file is compiled with the project's flags now.)
+//
+//   VT_GGUF_PREFAULT=0 VT_MOE_EXPERT_STREAM=1 VT_MOE_EXPERT_STREAM_SLOTS=8000
+//   VT_MOE_EXPERT_STREAM_STATS_EVERY=1 ./build/examples/expert-stream-device-w0e
+//   --model <shard-1>.gguf --device cuda --max-tokens 32 --max-num-seqs 1
+//   --prompt-ids 760,6511,314,9338,369
+//
+// `VT_GGUF_PREFAULT=0` is load-bearing for a model larger than memory, and
+// STATS_EVERY=1 is load-bearing for a short run: at the default 16 a healthy
+// 32-token run prints no periodic line at all, which looks exactly like a dead
+// lane.
+#include "vllm.h"
+
+#include <chrono>
+#include <cstdint>
+#include <cstdio>
+#include <cstdlib>
+#include <cstring>
+#include <string>
+#include <vector>
+
+namespace {
+
+struct Obs {
+  std::vector<double> t;                    // epoch seconds at each LP call
+  std::vector<std::vector<int32_t>> ids;    // ids seen at each LP call
+};
+
+double Now() {
+  return std::chrono::duration<double>(
+             std::chrono::system_clock::now().time_since_epoch())
+      .count();
+}
+
+// PURE OBSERVER: records and returns. It must not touch `logits`.
+void ObserveLp(const int32_t* token_ids, int32_t n, float* logits,
+               int32_t vocab, void* ud) {
+  (void)logits;
+  (void)vocab;
+  Obs* o = static_cast<Obs*>(ud);
+  o->t.push_back(Now());
+  o->ids.emplace_back(token_ids, token_ids + (n > 0 ? n : 0));
+  std::fprintf(stderr, "[w0e] lp_call=%zu t=%.6f n_ids=%d\n", o->t.size(),
+               o->t.back(), n);
+  std::fflush(stderr);
+}
+
+long ProcKb(const char* key) {
+  std::FILE* f = std::fopen("/proc/self/status", "r");
+  if (f == nullptr) return -1;
+  char line[512];
+  long v = -1;
+  while (std::fgets(line, sizeof(line), f) != nullptr) {
+    if (std::strncmp(line, key, std::strlen(key)) == 0) {
+      v = std::atol(line + std::strlen(key));
+      break;
+    }
+  }
+  std::fclose(f);
+  return v;
+}
+
+void PrintMem(const char* tag) {
+  std::printf("W0E_MEM %s VmRSS_kB=%ld VmHWM_kB=%ld\n", tag, ProcKb("VmRSS:"),
+              ProcKb("VmHWM:"));
+  std::fflush(stdout);
+}
+
+const char* NextArg(int argc, char** argv, int& i) {
+  if (i + 1 >= argc) {
+    std::fprintf(stderr, "w0e_gen: missing value after %s\n", argv[i]);
+    std::exit(2);
+  }
+  return argv[++i];
+}
+
+}  // namespace
+
+int main(int argc, char** argv) {
+  std::string model, prompt = "The capital of France is", tag = "run";
+  std::vector<int32_t> prompt_ids;
+  int32_t device = 0;  // 0 auto, 1 cpu, 2 cuda
+  int max_tokens = 33;
+  long long kv_bytes = 0;
+  int max_num_seqs = 1;
+
+  for (int i = 1; i < argc; ++i) {
+    const std::string f = argv[i];
+    if (f == "--model") {
+      model = NextArg(argc, argv, i);
+    } else if (f == "--prompt") {
+      prompt = NextArg(argc, argv, i);
+    } else if (f == "--tag") {
+      tag = NextArg(argc, argv, i);
+    } else if (f == "--prompt-ids") {
+      const std::string s = NextArg(argc, argv, i);
+      size_t p = 0;
+      while (p < s.size()) {
+        size_t q = s.find_first_of(", ", p);
+        if (q == std::string::npos) q = s.size();
+        if (q > p) prompt_ids.push_back(std::atoi(s.substr(p, q - p).c_str()));
+        p = q + 1;
+      }
+    } else if (f == "--max-tokens") {
+      max_tokens = std::atoi(NextArg(argc, argv, i));
+    } else if (f == "--kv-cache-memory") {
+      kv_bytes = std::atoll(NextArg(argc, argv, i));
+    } else if (f == "--max-num-seqs") {
+      max_num_seqs = std::atoi(NextArg(argc, argv, i));
+    } else if (f == "--device") {
+      const std::string d = NextArg(argc, argv, i);
+      if (d == "auto") {
+        device = 0;
+      } else if (d == "cpu") {
+        device = 1;
+      } else if (d == "cuda") {
+        device = 2;
+      } else {
+        std::fprintf(stderr, "w0e_gen: bad --device %s\n", d.c_str());
+        return 2;
+      }
+    } else {
+      std::fprintf(stderr, "w0e_gen: unknown flag %s\n", f.c_str());
+      return 2;
+    }
+  }
+  if (model.empty()) {
+    std::fprintf(stderr, "w0e_gen: --model required\n");
+    return 2;
+  }
+
+  std::printf("W0E_TAG=%s\nW0E_DEVICE=%d\nW0E_ABI=%d\nW0E_MAX_TOKENS=%d\n",
+              tag.c_str(), static_cast<int>(device),
+              static_cast<int>(vllm_abi_version()), max_tokens);
+  std::printf("W0E_PROMPT=%s\n", prompt.c_str());
+  std::fflush(stdout);
+  PrintMem("start");
+
+  vllm_model_params mp = vllm_model_params_default();
+  mp.model_path = model.c_str();
+  mp.device = device;
+  mp.max_num_seqs = max_num_seqs;
+  // 0.0 is the ABI's "unset" spelling; vllm-cli does the same so a plain run
+  // does not look like an explicit ask (#1165).
+  mp.gpu_memory_utilization = 0.0;
+  if (kv_bytes > 0) mp.kv_cache_memory_bytes = kv_bytes;
+
+  const double t_load0 = Now();
+  std::printf("W0E_LOAD_START=%.6f\n", t_load0);
+  std::fflush(stdout);
+  vllm_engine* eng = nullptr;
+  vllm_status st = vllm_engine_load(&mp, &eng);
+  const double t_load1 = Now();
+  if (st != VLLM_OK) {
+    std::printf("W0E_LOAD_STATUS=%d\nW0E_LOAD_ERROR=%s\n", static_cast<int>(st),
+                vllm_last_error());
+    std::printf("W0E_LOAD_SECS=%.3f\nW0E_RESULT=LOAD_FAILED\n",
+                t_load1 - t_load0);
+    std::fflush(stdout);
+    return 1;
+  }
+  std::printf("W0E_LOAD_SECS=%.3f\nW0E_LOAD_END=%.6f\n", t_load1 - t_load0,
+              t_load1);
+  std::fflush(stdout);
+  PrintMem("after_load");
+
+  Obs obs;
+  vllm_sampling_params sp = vllm_sampling_params_default();
+  sp.temperature = 0.0f;  // greedy
+  sp.max_tokens = max_tokens;
+  sp.logits_processor = &ObserveLp;
+  sp.logits_processor_user_data = &obs;
+
+  vllm_completion out{};
+  std::vector<int32_t> out_ids(256);
+  int32_t n_out_ids = 0;
+  const bool tokens_path = !prompt_ids.empty();
+  std::printf("W0E_ENTRY=%s\n", tokens_path ? "vllm_complete_tokens" : "vllm_complete");
+  if (tokens_path) {
+    std::printf("W0E_PROMPT_IDS_N=%zu\nW0E_PROMPT_IDS=", prompt_ids.size());
+    for (size_t i = 0; i < prompt_ids.size(); ++i) {
+      std::printf("%s%d", i == 0 ? "" : ",", prompt_ids[i]);
+    }
+    std::printf("\n");
+  }
+  std::fflush(stdout);
+  const double t_gen0 = Now();
+  if (tokens_path) {
+    st = vllm_complete_tokens(eng, prompt_ids.data(),
+                              static_cast<int32_t>(prompt_ids.size()), &sp,
+                              out_ids.data(),
+                              static_cast<int32_t>(out_ids.size()), &n_out_ids,
+                              &out);
+  } else {
+    st = vllm_complete(eng, prompt.c_str(), &sp, &out);
+  }
+  const double t_gen1 = Now();
+  std::printf("W0E_GEN_STATUS=%d\nW0E_GEN_SECS=%.3f\n", static_cast<int>(st),
+              t_gen1 - t_gen0);
+  if (st != VLLM_OK) {
+    std::printf("W0E_GEN_ERROR=%s\nW0E_RESULT=GEN_FAILED\n", vllm_last_error());
+    std::fflush(stdout);
+    vllm_engine_free(eng);
+    return 1;
+  }
+  std::printf("W0E_PROMPT_TOKENS=%d\nW0E_COMPLETION_TOKENS=%d\nW0E_FINISH=%s\n",
+              out.prompt_tokens, out.completion_tokens,
+              out.finish_reason != nullptr ? out.finish_reason : "(none)");
+  std::printf("W0E_TEXT_BEGIN\n%s\nW0E_TEXT_END\n",
+              out.text != nullptr ? out.text : "");
+
+  // Per-step arrival: LP call k fires when step k's logits are ready, i.e. at
+  // the moment token k becomes available. k == 1 is the end of prefill (TTFT).
+  std::printf("W0E_LP_CALLS=%zu\n", obs.t.size());
+  for (size_t k = 0; k < obs.t.size(); ++k) {
+    const double prev = (k == 0) ? t_gen0 : obs.t[k - 1];
+    std::printf("W0E_STEP k=%zu t=%.6f dt=%.6f n_ids=%zu\n", k + 1, obs.t[k],
+                obs.t[k] - prev, obs.ids[k].size());
+  }
+  // AUTHORITATIVE: the ids vllm_complete_tokens wrote.
+  if (tokens_path) {
+    std::printf("W0E_OUT_IDS_N=%d\nW0E_OUT_IDS=", n_out_ids);
+    for (int32_t i = 0; i < n_out_ids; ++i) {
+      std::printf("%s%d", i == 0 ? "" : ",", out_ids[static_cast<size_t>(i)]);
+    }
+    std::printf("\n");
+  }
+  // ADVISORY: the longest id prefix the processor's view showed. Recorded so a
+  // reader can see the lag rather than infer it; never the gate's evidence.
+  if (!obs.ids.empty()) {
+    size_t best = 0;
+    for (size_t k = 0; k < obs.ids.size(); ++k) {
+      if (obs.ids[k].size() >= obs.ids[best].size()) best = k;
+    }
+    const std::vector<int32_t>& ids = obs.ids[best];
+    std::printf("W0E_LP_IDS_N=%zu\nW0E_LP_IDS=", ids.size());
+    for (size_t i = 0; i < ids.size(); ++i) {
+      std::printf("%s%d", i == 0 ? "" : ",", ids[i]);
+    }
+    std::printf("\n");
+  }
+  PrintMem("after_gen");
+  std::printf("W0E_RESULT=OK\n");
+  std::fflush(stdout);
+
+  vllm_completion_free(&out);
+  vllm_engine_free(eng);
+  PrintMem("after_free");
+  std::fflush(stdout);
+  return 0;
+}
diff --git a/docs/BENCHMARKS.md b/docs/BENCHMARKS.md
index 2b4603beb..d05c07593 100644
--- a/docs/BENCHMARKS.md
+++ b/docs/BENCHMARKS.md
@@ -6,7 +6,7 @@
 |---|---|---|---|
 | **Weight load (`ENG-LOAD-DIRECT-UPLOAD`, #150)** | Qwen3.6-27B bf16 (50.098 GiB), GB10 Vulkan, same binary both arms, interleaved under one GPU lock | **Load phase 1.54x warm** (19.27 -> 12.48 s), **1.61x cold** (52.62 -> 32.75 s); bytes moved **100.196 -> 81.260 GiB**. Every ON leg beat every OFF leg | byte-identical; 6/6 token-exact |
 | **Expert streaming vehicle (`ENG-EXPERT-STREAM`, [#912](https://github.com/mudler/vllm.cpp/issues/912))** | Qwen3.8-2.4T-A95B `UD-Q1_0` (370 GiB) on ONE GB10, 119 GiB, c1 | **Loads and generates** (streaming OFF): resident **62 GiB**, decode **66.7 s/tok**. Experts BORROW the mmap; ~6.7 GB/token as 4 KiB faults, **~50x** off NVMe. Streaming-ON decode **VOID** (#912 F1); re-measure owed | correct answer; no oracle runs this |
-| **Device expert slices (`ENG-EXPERT-STREAM-DEVICE`, [#1124](https://github.com/mudler/vllm.cpp/issues/1124))** | Qwen3.8-2.4T-A95B `UD-Q1_0`, `--device cuda` vs `cpu`, one GB10, greedy, 32 tokens | Decode **PENDING**: W0e queued, no floor. W0a **MEASURED** (`dgx:gpu0`, `PAGEABLE_OK`): pageable=1, integrated=1, a kernel read+wrote a 2,490,368 B host-vector slot, 2.06-2.28x. [record](../.agents/benchmark-record.md) | 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, 2026-08-18, [record](../.agents/benchmark-record.md) | CPU **MEASURED**, replacing the VOID #912 F1 figure: **11.05 s/token** steady, 2 reps, live cache. CUDA has no number: it loads, then dies in forward 1, 0 steps (#1299). W0a probe `PAGEABLE_OK` | correct answer; no oracle runs this |
 | **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/STATUS.md b/docs/STATUS.md
index 9a94b3b55..731e24d4f 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) `<stem>.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 `<stem>.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<kIQ2_S/kMXFP4>`), 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<WType::kIQ2_S>` 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 measured `PAGEABLE_OK` on `dgx:gpu0`; W0b-W0d implemented and unit-gated; W0e still queued, so no decode number exists and none is claimed | A 369.96 GiB GGUF refuses on `--device cuda` against a 119.631 GiB pool (#1123). W0 serves `*_exps` slices from the HOST slot store, keyed on a PROBED `host_memory_is_device_addressable()` |
+| 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 | W0 removed the #1123 load refusal by serving `*_exps` from the HOST slot store on a PROBED `host_memory_is_device_addressable()`. What blocks it now is the DENSE half, not the arena: 0.15 GiB dies where 18.55 GiB does |
 | 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 462a7f1bc..63bbb6612 100644
--- a/docs/USAGE.md
+++ b/docs/USAGE.md
@@ -4350,13 +4350,17 @@ Three limits, stated plainly rather than left to be discovered.
   stage those towers, so charging the device for a slot arena instead would
   under-count what the load really needs and turn a correct refusal into an
   out-of-memory first forward.
-* **No speed claim is attached.** The decode measurement on the one machine that
-  answers capable has not run at the time of writing; `docs/BENCHMARKS.md`
-  carries it as `PENDING` and records what HAS been measured there, which is the
-  device probe itself. Device access to host-resident weights on that part 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 is a real possible outcome. Read
-  the benchmarks file before assuming the GPU is the faster arm here.
+* **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.
+  Read `docs/BENCHMARKS.md` before assuming the GPU is the faster arm here.
 
 ### The same thing as config, and which one wins
 
diff --git a/examples/CMakeLists.txt b/examples/CMakeLists.txt
index 1e57bf5e4..8b376377a 100644
--- a/examples/CMakeLists.txt
+++ b/examples/CMakeLists.txt
@@ -361,6 +361,21 @@ target_link_libraries(vocoder-conv-ab PRIVATE vllm::vllm)
 target_include_directories(vocoder-conv-ab PRIVATE ${CMAKE_SOURCE_DIR}/src)
 vllm_cpp_set_warnings(vocoder-conv-ab)
 
+# expert-stream-device-w0e: the W0e measurement harness for row
+# ENG-EXPERT-STREAM-DEVICE (#1124). It exists because one generation has to
+# yield three things together that no shipped command produces at once: the
+# generated token ids, a per-step arrival timestamp, and the expert-stream
+# counters at both ends of the run — and a live gate needs the DECODE-phase
+# delta of those counters, not the total, because prefill exhaustion here is
+# structural. Unlike the A/B harnesses above it is a pure client of
+# `include/vllm.h` and links `vllm::shared`, so it stays inside the ABI-client
+# rule and needs no `src/` include path. Not gated on CUDA: the CPU arm is the
+# baseline the CUDA arm is compared against.
+add_executable(expert-stream-device-w0e
+               ${CMAKE_SOURCE_DIR}/benchmarks/expert_stream_device_w0e.cpp)
+target_link_libraries(expert-stream-device-w0e PRIVATE vllm::shared)
+vllm_cpp_set_warnings(expert-stream-device-w0e)
+
 # vulkan-gemm-ab: the VK-C tactic A/B — cooperative matrix vs the portable scalar
 # GEMM, SAME BINARY with VT_VULKAN_COOPMAT as the only variable. Built only in a
 # Vulkan build, because it links the Vulkan context directly to report which