From c66c24c6bfaef1382585ccae74f353a57edf65a1 Mon Sep 17 00:00:00 2001 From: Richard Palethorpe Date: Sun, 9 Aug 2026 18:36:24 +0100 Subject: [PATCH 01/17] spec: define sm120 Qwen3.5 Pareto campaign Import the measured optimization history and bind issue #206 to a no-regression campaign for throughput, latency, and memory. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: Codex:gpt-5.6-sol [Codex] --- .agents/NOW.md | 5 +- .agents/benchmark-record.md | 412 ++++++++++++++++++ .agents/kernel-matrix.md | 11 + .agents/roadmap_v1.md | 1 + .../sm120-qwen35-conv-block256-2026-08-08.md | 151 +++++++ ...120-qwen35-conv-channel-tile-2026-08-08.md | 173 ++++++++ ...decode-bf16-vector-writeback-2026-08-09.md | 132 ++++++ ...0-qwen35-gdn-decode-bv-sweep-2026-08-09.md | 186 ++++++++ ...sm120-qwen35-gdn-decode-bv16-2026-08-08.md | 192 ++++++++ ...en35-gdn-decode-direct-store-2026-08-09.md | 185 ++++++++ ...sm120-qwen35-gdn-decode-regk-2026-08-09.md | 213 +++++++++ ...0-qwen35-gdn-decode-regstate-2026-08-09.md | 277 ++++++++++++ ...sm120-qwen35-gdn-decode-rpt2-2026-08-09.md | 126 ++++++ ...20-qwen35-gdn-decode-swizzle-2026-08-09.md | 139 ++++++ .../specs/sm120-qwen35-pareto-2026-08-09.md | 130 ++++++ ...0-qwen35-postconv-token-tile-2026-08-08.md | 170 ++++++++ docs/BENCHMARKS.md | 2 +- docs/STATUS.md | 2 +- 18 files changed, 2502 insertions(+), 5 deletions(-) create mode 100644 .agents/specs/sm120-qwen35-conv-block256-2026-08-08.md create mode 100644 .agents/specs/sm120-qwen35-conv-channel-tile-2026-08-08.md create mode 100644 .agents/specs/sm120-qwen35-gdn-decode-bf16-vector-writeback-2026-08-09.md create mode 100644 .agents/specs/sm120-qwen35-gdn-decode-bv-sweep-2026-08-09.md create mode 100644 .agents/specs/sm120-qwen35-gdn-decode-bv16-2026-08-08.md create mode 100644 .agents/specs/sm120-qwen35-gdn-decode-direct-store-2026-08-09.md create mode 100644 .agents/specs/sm120-qwen35-gdn-decode-regk-2026-08-09.md create mode 100644 .agents/specs/sm120-qwen35-gdn-decode-regstate-2026-08-09.md create mode 100644 .agents/specs/sm120-qwen35-gdn-decode-rpt2-2026-08-09.md create mode 100644 .agents/specs/sm120-qwen35-gdn-decode-swizzle-2026-08-09.md create mode 100644 .agents/specs/sm120-qwen35-pareto-2026-08-09.md create mode 100644 .agents/specs/sm120-qwen35-postconv-token-tile-2026-08-08.md diff --git a/.agents/NOW.md b/.agents/NOW.md index 4c3239905..4b13b45ee 100644 --- a/.agents/NOW.md +++ b/.agents/NOW.md @@ -21,7 +21,7 @@ Work: exact-chunks on main `1ce0d662b`; sm_120 measured at `3d2581551`. | MiniMax-H3 | **PRUNED ckpts RUN (#241): Q8_0 renders, seam 0.9941** | same-binary A/B | | Kimi-Linear-48B | 122/128 held; grouped router parallelised, e2e NOT ESTABLISHED | ckpt is tiktoken-only: no warm server | | 35B binding grid | @`a0fa12c7`: **flat 0.935-0.979x, NO c2/c8 weak cell** (CoV <0.81%); mem PSS 3.81x | Attribute the flat ~5% mid-band; TTFT c2 0.872x | -| Qwen3.5-4B sm_120 | Exact chunks ON: 3.072x kernel / +2.272% run; sealed-vLLM tput 1.021x PASS; latency/VRAM OPEN | Spike residual 1.609x conv gap | +| Qwen3.5-4B sm_120 | #206 spec; tput 1.021x PASS; TTFT/TPOT/VRAM OPEN | Transplant opt-ins; combined A/B | | RPi5 A76 CPU | **R5 asm GREEN; llama NOT MET**: 0.461x pf, 0.653x dec | W6: BF16 GEMM | | MXFP4 parity | c1 1.020, c2-c8 0.962-0.969. **#82 CLOSED: ptxas-lineage REFUTED** | TERMINAL: at parity | | SERVE-ASYNC-DENSE-MIRROR | **LANDED+VERIFIED** (`f9c969ae`): async mirror, dense Qwen3; SACRED 184/184 | Sibling scope one-liner | @@ -53,8 +53,7 @@ latency/memory on every axis, both gate models, reproduced 2–3x idle. See of roof. Dense-marlin +0.5%; Triton-AOT GDN a WASH. 2. **Spike the Parakeet encoder row** (vLLM: `nano_nemotron_vl.py`; the transducer half is NOT in vLLM: separate call). -3. **Qwen3.5-4B sm_120:** GREEN + reprofiled. Spike the residual 1.609x conv - gap; latency/VRAM and gate models stay open. +3. **Qwen3.5-4B #206:** transplant opt-ins; combined K4+postconv A/B, then trace. 2. **Merge the invocation-parity prevention** (CI guard + AGENTS.md checklist); CUDA build-verify the byte-exact `kGemvHeuristicAlgos` refactor on dgx. 4. **Restore `local-ai-worker`** on dgx at campaign end (`--restart=always`). diff --git a/.agents/benchmark-record.md b/.agents/benchmark-record.md index 48b408ad9..c1b9663f5 100644 --- a/.agents/benchmark-record.md +++ b/.agents/benchmark-record.md @@ -17619,3 +17619,415 @@ thermal / cache-drop / memory-return). Build contract: RelWithDebInfo, nvcc 4.5.0, TRITON=ON, BENCH_PROFILE_CONTROL=OFF. SACRED `test_qwen36_paged_engine` passed as the harness's own precondition. +## 2026-08-08 — sm_120 fused GDN post-conv 16-token tile: 1.859x kernel, byte-exact + +**Disposition:** IMPLEMENTED as opt-in `VT_GDN_POSTCONV_TOKEN_TILE=1`. +Locally positive and token-safe; default and release-model gates remain open. + +**Selection and falsification.** On exact-chunks `c3bb0f39a`, the accepted +Qwen3.5-4B/c32/1,280-block graph-node trace measured fast megablock +228.150171 ms and the existing per-V-head split 448.364941 ms across 1,728 +calls. Both token files were identical. The split is 1.965x slower and the +enclosing run 0.97% slower, so V-only launch decomposition is rejected. +Pinned vLLM's same-tool `_fused_post_conv_kernel` is 108.034870 ms across 1,923 +calls. Source comparison instead selected vLLM/FLA's 16-token, per-head, +four-warp schedule and Q/K register reuse. + +**Implementation and numerical mutant.** `GdnPostConvTokenTileKernel` maps one +block to `(16 tokens, one Q/K or V head)` and one warp to four tokens. Each Q/K +lane retains features `lane+{0,32,64,96}` across normalization. The first +implementation summed those four squares sequentially before a warp reduction: +kernel time improved 227.731960→122.472980 ms and enclosing throughput +6731.69→6773.85 tok/s, but the tile token SHA +`1d496ff0f989978155d8e900c7a5500a43db26816dead8e035310d0bf9cb9756` +did not match fast `83fcdc45...453545`; REJECTED. Reproducing the current +128-lane tree exactly—`(i+i+64)`, then `(i+i+32)`, then shuffle offsets +16/8/4/2/1—restored byte identity without restoring the reload/barrier costs. + +**Final same-binary profile.** One `/tmp/gpu` lock, 22/25 GiB user-systemd +scope, `--cuda-graph-trace=node`, identical production workload and binary: + +| Axis | fast megablock | token tile | change | +|---|---:|---:|---:| +| post-conv GPU total, 1,728 calls | 227.887066 ms | **122.587027 ms** | **1.858982x faster** | +| mean post-conv call | 131.879 us | **70.942 us** | **46.21% lower** | +| total throughput | 6734.82 tok/s | **6770.62 tok/s** | **+0.532%** | +| output throughput | 744.72 tok/s | **748.68 tok/s** | **+0.532%** | +| TTFT | 1024.14 ms | **1015.43 ms** | **-0.850%** | +| TPOT / ITL | 35.01 ms | **34.85 ms** | **-0.457%** | +| E2E | 5469.87 ms | **5440.81 ms** | **-0.531%** | + +The vLLM kernel residual is now **1.134699x**, down from 2.112x. Final token +files are identical, full SHA-256 +`83fcdc45f79ddb06a634c7d7d95eba3384543b3cd781a45a8db1fc4e2a453545`. +Portable flag/grid tests pass 6/6·50; CUDA GDN passes 67/67·4384 including +partial/exact tiles, packed BA non-zero views/wider row strides, byte-exact five +outputs, finiteness and norms; cached Qwen3.5-4B passes 3/3·1672. + +**Evidence.** Final fast/tile traces +`/tmp/qwen35-postconv-tile-exact-{fast,tile}.nsys-rep`, SHA-256 +`c75e2cb...bcadc` / `a0eb1808...f418`; token files beside them. Rejected +arithmetic traces `/tmp/qwen35-postconv-tile-wip-{fast,tile}.nsys-rep`, SHA-256 +`c2d8872e...b774` / `1c42d489...d66b`. Selection traces and exact recipes are +in [the spike/result](specs/sm120-qwen35-postconv-token-tile-2026-08-08.md). +One local profile is not extrapolated to the unavailable Qwen3.6-27B/35B gates; +the flag therefore remains opt-in. + +## 2026-08-08 — sm_120 causal-conv residual: K=4 specialization wins; 256-channel tile falsified + +**Disposition:** arm 1 is IMPLEMENTED, byte-exact and locally positive behind +`VT_CONV_CHANNEL_TILE=1`; arm 2 (`=2`) is retained as an explicit falsified +experiment. Unset/`0` remains the default pending repeated and release-model +gates. + +**Divide-and-conquer selection.** Fresh current-main/post-conv tracing measured +causal conv 234.255 ms versus pinned-vLLM 145.532 ms, while post-conv's +remaining excess was only 14.476 ms. Grouping causal-conv launches by grid +showed the 279/280-program waves consumed 136.189 ms (58.1%). Local used 64 +feature blocks and runtime width at 43 registers/thread; vLLM used `BLOCK_N=256`, +32 feature blocks and compile-time width at 32 registers/thread. The spike split +those differences into arm 1 (compile-time K=4, unchanged 64 blocks) and arm 2 +(K=4 plus two channels/thread, 32 blocks). + +**Correctness/review finding.** A serial-stripe arm 2 corrupted the second +stripe when an exact final-chunk block wrote state before stripe 2 loaded initial +history; the CUDA matrix caught it and the accepted kernel preloads both stripes. +Fresh mutation review then found byte comparisons could stay green when a whole +specialized dispatch branch was deleted. A fresh fix routes production through +the same portable callback dispatcher the tests mutate. Scoped re-review killed +arm-1 deletion, arm-2 deletion, relaxed arm-2 parsing and removed non-K4 fallback. +Final gates: portable 9/9·88, CUDA GDN 67/67·4631, paged-forward 4/4·8. + +**VOID series.** `/tmp/qwen35-conv-arm{0,1,2}-565a26fcc.*` is invalid for +selection: the test executables had rebuilt but `vllm-bench` had not relinked. +All three traces proved the old runtime kernel/grid/registers ran. This was +caught structurally before timing interpretation. + +**Accepted rebuilt same-binary profile.** One GPU lock, 22/25 GiB user-systemd +scope, exact c32 workload and `--cuda-graph-trace=node`: + +| Axis | arm 0 runtime | arm 1 K4x1 | arm 2 K4x2 | +|---|---:|---:|---:| +| kernel / grid / registers | runtime / 64 / 43 | K4<1> / 64 / 52 | K4<2> / 32 / 58 | +| causal-conv total, 1,728 calls | 234.604587 ms | **219.506425 ms** | 228.400830 ms | +| 279-program mean | 149.546 us | **140.133 us** | 145.586 us | +| 280-program mean | 149.480 us | **139.982 us** | 145.468 us | +| total throughput | 6759.39 tok/s | **6767.62 tok/s** | 6757.19 tok/s | +| output throughput | 747.43 tok/s | **748.34 tok/s** | 747.19 tok/s | +| TTFT | 1016.69 ms | **1013.82 ms** | 1017.61 ms | +| TPOT / ITL | 34.91 ms | **34.88 ms** | 34.91 ms | +| E2E | 5449.87 ms | **5443.16 ms** | 5451.66 ms | + +Arm 1 improves conv **6.4356%**, with every enclosing axis positive but small. +It leaves a **1.5083x** same-tool vLLM conv residual. Register count rises rather +than falls, falsifying the occupancy rationale; compile-time removal of runtime +width work is the supported cause. Arm 2 improves only 2.6443% versus baseline, +is **4.0520% slower than arm 1**, and is neutral/slightly negative end to end: +halving blocks does not repay duplicated channel-local register state on sm_120. + +All accepted token files SHA-256 +`83fcdc45f79ddb06a634c7d7d95eba3384543b3cd781a45a8db1fc4e2a453545`. +Report SHA-256 arm0/arm1/arm2: +`39d383dd878fc340a3cfaaee79a4addcb4eccb181439e9b4725f724f4569a6eb`, +`c8799ac0b4cdf997d383fe8a690b223be882dce3b1ee1a6fff35a62d75f7cf85`, +`3d38793571539864b23688fd9a85966debbf1e7c48fe8a1a2509438a45ee0452`. +Full recipe and decision: +[structured spike/result](specs/sm120-qwen35-conv-channel-tile-2026-08-08.md). + +## 2026-08-09 — sm_120 block-256 causal conv and four-warp GDN decode refuted + +**Block-256 causal conv (`VT_CONV_CHANNEL_TILE=3`) — REJECTED and removed.** +At production head `84c7e23b0`, one rebuilt binary, one GPU lease/lock and the +standard c32 Qwen3.5-4B scope gave byte-identical arm-1/arm-3 token files +(SHA-256 `83fcdc45...453545`). The intended launch change executed: K4 arm 1 +used grid-x 64 / block 128 / 52 registers, while arm 3 used grid-x 32 / block +256 / 52 registers. Nevertheless the selected 279+280 programs regressed +127.645906→142.936163 ms over 912 calls (**+11.9802%**) and all causal conv +regressed 219.400955→245.421363 ms over 1,728 calls (**+11.8598%**). The +enclosing single profile also lost: total/output throughput +6772.76/748.91→6756.52/747.12 tok/s, TTFT 1014.42→1016.46 ms, TPOT/ITL +34.84→34.93 ms and E2E 5439.06→5452.21 ms. Halving block count does not repay +the residency cost of twice as many threads at this register footprint. Per the +spike stop condition, the dormant arm-3 selector, instantiation and tests were +removed; accepted arm 1 and the falsified arm-2 control are unchanged. Traces: +`/tmp/qwen35-conv-block256-arm{1,3}-84c7e23b0.nsys-rep`, SHA-256 +`39d333c5...344b9` / `62ee50dd...3f4e9`. + +**Four-warp fused GDN decode (`VT_GDN_DECODE_NW=4`) — correctness REJECTED, +no timing claim.** The immediately following existing-selector discriminator +changed the production token SHA from default NW8 +`83fcdc45...453545` to `93003ca1...8eebe`: 24/128 requests and 678/16,384 +aligned token positions differed, with equal lengths. Correctness failed before +timing acceptance, so NW8 stays unchanged and no selector code changed. The NW4 +diagnostic-only profile was total/output 6730.55/744.25 tok/s, TTFT 1017.55 ms, +TPOT/ITL 35.08 ms and E2E 5473.25 ms. Trace +`/tmp/qwen35-gdn-decode-nw4-84c7e23b0.nsys-rep`, SHA-256 +`8680008e0261e9aa70bb9e001e2b31165e7999341ae3f963bc0b667caff7fb9a`. +Reducing decode warps is therefore closed for exact greedy inference unless a +separate numerical-order-preserving implementation is spiked. + +## 2026-08-09 — sm_120 fused GDN decode BV16 accepted opt-in + +**Disposition:** ACCEPTED OPT-IN at product commit `92256e6a9`. The exact +`VT_GDN_DECODE_BV=16` selector changes only the independent value tile; unset +keeps BV32. Default and Qwen3.6-27B/35B release gates remain open. + +**Correctness and geometry.** Operator gates pass: portable selector/geometry +3/3 · 49, CUDA GDN 68/68 · 4,699, and Qwen3.5 paged-forward 4/4 · 8. Every +profile and memory-run token file has SHA-256 +`83fcdc45f79ddb06a634c7d7d95eba3384543b3cd781a45a8db1fc4e2a453545`. +The profiler confirms the intended grid-x 4 / block 256 / 52-register BV32 +launch becomes grid-x 8 / block 128 / 52-register BV16, preserving NW8 and its +arithmetic order. + +**Counterbalanced same-binary graph-node series.** Order was +BV32a→BV16a→BV16b→BV32b on the identical 128-request Qwen3.5-4B BF16 c32 +workload: + +| Arm | all GDN decode, 1,704 calls | grid-y 800, 816 calls | total / output tok/s | TTFT | TPOT / ITL | E2E | +|---|---:|---:|---:|---:|---:|---:| +| BV32a | 295.608238 ms | 149.887342 ms | 6763.51 / 747.89 | 1014.64 ms | 34.90 ms | 5446.55 ms | +| BV16a | 265.417816 ms | 133.690094 ms | 6766.91 / 748.27 | 1013.00 ms | 34.89 ms | 5443.79 ms | +| BV16b | 264.152452 ms | 133.030633 ms | 6779.94 / 749.71 | 1011.37 ms | 34.82 ms | 5433.25 ms | +| BV32b | 295.123586 ms | 149.627524 ms | 6763.94 / 747.94 | 1015.48 ms | 34.89 ms | 5446.15 ms | + +Counterbalanced means are **295.365912→264.785134 ms (-10.3535%)** across all +1,704 decode calls and **149.757433→133.360364 ms (-10.9491%)** for the +dominant grid-y 800 shape. That shape falls from 183.526 to 163.432 us/call. +Pinned vLLM remains faster at 128.061 us/call, a local BV16 residual of about +**1.276x**; unequal all-call composition prevents claiming an all-call oracle +ratio. + +**Enclosing and memory reproduction.** A third pair sampled peak memory. +Across all three pairs, BV32→BV16 means are total throughput +6783.437→6791.063 tok/s (**+0.1124%**), output throughput +750.093→750.940 tok/s (**+0.1129%**), TTFT 1013.237→1011.490 ms +(**-0.1724%**), TPOT/ITL 34.7867→34.7500 ms (**-0.1054%**), and E2E +5430.777→5424.630 ms (**-0.1132%**). The memory pair itself is +6822.86/754.45 tok/s, 1009.59 ms TTFT, 34.57 ms TPOT/ITL and 5399.63 ms E2E +for BV32 versus 6826.34/754.84, 1010.10, 34.54 and 5396.85 for BV16. Peak GPU +memory is 13058→13054 MiB and peak PSS 2,273,490→2,170,880 KiB. + +**Evidence.** Traces, exported SQLite databases and tokens are +`/tmp/qwen35-gdn-bv{32a,16a,16b,32b}-92256e6a9.*`; trace SHA-256 values are +`198d1bf843443f69fd1491131a13e7aa95cc136c56e96441aa43c48f267958f5`, +`2f5e47f54faafa70ed8494c3afaa2160e31d311d90e7412cbf9777f2e93e3d18`, +`3b0303a97e2335d5ca2f88eacf5c2e5f063d7d899cb44db3dd7d6a170938d117`, +and `3d7b4a214b14976c95e9e4c944540e91163c7f10e820bb63bf7d0cb9fc9d8b11`. +Memory JSONL/log/token evidence is +`/tmp/qwen35-gdn-bv{32,16}-mem-92256e6a9.*`. The accepted result is local +evidence only; it is not extrapolated to unavailable release models and does +not change the default. + +## 2026-08-09 — sm_120 fused GDN decode BV8/BV24 rejected + +**Disposition:** REJECTED AND REMOVED at measured implementation `e102a14de` +and cleanup `634ccba70`; BV16 remains the sole exact opt-in and BV32 remains +default. Default and Qwen3.6-27B/35B release gates remain open. + +Operator runtime gates pass: portable selector/geometry **3/3 · 100**, CUDA GDN +**69/69 · 4,850** including the public production-graph geometry case, and +Qwen3.5 paged-forward **4/4 · 8**. The prescribed same-binary order was +BV32a→BV16a→BV8a→BV24a→BV24b→BV8b→BV16b→BV32b. Every token file has SHA-256 +`83fcdc45f79ddb06a634c7d7d95eba3384543b3cd781a45a8db1fc4e2a453545`. + +| Arm mean | all fused decode, 1,704 calls | grid-y 800, 816 calls | total / output tok/s | TTFT | TPOT / ITL | E2E | +|---|---:|---:|---:|---:|---:|---:| +| BV32 | 302.7940655 ms | 153.7994265 ms | 6730.700 / 744.265 | 1022.205 ms | 35.050 ms | 5473.235 ms | +| BV16 | 271.5899630 ms | 136.9721265 ms | 6740.010 / 745.290 | 1021.150 ms | 34.995 ms | 5465.690 ms | +| BV8 | 278.4436445 ms | 141.1102230 ms | 6739.720 / 745.255 | 1021.195 ms | 35.000 ms | 5465.840 ms | +| BV24 | 279.7297115 ms | 141.5291235 ms | 6732.840 / 744.500 | 1022.340 ms | 35.035 ms | 5471.455 ms | + +Against BV16, BV8 regresses all fused decode **2.5235%** and the dominant shape +**3.0211%**; BV24 regresses **2.9971%** and **3.3270%** respectively. Both +samples of each candidate have the same losing direction. They fail the ≥1% +dominant-shape win and all-fused non-regression bars, so the enclosing means +are recorded but confer no selection credit. BV16 repeats its win over BV32 in +this series: **-10.3054%** across all fused calls and **-10.9411%** at grid-y +800. + +Profiler launch contracts are exact at Dv=Dk=128 and 52 registers/thread: +BV8 grid-x 16 / block 64 / 5,152 shared bytes; BV16 8 / 128 / 9,280; BV24 6 / +192 / 13,408; BV32 4 / 256 / 17,536. Cleanup removes the losing selectors, +arms, instantiations and expectations; `VT_GDN_DECODE_BV=8` and `=24` fall back +to BV32, while exact `=16` remains opt-in. + +The current BV16 dominant mean is **167.858 us/call** versus pinned vLLM +**128.061 us/call**, about **1.311x** slower. The prior accepted series' 1.276x +ratio remains valid for that run; the difference is run variance, not evidence +of a new oracle or a closed residual. + +Evidence is +`/tmp/qwen35-gdn-bvsweep-{32a,16a,8a,24a,24b,8b,16b,32b}-e102a14de.{nsys-rep,sqlite,log,tokens.json}`. +The per-leg values and eight trace SHA-256 values are preserved in the +[structured result](specs/sm120-qwen35-gdn-decode-bv-sweep-2026-08-09.md). + +## 2026-08-09 — sm_120 fused GDN decode RPT2 rejected and removed + +**Disposition:** REJECTED AND REMOVED. Measured product `6ac8bf390` passed +portable **4/4 · 125**, CUDA GDN **69/69 · 5,046**, and Qwen3.5 paged-forward +**4/4 · 8**. All four production token files have SHA-256 +`83fcdc45f79ddb06a634c7d7d95eba3384543b3cd781a45a8db1fc4e2a453545`. +Cleanup `f6a0c879141a1887f9fcf8c85a9eb3a8cf8a6275` removes the losing selector, specialization and +tests; BV16 remains the sole opt-in and BV32 remains default. + +The intended geometry executed at Dv=Dk=128: BV16 grid-x 8 / block 128 / +9,280 shared bytes / 50 registers, versus RPT2 grid-x 4 / block 128 / 17,536 +shared bytes / 56 registers. Across the counterbalanced +`BV16a -> RPT2a -> RPT2b -> BV16b` series, the 1,704-call all-fused mean +regressed **271.577988 -> 372.4151745 ms (+37.1301%)** and the 816-call y800 +mean regressed **136.9415415 -> 190.121047 ms (+38.8337%)**. + +| Arm mean | all fused decode, 1,704 calls | grid-y 800, 816 calls | total / output tok/s | TTFT | TPOT / ITL | E2E | +|---|---:|---:|---:|---:|---:|---:| +| BV16 | 271.577988 ms | 136.9415415 ms | 6732.58 / 744.47 | 1026.28 ms | 35.00 ms | 5471.38 ms | +| RPT2 | 372.4151745 ms | 190.121047 ms | 6709.16 / 741.88 | 1027.645 ms | 35.145 ms | 5490.87 ms | + +RPT2 fails both micro bars and every enclosing mean despite exact correctness. +BV16's current dominant call is **167.82 us/call** versus pinned vLLM +**128.061 us/call**, about **1.310x** slower. Default/release and the +hardware-unavailable 27B/35B gates remain open. Evidence is +`/tmp/qwen35-gdn-rpt2-{bv16a,rpt2a,rpt2b,bv16b}-6ac8bf390.*`; full disposition +is in the [structured result](specs/sm120-qwen35-gdn-decode-rpt2-2026-08-09.md). + +## 2026-08-09 — sm_120 BV16/NW8 GDN decode shared swizzle accepted opt-in + +**Disposition:** ACCEPTED OPT-IN at spike `bdfaf823e` and product `824370396`. +Operator gates pass: portable **6/6 · 633**, CUDA GDN **69/69 · 4,742**, and +Qwen3.5 paged-forward **4/4 · 8**; parser/mapping negative mutations fail 5/132. +All six token files have SHA-256 +`83fcdc45f79ddb06a634c7d7d95eba3384543b3cd781a45a8db1fc4e2a453545`. +Both arms launch gx8/bx128/reg56; shared bytes change 9,280 -> 9,728. + +| Arm mean | all fused decode, 1,704 calls | grid-y 800, 816 calls | total / output tok/s | TTFT | TPOT / ITL | E2E | +|---|---:|---:|---:|---:|---:|---:| +| BV16 | 289.8737275 ms | 146.959362 ms | 6724.18 / 743.54 | 1027.34 ms | 35.045 ms | 5478.205 ms | +| SWZ | 268.254130 ms | 135.9178535 ms | 6739.86 / 745.275 | 1023.005 ms | 34.98 ms | 5465.725 ms | + +The prescribed `BV16a -> SWZa -> SWZb -> BV16b` series improves all fused +decode **7.4555%** and y800 **7.5123%**, with both samples winning and every +enclosing throughput/latency mean positive. The memory pair is likewise +positive on user-visible/device axes: total 6784.64 -> 6793.05 tok/s, output +750.23 -> 751.16 tok/s, TTFT 1020.77 -> 1018.67 ms, TPOT 34.72 -> 34.68 ms, +E2E 5430.24 -> 5423.44 ms, and GPU 13,060 -> 13,054 MiB. + +Peak host PSS rises **1,909,259 -> 2,038,631 KiB (+6.78%)**, while the supplied +available-memory drop is **1,641,476 -> 1,529,724 KiB**. The isolated PSS move +is recorded as a noisy transient host-load caveat; the candidate stays explicit +opt-in and receives no default/release credit. Accepted swizzle y800 is +**166.566 us/call** versus pinned vLLM **128.061 us/call** (**1.301x** slower). + +Evidence is +`/tmp/qwen35-gdn-swizzle-{bv16a,swza,swzb,bv16b}-8243703.{nsys-rep,sqlite,log,tokens.json}` +and `/tmp/qwen35-gdn-swizzle-{bv16,swz}-mem-8243703.{jsonl,log,tokens.json}`. +Trace SHA-256 prefixes: BV16a `59bb6081`, SWZa `082cd6d4`, SWZb `b44792c8`, +BV16b `633fad31`. Default/release and 27B/35B gates remain open. + +## 2026-08-09 — sm_120 GDN decode register-state measured provisional opt-in + +**Disposition:** `MEASURED/PROVISIONAL` at product commit `7476818c1`. The +strict `VT_GDN_DECODE_REGSTATE=1` specialization retains each lane's 16 F32 +state values across the two recurrence loops while preserving the accepted +BV16/NW8/shared-swizzle mapping. It remains explicit opt-in: exact static SASS +size and NCU attribution are pending external tool/download authority, so this +is not fully `ACCEPTED` and changes no default, release, or 27B/35B claim. + +**Correctness, review and resources.** The full canonical preflight is green; +fresh static and targeted mutation re-review is `PASS`. Operator gates pass: +portable selector/mapping **10/10 · 1,962**, CUDA GDN **70/70 · 4,830**, and +cached Qwen3.5 paged-forward **4/4 · 8**, all observed diffs zero. Every token +file has SHA-256 +`83fcdc45f79ddb06a634c7d7d95eba3384543b3cd781a45a8db1fc4e2a453545`. +The intended graph executes at gx8/gy800/bx128, 9,728 dynamic shared bytes, 56 +registers and zero local bytes. Same-tool ptxas evidence at +`/tmp/cuda_gdn-regstate-candidate-7476818c1.log` reports the exact +`bf16,bf16,float,NW8,SWIZZLED=true,REGSTATE=true` specialization at 56 +registers, one barrier, zero stack and zero spills; the accepted SWZ baseline +also uses 56 registers. The missing exact-specialization static SASS-size +<=5% and shared-instruction count are not inferred from these resource facts. +The retained nvcc 12.9 build `/tmp/cuda-gdn-keep.gXumAL/` supplies additional +non-substitute structure: exact cubin ELF symbol size is 27,008 bytes SWZ versus +26,496 REGSTATE (**-1.896%**), while exact PTX counts are 922/894 total +instructions, 69/53 shared loads and 47/31 shared stores, with unchanged 19/7 +global loads/stores and two barriers. The exact -16/-16 PTX shared operations +support the mechanism and the cubin byte-size direction, but neither is the +required SASS instruction/shared-op count; cuobjdump remains pending. PTX/cubin +SHA-256 values are `dc500503...e2e6` / `5169611e...d913`. + +**Counterbalanced same-binary graph-node series.** The prescribed order was +`SWZa -> REGa -> REGb -> SWZb` on the standard 128-request Qwen3.5-4B BF16 c32 +workload. Both REGSTATE raw legs beat both bracketing SWZ legs on grid-y 800 and +all fused decode. + +| Arm mean | all fused decode, 1,704 calls | grid-y 800, 816 calls | total / output tok/s | TTFT | TPOT / ITL | E2E | +|---|---:|---:|---:|---:|---:|---:| +| SWZ | 268.1141915 ms | 135.7758125 ms | 6726.240 / 743.770 | 1026.800 ms | 35.000 ms | 5471.960 ms | +| REGSTATE | 264.7265645 ms | 134.2911050 ms | 6739.125 / 745.195 | 1021.745 ms | 35.000 ms | 5466.420 ms | + +REGSTATE improves all fused decode **1.263502%** and the dominant y800 shape +**1.093499%**, with total/output throughput **+0.191563%/+0.191591%**, TTFT +**-0.492306%**, unchanged TPOT/ITL, and E2E **-0.101243%**. This is a local +same-binary discriminator, not a new pinned-vLLM ratio. + +**Memory and evidence.** The separate non-profiled memory pair reports peak GPU +allocation 13,058 -> 13,060 MiB and peak PSS 2,436,964 -> 1,989,479 KiB. The +candidate's host PSS is lower, so no repeat is owed; timing from the memory pair +is non-binding. Four-leg roots are +`/tmp/qwen35-gdn-regstate-{swza,rega,regb,swzb}-7476818c1.*`; memory roots are +`/tmp/qwen35-gdn-regstate-{swz,reg}-mem-7476818c1.*`. `cuobjdump`/NCU downloads +were separately requested but are not authorized and the Ordino service is +unavailable. Preserve this candidate as provisional opt-in until those exact +diagnostics can run. + +## 2026-08-09 — sm_120 GDN decode REGK structurally rejected and removed + +**Disposition:** `REJECTED/REMOVED` before product commit or GPU work. The +red-first missing-contract compile was captured, then the candidate selector, +specialization and tests existed only as uncommitted experiment changes and +were fully restored. Clean disposition head +`28d3a9334d50a80e16f0f3d2ccc0e0cc1b9a4593` has no REGK code or test residue. +The restored portable gate passed **13/13 · 2,005**. + +The contract negatives were live: a parser-prefix mutation killed **0/1** with +two failures, regular-path eligibility died at `-Werror`, and both compile-time +assertions killed the shared-column mapping mutation. The required recurrence +discriminator failed structurally. Same-tool nvcc 12.9 exact +`bf16,bf16,float,NW8,SWIZZLED=true,REGSTATE=true,REGK=true` PTX is byte-identical +for the baseline, first-loop-to-`bk`, and second-loop-to-`bk` variants: exact +symbol **848 instructions**, **53 `ld.shared`**, full PTX SHA-256 +`1fec62b7f48dc5fa2a33bfe414ce6a10aa01979c41489375a8d67b07c88a97a0`; both +loop comparisons are identical. Current REGSTATE without REGK already has +**53 `ld.shared`**, proving the compiler performs the intended K reuse and the +spec's mandatory shared-load reduction cannot be met. + +No GPU, timing, performance or memory run was made and no such claim follows. +REGSTATE's provisional exact result is preserved; exact SASS and NCU remain +pending. The next candidate must be a new spike for direct register-to-global +REGSTATE writeback, removing the final `rr` shared write, synchronization and +shared reread while treating global-store coalescing as a falsifiable tradeoff. +Do not pursue q+k caching: q is single-use and K caching is a compiler no-op. + +## 2026-08-09 — sm_120 GDN decode direct store rejected and removed + +**Disposition:** `REJECTED/REMOVED`. Product experiment `9485a5514` was +reverted in the records/cleanup change; all seven product, test, env and +pending-doc paths equal parent `a9d778957`. Every token SHA-256 is exact and +the candidate graph is gx8/gy800/bx128/smem9728. REGSTATE/DIRECT STORE report +56/55 registers and zero local storage. + +| Arm mean | all fused decode | grid-y 800 | total / output tok/s | TTFT | TPOT / ITL | E2E | +|---|---:|---:|---:|---:|---:|---:| +| REGSTATE | 264.656741 ms | 134.3005475 ms | 6725.73 / 743.715 | 1029.48 ms | 35.02 ms | 5477.21 ms | +| DIRECT STORE | 388.7906645 ms | 196.3294475 ms | 6694.67 / 740.28 | 1031.48 ms | 35.21 ms | 5502.83 ms | + +DIRECT STORE regresses y800 **46.1866%** (164.584 -> 240.600 us/call) and all +fused decode **46.904%**; both raw DST legs lose. Total/output throughput fall +**0.4618%/0.4619%**; TTFT, TPOT/ITL and E2E worsen +**0.1943%/0.5425%/0.4678%**. Evidence roots are +`/tmp/qwen35-gdn-direct-store-{rega,dsta,dstb,regb}-9485a5514.*`. + +Static PTX moved in the intended direction—848 -> 694 instructions, shared +loads 53 -> 48, shared stores 31 -> 15, barriers 2 -> 1—but global stores rose +7 -> 18 and the timing proves lost coalescing dominates. No memory pair is +needed after the hard timing rejection. No acceptance, default, release, +pinned-vLLM, 27B or 35B claim changes. diff --git a/.agents/kernel-matrix.md b/.agents/kernel-matrix.md index 0a0cf92ec..55b226ef3 100644 --- a/.agents/kernel-matrix.md +++ b/.agents/kernel-matrix.md @@ -172,6 +172,17 @@ and the profiled enclosing workload improves **2.272%**; pinned vLLM remains Lifecycle stays `INVENTORIED` because generic Mamba coverage and the 27B/35B release gates are unchanged. [Spec and evidence](specs/sm120-qwen35-conv-chunking-2026-08-07.md). +**2026-08-09 `KERNEL-SSM-MAMBA` sm_120 campaign anchor.** +[#206](https://github.com/mudler/vllm.cpp/issues/206) tracks the RTX 5070 Ti +Qwen3.5-4B Pareto campaign. The current sealed result passes throughput +(**6784.743 vs 6643.593 tok/s**) while TTFT (**1018.040 vs 937.584 ms**), +TPOT/ITL (**34.740 vs 33.906 ms**) and VRAM (**13053.3 vs 12820 MiB**) remain +open. The first ordered gate transplants the historical opt-ins and measures +combined K=4 causal-conv plus 16-token post-conv in a counterbalanced A/B; +same-tool local/vLLM tracing then selects the largest residual. Lifecycle +stays `INVENTORIED`; generic Mamba and 27B/35B coverage are unchanged. +[Campaign spec](specs/sm120-qwen35-pareto-2026-08-09.md). + ## Count invariants - This table has exactly 35 practical kernel-family rows. diff --git a/.agents/roadmap_v1.md b/.agents/roadmap_v1.md index d8ae67cca..3269e645f 100644 --- a/.agents/roadmap_v1.md +++ b/.agents/roadmap_v1.md @@ -43,6 +43,7 @@ issue is not yet placed. Keyed record: update in place, never append. | [#230](https://github.com/mudler/vllm.cpp/issues/230) | — | `test_agent_record.py`: 7 issue-intake tests never run, and error when they do | bug | | [#222](https://github.com/mudler/vllm.cpp/issues/222) | `SERVE-GATE-ONLINE` | The online-serving gate aborts on current main: harness looks for `examples/server`, the build produces `examples/vllm-server` — paths repaired by `2b262622`/`8fce04d3`; residue is the drift guard, the `27n` key and a model gate that recorded a skip as passed | bug | | [#224](https://github.com/mudler/vllm.cpp/issues/224) | `ENG-DOCS-SITE` | Publish `docs/` as a GitHub Pages site without owning a second copy | feature | +| [#206](https://github.com/mudler/vllm.cpp/issues/206) | `KERNEL-SSM-MAMBA` | RTX 5070 Ti: close Qwen3.5-4B TTFT, TPOT, and VRAM gaps vs vLLM | perf | | [#203](https://github.com/mudler/vllm.cpp/issues/203) | `BACKEND-VULKAN` | Vulkan on unified memory holds TWO copies of the weights: 27B peaks at 100.8 GiB RSS and OOM-reboots a Spark | bug | | [#201](https://github.com/mudler/vllm.cpp/issues/201) | `BACKEND-ROCM` | `hipblasGemmEx` overload mismatch in `rocm_matmul_hipblaslt.hip` | bug | | [#199](https://github.com/mudler/vllm.cpp/issues/199) | `BACKEND-METAL-MLX` | macOS MLX build fails on `-Werror` in MLX headers | bug | diff --git a/.agents/specs/sm120-qwen35-conv-block256-2026-08-08.md b/.agents/specs/sm120-qwen35-conv-block256-2026-08-08.md new file mode 100644 index 000000000..84ffca586 --- /dev/null +++ b/.agents/specs/sm120-qwen35-conv-block256-2026-08-08.md @@ -0,0 +1,151 @@ +# sm_120 Qwen3.5 causal-conv 256-thread block — discriminator spike + +**Rows:** `KERNEL-SSM-MAMBA`, feeding `ROAD-V1-C2-LOCAL-BF16`. +**Lifecycle:** `REJECTED` on sm_120 and removed after the committed experiment. +The generic Mamba row remains `INVENTORIED` and the local-4B campaign remains +`GATING` on the accepted arm-1 residual. + +## Current-head revalidation + +The campaign stack was rebased onto `upstream/main` `882ed8fa5`. At immutable +head `23055f49225f9ba4b87781672baad446d31f11eb`, the production Triton-AOT/CUTLASS +binary was rebuilt and the portable prefill contract passed **9/9·88**, the full +CUDA GDN suite **67/67·4631**, and Qwen3.5 paged forward **4/4·8**. + +One graph-node `nsys` run used the same cached Qwen3.5-4B BF16 checkpoint and +closed-loop workload as the existing campaign: 128 ShareGPT requests, 128 output +tokens, concurrency 32, `max_num_batched_tokens=2048`, 1,280 KV blocks and +greedy sampling. It ran under the Ordino `host:ledbx/gpu` lease, `/tmp/gpu` +lock, and a user-systemd 22/25 GiB memory scope. Arm 1 and the post-conv tile +were enabled (`VT_CONV_CHANNEL_TILE=1`, `VT_GDN_POSTCONV_TOKEN_TILE=1`). + +The rebuilt arm-1 `CausalConv1dFwdRegK4Kernel<1>` baseline is: + +| Exact programs | Calls | Total | Mean | Grid / block / registers | +|---:|---:|---:|---:|---:| +| 279 | 576 | 80.655282 ms | 140.026531 us | 64 / 128 / 52 | +| 280 | 336 | 46.974371 ms | 139.804676 us | 64 / 128 / 52 | +| **279+280** | **912** | **127.629653 ms** | **139.944795 us** | **64 / 128 / 52** | +| all causal-conv shapes | 1,728 | 219.357814 ms | 126.943 us | mixed | + +The already-sealed same-tool vLLM trace totals 80.600586 ms over the same 912 +dominant programs (88.378 us mean), leaving the selected local micro-metric +**1.5835x slower** and 47.029067 ms excess. The rebase therefore preserved the +gap and the prior arm-1 result; this is not a stale pre-rebase hypothesis. + +Artifacts: + +- `/tmp/qwen35-conv-block256-baseline-23055f4.nsys-rep`, SHA-256 + `b5341c5bf58e30842432735390f2ddc5513dc17c7172536db8fdc62f4b43fc36`; +- SQLite export `/tmp/qwen35-conv-block256-baseline-23055f4.sqlite`; +- token output `/tmp/qwen35-conv-block256-baseline-23055f4.tokens.json`, + SHA-256 `83fcdc45f79ddb06a634c7d7d95eba3384543b3cd781a45a8db1fc4e2a453545`; +- pinned-vLLM SQLite `/tmp/qwen35-async-3f35356e0-vllm.sqlite`. + +## Isolated hypothesis + +Pinned vLLM uses exact eight-token programs and `BLOCK_N=256` at +`${VLLM_SOURCE}/vllm/model_executor/layers/mamba/ops/causal_conv1d.py:15-79` +and `:689-744`. Its traced launch is grid feature extent 32, block 128 and 32 +registers/thread: Triton maps 256 feature elements across four warps. + +The local K=4 arm 1 at `src/vt/cuda/cuda_gdn.cu:793-960` owns one channel per +thread, so block 128 requires 64 feature blocks. The rejected arm 2 halves that +grid with two channels per thread, but duplicates weights/windows, raises the +trace-visible register count 52→58 and is 4.052% slower than arm 1. + +This discriminator instead gives **one channel to each of 256 threads**. It +also produces 32 feature blocks, but does not duplicate any lane's weights, +window, bias or state pointers. Relative to two arm-1 blocks, it keeps the same +total 256 threads/eight warps and channel-local arithmetic while halving +block-level descriptor and sequence/chunk control. It is deliberately a +single-variable test of block width, not a claim that CUDA's lane mapping is +identical to Triton's vectorized four-warp program. + +The counter-hypothesis is explicit: 52 registers × 256 threads consumes roughly +twice the registers per block, and the larger block may reduce residency enough +to erase the saved control. Nsight Compute is unavailable, so occupancy is not +asserted; the per-shape timing and CUPTI launch metadata decide the experiment. + +## Port and rollback + +Extend the existing same-binary selector with exact arm `3`: + +- `0`: runtime-width, block 128, one channel/thread; +- `1`: K=4, block 128, one channel/thread (current winner/control); +- `2`: K=4, block 128, two channels/thread (retained falsified control); +- `3`: K=4, block 256, one channel/thread (this discriminator). + +Unset and invalid spellings remain arm 0. Non-K4 and non-positive channel shapes +fall back to arm 0. Generalize the K4 kernel/launcher by compile-time block +threads and channels-per-thread so arms 1 and 2 instantiate their existing +128-thread contracts and arm 3 instantiates `<256,1>`. Channel indexing and +feature-block calculation must use the selected compile-time block width; no +metadata, state layout, tensor stride, scheduler, model, post-conv or default +path changes are in scope. + +Every active channel must retain exactly the shipped order: bias; K=4 taps in +ascending order; optional SiLU; output store; window shift; raw-input final-state +writeback. The arm stays opt-in even if locally positive because the repeated +local and hardware-unavailable 27B/35B default gates remain open. + +## Red tests and gates + +Before implementation, extend the portable contract so it fails to compile or +fails assertions until arm 3 exists. It must prove: + +1. only exact string `"3"` selects the new arm; `"30"`, `"3garbage"`, empty, + unset and all other values preserve arm 0; +2. C=8192,K=4 resolves arm 3 to 32 feature blocks and 256 threads; partial + channels round up; unsupported widths resolve to arm 0/64/128; +3. the shared production dispatcher invokes the arm-3 callback exactly once; + deleting the callback branch or mapping it to arm 1 is RED; +4. CUDA arm 3 is byte-identical to arm 1 for output and final state over the + existing BF16/F32, initial/fresh state, unequal exact chunks, `T / 64 / 52 | **219.506 ms** | **140.133 / 139.982 us** | **6767.62 tok/s** | **1013.82 ms** | **34.88 ms** | +| 2 K=4, two channels | K4<2> / 32 / 58 | 228.401 ms | 145.586 / 145.468 us | 6757.19 tok/s | 1017.61 ms | 34.91 ms | + +Arm 1 improves the selected kernel total **6.436%** and every observed +enclosing axis (total/output +0.122%, TTFT -0.282%, TPOT/ITL -0.086%, E2E +-0.123%). Contrary to the initial register hypothesis, specialization raises +the CUPTI register count 43→52; the win is therefore dead runtime-width work, +not higher occupancy. Arm 2 halves feature blocks but raises registers to 58, +is **4.052% slower than arm 1**, and is neutral/slightly negative end to end; +the 256-channel hypothesis is falsified on sm_120. + +All rebuilt-series token files have SHA-256 +`83fcdc45f79ddb06a634c7d7d95eba3384543b3cd781a45a8db1fc4e2a453545`. +Portable tests pass 9/9·88, CUDA GDN 67/67·4631, and Qwen3.5 paged-forward +4/4·8. Arm 1 remains opt-in because its enclosing movement is small and the +repeated local plus hardware-unavailable 27B/35B default gates remain open. + +Final reports: + +- arm 0 `/tmp/qwen35-conv-rebuilt-arm0-565a26fcc.nsys-rep`, SHA-256 + `39d383dd878fc340a3cfaaee79a4addcb4eccb181439e9b4725f724f4569a6eb`; +- arm 1 `/tmp/qwen35-conv-rebuilt-arm1-565a26fcc.nsys-rep`, SHA-256 + `c8799ac0b4cdf997d383fe8a690b223be882dce3b1ee1a6fff35a62d75f7cf85`; +- arm 2 `/tmp/qwen35-conv-rebuilt-arm2-565a26fcc.nsys-rep`, SHA-256 + `3d38793571539864b23688fd9a85966debbf1e7c48fe8a1a2509438a45ee0452`. diff --git a/.agents/specs/sm120-qwen35-gdn-decode-bf16-vector-writeback-2026-08-09.md b/.agents/specs/sm120-qwen35-gdn-decode-bf16-vector-writeback-2026-08-09.md new file mode 100644 index 000000000..7e6b643dd --- /dev/null +++ b/.agents/specs/sm120-qwen35-gdn-decode-bf16-vector-writeback-2026-08-09.md @@ -0,0 +1,132 @@ +# sm_120 Qwen3.5 GDN decode BF16 vector-writeback discriminator + +**Lifecycle:** `SPIKED`; implementation and measurement pending + +**Owner rows:** `KERNEL-SSM-MAMBA`, `ROAD-V1-C2-LOCAL-BF16` + +**Hardware/workload:** RTX 5070 Ti (`sm_120`), Qwen3.5-4B plain BF16 + +**Scope/authority:** spec only. No product, tests, GPU work, remote operation, +default, release, acceptance, or 27B/35B claim is part of this checkpoint. + +## Ground and falsifiable hypothesis + +Exact provisional REGSTATE `7476818c1` retains `rr[16]` per lane but writes it +to swizzled shared, synchronizes, and rereads it for coalesced global state +writeback (`src/vt/cuda/cuda_gdn.cu:2789-2865`). Its y800 mean is 134.291105 ms +in the accepted comparison and exact tokens hash to +`83fcdc45f79ddb06a634c7d7d95eba3384543b3cd781a45a8db1fc4e2a453545`. + +The removed scalar-direct experiment `9485a5514` proved the final shared phase +is removable but the replacement transport was wrong. PTX fell 848->694, +shared loads 53->48, shared stores 31->15 and barriers 2->1, yet global stores +rose 7->18; y800/all-fused regressed 46.1866%/46.904%. Product was removed at +`d357be92b`. Do not resurrect that scalar loop. + +For exact Qwen production state `TState=bf16`, lane `wk` already owns logical +columns `wk*16..wk*16+15`: a contiguous 32-byte range. The state allocation and +128-element BF16 row stride are at least 16-byte aligned, and each lane start +adds a multiple of 32 bytes. The existing scalar `Store(__nv_bfloat16*,...)` +uses `__float2bfloat16`, round-to-nearest-even (`cuda_gdn.cu:209-219`). + +**Hypothesis:** convert `rr[0..15]` elementwise with that exact operation, pack +each group of eight into a 16-byte value, and issue two aligned vector stores +per lane directly to its contiguous logical columns. This can remove the final +shared write/barrier/reread without the scalar-direct arm's sixteen store +instructions per lane. If nvcc scalarizes or spills the pack, alignment is not +proved, state bytes differ, or either timing leg loses, reject and remove it. + +Pinned vLLM remains grounded by its FLA register tile and direct persistent +store at `${VLLM_SOURCE}/vllm/third_party/flash_linear_attention/ops/fused_sigmoid_gating.py:156-170`. +That establishes register-to-state lifetime, not this transport schedule. + +## Exact candidate contract + +Add strict default-OFF `VT_GDN_DECODE_BF16_VECSTORE=1`; only exact `"1"` +selects it. Eligibility is the complete regular REGSTATE production path: +BV16, shared swizzle, Dv=Dk=128, NW8, non-speculative fused decode, and +`TState=__nv_bfloat16`. F16/F32 state, partial shapes, alternate BV/NW/layout, +sequential fallback, invalid/null index paths and every unset/invalid spelling +remain incumbent REGSTATE. + +Use a host/device constexpr mapping/alignment contract for lane, pack and +element: + +```text +logical_column(wk, pack, element) = wk*16 + pack*8 + element +pack in {0,1}, element in [0,8) +byte offset from state value-row = logical_column * 2 +``` + +The 128 columns must be a bijection; each pack begins at a 16-byte boundary, +the two packs cover exactly the lane's current `rr[16]`, and four 8-lane +subgroups in a warp address four independent value rows. Form each packed value +with the same per-element BF16 RN conversion as scalar `Store`; use a bit-safe +aligned aggregate (`int4` or equivalent), never type-pun through an unaligned +object. Static assertions must bind element size, pack size, Dk and alignment. + +Only replace REGSTATE's final `rr` shared write, final barrier, and shared +reread/writeback. Preserve recurrence expression/order, shuffles, output store, +initial coalesced state load/barrier, q/k staging, state row/index/null +resolution, conversions, gx8/bx128/smem9728 and every fallback. + +An 8-lane register transpose is a **fallback discriminator**, not the first +implementation. If the vector arm fails only because compiler output cannot +issue aligned vector stores, write a new spike for two 8x8 half-transposes that +make each scalar store instruction address adjacent subgroup columns. Do not +combine transpose and vector packing in this row: the extra shuffles obscure +which transport won. + +## Red-first tests and mutations + +Before product work, add portable tests for strict parsing, every eligibility +predicate, exactly one incumbent/vector callback, the complete 128-column +mapping bijection, two disjoint aligned packs per lane, and BF16 pack bytes +matching sixteen independent reference RN conversions. Capture the intended +red result before adding selector/product support. + +Extend CUDA REGSTATE-versus-vector coverage to require byte-exact output and +the complete persistent state for compact/indexed storage, valid/null indices, +production BF16 input/output/state and every ineligible dtype/shape fallback. +The graph test must distinguish the vector callback and still observe exactly +one gx8/gy800/bx128/smem9728 node. Cached Qwen tokens must match the hash above. + +Scratch mutations must kill parser, each eligibility predicate, lane/pack +addressing, either pack store, BF16 conversion, indexed row, null handling and +callback identity. Replacing vector writeback with incumbent shared writeback +must fail compiler discrimination. Restore the tree byte-for-byte after each +mutation. + +## Compiler, correctness, and performance gates + +Compare exact REGSTATE and vector specializations from one immutable nvcc +binary. Record compiler/binary hashes, PTX/cubin/SASS symbol identity, +registers, local/stack/spills, barriers, total instructions, shared/global +load/store instructions, and actual-function occupancy. The vector arm must: + +- contain two 16-byte vector global state stores per lane (or an equivalent + lower instruction form), with no scalarized sixteen-store sequence; +- materially reduce global-store instructions versus removed scalar-direct + (18) and reduce shared loads/stores plus barriers versus REGSTATE; +- have zero local/stack/spills, unchanged geometry/shared bytes, no occupancy + loss over 12.5%, and no total SASS growth over 5%. + +Missing symbols, unproved alignment, scalarization, compiler no-op, or any +resource miss rejects/removes the candidate before performance credit. When +authorized, same-tool NCU on matched y800 nodes records global sectors and +transactions, shared traffic, barrier/scoreboard stalls, DRAM and occupancy for +REGSTATE/vector/pinned-vLLM; counters diagnose but never substitute for timing. + +After exact portable, CUDA, graph and cached-Qwen correctness, hold one GPU lock +and run same-binary `REGa -> VECa -> VECb -> REGb`. Both VEC raw y800 legs must +beat both controls; counterbalanced y800 must improve at least 1.00%; all fused +decode must improve; total/output throughput, TTFT, TPOT/ITL, E2E, GPU memory +and material host PSS must not regress. Reproduce a passing result with the +same binary. Any miss removes selector/product/tests and records the rejection. +A pass remains explicit opt-in: it grants no default, release, pinned-vLLM or +unavailable 27B/35B acceptance. + +## First implementation step + +Add only portable parser/eligibility/mapping/pack-reference tests and capture +their red failure. Do not edit CUDA recurrence code until that evidence exists. diff --git a/.agents/specs/sm120-qwen35-gdn-decode-bv-sweep-2026-08-09.md b/.agents/specs/sm120-qwen35-gdn-decode-bv-sweep-2026-08-09.md new file mode 100644 index 000000000..61f429130 --- /dev/null +++ b/.agents/specs/sm120-qwen35-gdn-decode-bv-sweep-2026-08-09.md @@ -0,0 +1,186 @@ +# sm_120 Qwen3.5 fused GDN decode value-tile sweep + +**Rows:** `KERNEL-SSM-MAMBA`, feeding `ROAD-V1-C2-LOCAL-BF16`. +**Lifecycle:** `COMPLETE` local discriminator. BV8 and BV24 were byte-exact but +performance-rejected and removed; BV16 remains the one opt-in incumbent and +BV32 remains shipped default. The generic Mamba row remains `INVENTORIED`; +default and Qwen3.6-27B/35B release gates stay open. + +## Intake, evidence and fixed scope + +This records-only spike is based on current head `26122c7e5debaf152b29f589fc778e532f6e5eb9` +and `upstream/main` `ba8d867c9ea1aab0896af9c741924ba9855e0400`. +The accepted BV16 result at product commit `92256e6a9` keeps NW8 and every +per-row operation byte-ordered while changing the production launch from +grid-x 4 / block 256 to grid-x 8 / block 128. On the dominant grid-y 800 shape +it improves **183.526→163.432 us/call (-10.95%)**, but remains **1.276x** behind +the pinned same-tool vLLM result, **128.061 us/call**. BV16 is the incumbent; +shipped BV32 is the default control. The correctness-invalid NW4 experiment +changed the Dk partition/reduction order and is forbidden here. + +The local execution chain is `GdnDecodeFusedCuda` → +`DispatchGdnDecodeValueTile` → `LaunchGdnDecodeFused{S,NW}` → +`GdnDecodeFusedKernel` at `src/vt/cuda/cuda_gdn.cu:2700-2931`, with the +CUDA-free parser, ceil geometry and callback dispatcher at +`src/vt/cuda/gdn_decode_fused.h:11-67`. The same-process BV16/BV32 byte matrix +is at `tests/vt/test_ops_gdn.cpp:1865-1951,3464-3481`; its portable contract is +`tests/vt/test_gdn_decode_fused.cpp:14-85`. + +Pinned vLLM `5559679229bc961848b121ccdeaa8fa5d79bec98` vendors FLA at +`${VLLM_SOURCE}/vllm/third_party/flash_linear_attention/ops/fused_sigmoid_gating.py:24-277`: +the recurrence is `:122-170`, `BV=min(next_power_of_2(V),32)` is `:205-212`, +grid `(NK,NV,N*HV)` is `:241`, and `num_warps=4` is `:212,275`. This sweep is +a local schedule discriminator, not a claim that BV8/BV16/BV24 mirrors FLA's +BV32 choice. No CUTLASS, cuBLASLt or DeepGEMM layer executes this recurrence. + +## Exact design and launch contracts + +Change only independent value-row tiling. Keep `NW=8`, each row's eight +consecutive Dk-slice lanes, `c0/c1`, shuffle offsets, decay, dot/update/output +arithmetic, tail guards, output store and state writeback unchanged. Do not +specialize, duplicate or reorder the recurrence. For Dv=Dk=128: + +| Candidate | grid-x `ceil(128/BV)` | block `BV*NW` | dynamic shared bytes `(2*Dk + BV*(Dk+1))*4` | +|---|---:|---:|---:| +| BV8 | 16 | 64 | 5,152 | +| BV16 incumbent | 8 | 128 | 9,280 | +| BV24 | 6 | 192 | 13,408 | +| BV32 shipped | 4 | 256 | 17,536 | + +The BV24 sixth block is a guarded eight-row tail; grid rounding must be ceiling, +never floor. General geometry remains `BV=min(Dv,requested)`, +`NV=ceil(Dv/BV)`, `NW_eff=(Dv>=32 ? requested_nw : 1)`, block +`BV*NW_eff`, and the shared-byte formula above. Zero Dv/Dk is the existing +no-launch contract and the existing >48 KiB sequential fallback is unchanged. + +Extend `VT_GDN_DECODE_BV` with strict full-string parsing. Exactly `"8"`, +`"16"`, `"24"` and `"32"` resolve BV8/BV16/BV24/BV32 respectively. Unset, +empty, whitespace, signed or zero-padded forms, prefixes/suffixes, overflow and +every other value resolve shipped BV32. No `atoi` or prefix parsing. Production +must use the same four-callback portable dispatcher that tests exercise. + +## RED-first correctness and mutation contract + +Before product edits, extend the portable test so it fails on the current +two-arm contract. It must prove exact parser/fallback semantics; all four +production geometries/shared sizes; BV24 ceiling/tail and partial dimensions; +the Dv<32 NW1 corner; zero-dimension no-launch; and exactly one callback for +each candidate. Deleting or rerouting either new arm, relaxing parsing, changing +ceil to floor, or corrupting a block/shared calculation must make it RED. + +Expand the CUDA matrix to run BV8, BV16, BV24 and BV32 in one process and +require raw-byte equality of output and final state against BV32. Cover compact +and indexed state, BF16 and F32 state, production Dv=Dk=128, a partial Dv, and +an indexed null row whose output is zero and unreferenced state stays untouched. +Mutate each new dispatch, the BV24 last-tile guard, row mapping, and both +output/state stores; each mutation must fail focused coverage and restoration +must return the tree byte-for-byte. + +Build and run `test_gdn_decode_fused`, full CUDA `test_ops_gdn`, +`test_qwen35_paged_forward`, and `vllm-bench` from one CUDA/Triton/CUTLASS +build. The production 128-request greedy token output for every arm must equal +SHA-256 `83fcdc45f79ddb06a634c7d7d95eba3384543b3cd781a45a8db1fc4e2a453545`. +Any byte mismatch stops performance work. Run focused, staged and full gates, +then fresh static and targeted-mutation review; the operator reruns the declared +gate independently. + +## Same-binary trace and selection gate + +Under one GPU lease and `flock /tmp/gpu`, use one rebuilt binary and the cached +Qwen3.5-4B BF16 workload: 128 ShareGPT requests, 128 output tokens, c32, +`max_num_batched_tokens=2048`, 1,280 KV blocks and greedy sampling. Keep every +other selector identical. Run one mirrored eight-leg series: + +`BV32a → BV16a → BV8a → BV24a → BV24b → BV8b → BV16b → BV32b`. + +Collect `nsys --cuda-graph-trace=node`, exported SQLite, raw summaries, token +hashes, binary/build SHA, clocks/contention, peak VRAM and peak PSS for every +leg. Launch evidence must show all four exact grid/block/shared contracts above, +NW8, the fused kernel (not scan fallback), identical candidate call counts, and +the dominant grid-y 800 group separately from all fused-decode calls. + +A new candidate is eligible only when its counterbalanced mean grid-y 800 time +is at least **1.00% lower than BV16**, its all-fused-decode mean does not regress +against BV16, both samples have the same direction, and every token hash is +identical. If both qualify, select the lower grid-y 800 mean; a noisy/mixed +ordering retains BV16. The selected opt-in then must not regress against BV16 +on any enclosing mean: total and output throughput may not decrease; TTFT, +TPOT, ITL, E2E, peak VRAM and peak PSS may not increase. Record values and +ratios for every axis. Default BV32 receives no displacement credit. + +At disposition, retain exactly shipped BV32 plus one opt-in incumbent: BV16 if +neither new candidate clears every bar, otherwise the selected BV8 or BV24. +Remove every losing dormant selector, enum/dispatcher arm, instantiation and +test expectation in the same implementation change; the removed selector +strings must then fall back to BV32. Preserve rejected measurements in the +append-only benchmark record and measured result section here. Even a local win +stays opt-in: default/release gates remain open and the 27B/35B vehicles are +hardware-unavailable on this host. + +## Measured result — BV8/BV24 rejected, BV16 retained opt-in + +Sweep implementation `e102a14de` passed operator runtime verification: portable +selector/geometry **3/3 · 100**, CUDA GDN **69/69 · 4,850** (including public +production-graph geometry) and Qwen3.5 paged-forward **4/4 · 8**. The standard +128-request greedy workload ran in the prescribed order +BV32a→BV16a→BV8a→BV24a→BV24b→BV8b→BV16b→BV32b. All eight token files have +SHA-256 `83fcdc45f79ddb06a634c7d7d95eba3384543b3cd781a45a8db1fc4e2a453545`. + +| Arm | all fused decode, 1,704 calls | grid-y 800, 816 calls | total / output tok/s | TTFT | TPOT / ITL | E2E | +|---|---:|---:|---:|---:|---:|---:| +| BV32a | 302.606110 ms | 153.728820 ms | 6733.45 / 744.57 | 1021.35 ms | 35.04 ms | 5471.00 ms | +| BV16a | 271.472694 ms | 136.909597 ms | 6743.80 / 745.71 | 1020.07 ms | 34.98 ms | 5462.59 ms | +| BV8a | 277.729795 ms | 140.484486 ms | 6741.12 / 745.41 | 1021.16 ms | 34.99 ms | 5464.71 ms | +| BV24a | 279.819526 ms | 141.534196 ms | 6734.28 / 744.66 | 1022.07 ms | 35.03 ms | 5470.34 ms | +| BV24b | 279.639897 ms | 141.524051 ms | 6731.40 / 744.34 | 1022.61 ms | 35.04 ms | 5472.57 ms | +| BV8b | 279.157494 ms | 141.735960 ms | 6738.32 / 745.10 | 1021.23 ms | 35.01 ms | 5466.97 ms | +| BV16b | 271.707232 ms | 137.034656 ms | 6736.22 / 744.87 | 1022.23 ms | 35.01 ms | 5468.79 ms | +| BV32b | 302.982021 ms | 153.870033 ms | 6727.95 / 743.96 | 1023.06 ms | 35.06 ms | 5475.47 ms | + +The counterbalanced means are: + +| Arm | all fused decode | vs BV16 | grid-y 800 | vs BV16 | +|---|---:|---:|---:|---:| +| BV32 | 302.7940655 ms | +11.4894% | 153.7994265 ms | +12.2852% | +| BV16 | 271.5899630 ms | control | 136.9721265 ms | control | +| BV8 | 278.4436445 ms | **+2.5235%** | 141.1102230 ms | **+3.0211%** | +| BV24 | 279.7297115 ms | **+2.9971%** | 141.5291235 ms | **+3.3270%** | + +Both samples of each new candidate are slower than their corresponding BV16 +sample. BV8 and BV24 therefore fail both the required ≥1% dominant-shape win +and the all-fused non-regression bar; neither reaches downstream selection. +BV16 remains a repeatable win over BV32 in this series: **-10.3054%** across all +fused calls and **-10.9411%** on grid-y 800. + +| Arm mean | total tok/s | output tok/s | TTFT | TPOT / ITL | E2E | +|---|---:|---:|---:|---:|---:| +| BV32 | 6730.700 | 744.265 | 1022.205 ms | 35.050 ms | 5473.235 ms | +| BV16 | 6740.010 | 745.290 | 1021.150 ms | 34.995 ms | 5465.690 ms | +| BV8 | 6739.720 | 745.255 | 1021.195 ms | 35.000 ms | 5465.840 ms | +| BV24 | 6732.840 | 744.500 | 1022.340 ms | 35.035 ms | 5471.455 ms | + +The profiler confirms exact Dv=Dk=128 launches at 52 registers/thread: BV8 +grid-x 16 / block 64 / 5,152 shared bytes; BV16 8 / 128 / 9,280; BV24 6 / +192 / 13,408; and BV32 4 / 256 / 17,536. Disposition cleanup commit +`634ccba70` removes the rejected BV8/BV24 selectors, dispatch arms, +instantiations and expectations, so exact strings `"8"` and `"24"` again fall +back to BV32. Exact `"16"` remains the sole opt-in and unset remains BV32. + +The current series' dominant BV16 mean is **136.9721265 ms / 816 = +167.858 us/call** versus pinned vLLM **128.061 us/call**, about **1.311x** +slower. The prior accepted BV16 series reported **1.276x**; that separate run is +preserved rather than overwritten, and the difference is treated as run +variance rather than a new oracle claim. The residual, default gate and +hardware-unavailable 27B/35B release gates remain open. + +Evidence is preserved as +`/tmp/qwen35-gdn-bvsweep-{32a,16a,8a,24a,24b,8b,16b,32b}-e102a14de.{nsys-rep,sqlite,log,tokens.json}`. +Trace SHA-256 values in that order are +`323bf5a4c28b46bc2b4bdd1ee7469abb0ad3f2205d6866c11c8d07fff17942e8`, +`21a48aa428eebf5c85ded20eb309a7f8ca5272a3e6c8d34aa5ca18e66f336aef`, +`8266020bd8dae1a1580484e9e8bf6ce3ac427598d588393704793d47b7e1e260`, +`bd36b8470dd3f563cc494e9e9e5215b846cd38ec0e2f34d104406586ff0c3229`, +`ed046407514480c281e6de2511c27a566c7da7504970159b0b0f574605f003f1`, +`f7c9931f3953a836f612d8d940ef91b8657b3f90940e66a75e95cc4f744ee374`, +`e22aa55a1094ef9282a28a107d26a3297409ff5f74d24ce6f138380ce2a21db8` +and `7a6cbb1281acf9fd23c9923156d744d76d54af1765b7752345e18dc4c50cfcce`. diff --git a/.agents/specs/sm120-qwen35-gdn-decode-bv16-2026-08-08.md b/.agents/specs/sm120-qwen35-gdn-decode-bv16-2026-08-08.md new file mode 100644 index 000000000..a9df9489c --- /dev/null +++ b/.agents/specs/sm120-qwen35-gdn-decode-bv16-2026-08-08.md @@ -0,0 +1,192 @@ +# sm_120 Qwen3.5 fused GDN decode BV16 — discriminator spike + +**Rows:** `KERNEL-SSM-MAMBA`, feeding `ROAD-V1-C2-LOCAL-BF16`. +**Lifecycle:** `READY` as an opt-in local discriminator. The generic Mamba row +remains `INVENTORIED`; Qwen3.6-27B/35B release gates are unavailable here. + +## Current-head and intake revalidation + +This spike is based on immutable campaign head +`96382c7f5b625f5e5cdf66feebd49d76d334efae`, itself based on +`upstream/main` `ba8d867c9ea1aab0896af9c741924ba9855e0400`. The active local +stack is already bound to PR #155; this records the next discriminator without +changing product code, tests, claims or remotes. A current-tree search finds no +`VT_GDN_DECODE_BV` selector or BV16 implementation. + +The local graph-node trace +`/tmp/qwen35-conv-block256-arm1-84c7e23b0.sqlite` resolves the executing kernel +to `GdnDecodeFusedKernel<__nv_bfloat16,__nv_bfloat16,float,8>`. Across all +1,704 calls it takes **295.924017 ms** (**173.664 us** mean), with grid-y 800 +dominant at 816 calls / **150.043436 ms** / **183.877 us**. Its launch is +grid `(4,N*Hv,1)`, block 256, 52 registers/thread and 17,536 bytes dynamic +shared memory. + +The same-tool pinned-vLLM trace `/tmp/qwen35-async-3f35356e0-vllm.sqlite` +resolves `fused_sigmoid_gating_delta_rule_update_kernel`: all 1,873 calls take +**229.354660 ms** (**122.453 us** mean), and extent 800 takes 840 calls / +**107.571496 ms** / **128.061 us**. It launches grid `(1,4,N*Hv)`, block 128, +96 registers/thread and 512 bytes dynamic shared memory. The unequal total call +counts prohibit treating the all-call totals as a throughput ratio; the shape +and matched workload establish a concrete schedule difference to test. + +## Whole execution chain + +The local direct-CUDA path is `GdnDecodeFusedCuda` through +`LaunchGdnDecodeFused{S,NW}` to `GdnDecodeFusedKernel` at +`src/vt/cuda/cuda_gdn.cu:2699-2925`. It caps `BV` at 32, computes +`NV=ceil(Dv/BV)`, and launches `BV*NW` threads. For production `Dv=Dk=128` and +the accepted `NW=8`, that is BV32, grid-x 4 and block 256. Each value row owns +eight consecutive lanes; those lanes retain the exact Dk slices and +`__shfl_xor` reduction tree. + +Pinned vLLM vendors the Flash Linear Attention Triton operation at +`${VLLM_SOURCE}/vllm/third_party/flash_linear_attention/ops/fused_sigmoid_gating.py:24-277`: +`BV=min(next_power_of_2(V),32)` at `:208`, grid `(NK,NV,N*HV)` at `:241`, and +`num_warps=4` at `:212,275`. The profiler confirms that generated Triton kernel +executed. This operation does not route through CUTLASS, cuBLASLt or DeepGEMM; +the relevant dependency boundary is the vendored FLA Triton body and its +resolved profiler launch. + +The existing `VT_GDN_DECODE_NW=4` probe is not this design. It changed each +row's Dk partition and reduction tree; it was both slower locally and +correctness-invalid: 24/128 streams and 678/16,384 aligned token positions +differed, with SHA-256 `93003ca1...8eebe` instead of the accepted +`83fcdc45...453545`. NW remains 8. + +## Isolated hypothesis and design + +Keep `NW=8` and every per-row lane, slice, shuffle, arithmetic and store order +unchanged. Change only the independent value tile from BV32 to BV16: + +- production grid-x doubles 4→8 while block width halves 256→128; +- one block still has exactly eight lanes for each value row; +- each row's `c0/c1`, decay, dot, update, output reduction and state writeback + remain byte-for-byte ordered as BV32; +- the shared state slice halves from `32*(Dk+1)` to `16*(Dk+1)` floats, giving + 9,280 bytes at Dk=128 rather than 17,536 bytes; +- value rows are independent, so only value-tile scheduling changes. + +The counter-hypothesis is that twice as many blocks repeat q/k staging and +block setup, erasing any occupancy/latency benefit. Measurement decides; no +occupancy claim is inferred from block width alone. + +## Selector, geometry and rollback + +Add exact opt-in `VT_GDN_DECODE_BV=16`. Only the exact string `"16"` selects +BV16. Unset, empty, whitespace, prefixes/suffixes, numeric lookalikes and every +other value select the shipped BV32 default. This is intentionally stricter +than `atoi` parsing and leaves existing deployments unchanged. + +Factor a CUDA-free selector/geometry contract and shared callback dispatcher +for portable tests and the production launcher. For requested width `R`, use +`BV=min(Dv,R)`, `NV=ceil(Dv/BV)`, the existing `NW` corner rule, block threads +`BV*NW_eff`, and shared bytes `(2*Dk + BV*(Dk+1))*sizeof(float)`. Zero dimensions +remain the existing no-launch path. Partial `Dv` uses the existing tail guards; +`Dv<32` retains the existing `NW_eff=1`, so no row group crosses a warp. + +The implementation may parameterize or pass BV at runtime, but it must not +duplicate or reorder the kernel recurrence. `VT_GDN_DECODE_BV` is removed with +its code/tests if the discriminator fails. It stays opt-in even on a local win +until repeated local evidence and the unavailable 27B/35B gates close. + +## RED-first tests and mutations + +Before product implementation, add a CPU-runnable selector/geometry test that +fails until the new contract exists. It must prove: + +1. exact `"16"` selects BV16, while unset, `""`, `"16x"`, `"016"`, `" 16"`, + `"32"` and invalid values select BV32; +2. Dv/Dk=128 with NW8 resolves BV32 to grid-x 4, block 256, 17,536 bytes and + BV16 to grid-x 8, block 128, 9,280 bytes; +3. partial dimensions round up without dropping rows, and `Dv<32` preserves + the current one-lane-per-row corner contract; +4. the shared dispatcher invokes exactly one BV16 or BV32 callback, so deleting + the selector arm or mapping `"16"` to the default is RED. + +Extend CUDA GDN coverage with BV16-versus-BV32 byte comparisons for both output +and final state. Cover compact and indexed state, BF16 and F32 state, a partial +value dimension and production Dv=Dk=128; include a null indexed row. The matrix +must exercise both selector arms in one process. Mutations to BV selection, +grid rounding, block width, row mapping, tail guards or either state/output +store must make a focused test fail, followed by byte-for-byte restoration. + +## Build and correctness gates + +Build the portable selector test, `test_ops_gdn`, `test_qwen35_paged_forward` +and the production `vllm-bench` from the CUDA/Triton/CUTLASS build. Run the +portable test first, then the full CUDA GDN suite and Qwen3.5 paged-forward +gate under the operator's GPU lease, `/tmp/gpu` lock and 22/25 GiB user-systemd +scope. Run staged and full preflight from `nix develop .#cuda`. + +The production 128-request greedy token file must remain byte-identical to +SHA-256 `83fcdc45f79ddb06a634c7d7d95eba3384543b3cd781a45a8db1fc4e2a453545`. +Any mismatch stops performance work immediately. + +## Same-binary performance acceptance + +Use one rebuilt binary, one GPU lease/lock and the existing identical cached +Qwen3.5-4B BF16 workload: 128 ShareGPT requests, 128 output tokens, concurrency +32, `max_num_batched_tokens=2048`, 1,280 KV blocks and greedy sampling. Keep the +accepted causal-conv and post-conv selectors identical in both arms. Run the +counterbalanced BV32→BV16→BV16→BV32 order and record graph-node traces, raw +summaries, token hashes, build SHA, clocks/contention and peak VRAM. + +BV16 is retained only if the grid-y 800 GdnDecode total improves by **at least +1%** against BV32 and the all-GdnDecode total does not regress. It then owes an +enclosing repeated A/B in which total/output throughput do not decrease and +TTFT, TPOT/ITL, E2E and peak VRAM do not increase. Record both values and ratios +for every axis; a noisy or mixed result fails rather than becoming a speed claim. + +If either micro bar or any enclosing axis fails, remove selector, product and +test code and preserve the refutation in the append-only benchmark record and +this spec. If all local bars pass, retain the arm opt-in, complete fresh +implementer/reviewer mutation review, and keep release/default status open on +the hardware-unavailable 27B/35B gates. + +## Measured result — ACCEPTED OPT-IN + +Product commit `92256e6a9` passes the local acceptance gate. Operator runtime +verification is green: portable selector/geometry **3/3 · 49**, CUDA GDN +**68/68 · 4,699**, and Qwen3.5 paged-forward **4/4 · 8**. Every production +token file in the four-run profile and the memory pair has SHA-256 +`83fcdc45f79ddb06a634c7d7d95eba3384543b3cd781a45a8db1fc4e2a453545`. + +The counterbalanced same-binary series ran BV32a → BV16a → BV16b → BV32b: + +| Arm | all GDN decode, 1,704 calls | grid-y 800, 816 calls | total / output tok/s | TTFT | TPOT / ITL | E2E | +|---|---:|---:|---:|---:|---:|---:| +| BV32a | 295.608238 ms | 149.887342 ms | 6763.51 / 747.89 | 1014.64 ms | 34.90 ms | 5446.55 ms | +| BV16a | 265.417816 ms | 133.690094 ms | 6766.91 / 748.27 | 1013.00 ms | 34.89 ms | 5443.79 ms | +| BV16b | 264.152452 ms | 133.030633 ms | 6779.94 / 749.71 | 1011.37 ms | 34.82 ms | 5433.25 ms | +| BV32b | 295.123586 ms | 149.627524 ms | 6763.94 / 747.94 | 1015.48 ms | 34.89 ms | 5446.15 ms | + +The counterbalanced means improve all decode calls +**295.365912→264.785134 ms (-10.3535%)** and the matched grid-y 800 shape +**149.757433→133.360364 ms (-10.9491%)**. The intended launch executed: +grid-x 4 / block 256 / 52 registers became grid-x 8 / block 128 / 52 +registers. Local BV16 is still about **1.276x slower** than pinned vLLM on the +matched dominant shape (163.43 versus 128.061 us/call), so that gap stays open. + +Including the separately sampled memory pair, the three-pair means move in the +accepted direction on every enclosing axis: total throughput +**6783.437→6791.063 tok/s (+0.1124%)**, output throughput +**750.093→750.940 tok/s (+0.1129%)**, TTFT +**1013.237→1011.490 ms (-0.1724%)**, TPOT/ITL +**34.7867→34.7500 ms (-0.1054%)**, and E2E +**5430.777→5424.630 ms (-0.1132%)**. The memory pair is +BV32→BV16 **13058→13054 MiB** peak GPU memory and +**2,273,490→2,170,880 KiB** peak PSS, both non-regressing. + +Evidence lives at +`/tmp/qwen35-gdn-bv{32a,16a,16b,32b}-92256e6a9.{nsys-rep,sqlite,tokens.json}`; +trace SHA-256 values in that order are +`198d1bf843443f69fd1491131a13e7aa95cc136c56e96441aa43c48f267958f5`, +`2f5e47f54faafa70ed8494c3afaa2160e31d311d90e7412cbf9777f2e93e3d18`, +`3b0303a97e2335d5ca2f88eacf5c2e5f063d7d899cb44db3dd7d6a170938d117`, +and `3d7b4a214b14976c95e9e4c944540e91163c7f10e820bb63bf7d0cb9fc9d8b11`. +Memory evidence is +`/tmp/qwen35-gdn-bv{32,16}-mem-92256e6a9.{jsonl,log,tokens.json}`. + +Disposition is therefore **ACCEPTED OPT-IN**. `VT_GDN_DECODE_BV=16` remains an +explicit rollback-safe experiment; unset retains BV32. No default or release +credit is claimed while the Qwen3.6-27B/35B gates are hardware-unavailable. diff --git a/.agents/specs/sm120-qwen35-gdn-decode-direct-store-2026-08-09.md b/.agents/specs/sm120-qwen35-gdn-decode-direct-store-2026-08-09.md new file mode 100644 index 000000000..21752d51c --- /dev/null +++ b/.agents/specs/sm120-qwen35-gdn-decode-direct-store-2026-08-09.md @@ -0,0 +1,185 @@ +# sm_120 Qwen3.5 fused GDN decode direct-state-store discriminator + +**Lifecycle:** `REJECTED/REMOVED`; decisive timing loss, product/tests restored +to parent `a9d7789579582872d5063620acd9eb33b73ed1a3` + +**Owner rows:** `KERNEL-SSM-MAMBA`, `ROAD-V1-C2-LOCAL-BF16` + +**Hardware/workload:** RTX 5070 Ti (`sm_120`), Qwen3.5-4B plain BF16 + +**Authority:** records-only rejection/cleanup. No retained product or test code, +new GPU run, remote operation, default, release, acceptance, or 27B/35B claim. + +## Gap, evidence, and falsifiable hypothesis + +This spike is based on clean disposition head +`63842df4b85a5eaf3b703ff7d3627d04340a881d`. REGK was structurally rejected +and removed: same-tool nvcc 12.9 emitted byte-identical exact PTX for its +baseline and either recurrence loop restored to `bk`, proving the compiler +already keeps K live. Do not resurrect REGK or add q+k caching; q is single-use. + +The incumbent is exact, provisional `REGSTATE` product `7476818c1`. Its y800 +mean is **164.57 us/call**, **-1.093499%** versus SWZ but about **1.285x** +pinned vLLM's 128.061 us/call. Its exact compiler structure is 894 PTX +instructions, 53 `ld.shared`, 31 `st.shared`, two barriers, 56 registers and +zero stack/spills; gx8/bx128/smem9728 is unchanged from SWZ. Exact token +SHA-256 is +`83fcdc45f79ddb06a634c7d7d95eba3384543b3cd781a45a8db1fc4e2a453545`. + +In `src/vt/cuda/cuda_gdn.cu:GdnDecodeFusedKernel`, REGSTATE finishes the +recurrence with each lane's 16 updated F32 values in `rr`. It then writes +`rr[j]` back to swizzled shared, executes a second `__syncthreads()`, and +cooperatively rereads shared before converting/storing persistent state. The +sequence exists only to recover coalesced global stores. + +Pinned vLLM `5559679229bc961848b121ccdeaa8fa5d79bec98` directly stores its F32 +register state tile: vendored FLA +`vllm/third_party/flash_linear_attention/ops/fused_sigmoid_gating.py:156-170` +forms `p_ht` from the resolved state index and executes +`tl.store(p_ht, b_h.to(p_ht.dtype.element_ty), mask=mask_h)`. This grounds the +direct register-to-persistent-state lifetime, not FLA's layout or schedule. + +**Hypothesis:** on the exact full REGSTATE production tile, each thread can +store its own `rr[16]` directly to logical persistent-state columns, removing +the final shared writes, barrier and rereads while preserving every state byte. +The deliberate counter-cost is less-coalesced global traffic. Static structure +and same-binary timing decide; inferred barrier or coalescing credit is invalid. + +## Exact selector, eligibility, and mapping + +Add strict default-OFF `VT_GDN_DECODE_DIRECT_STORE=1`. Only complete string +`"1"` selects it. Unset, empty, whitespace, prefixes/suffixes, numeric and +truthy lookalikes preserve the current path. + +Eligibility is only the regular, non-speculative, non-fallback, full-tile +`REGSTATE=true` production specialization: exact BV16+SWZ, `Dv==Dk==128`, +resolved NW8. The flag never promotes another path. Invalid/unset or any failed +predicate retains current REGSTATE when otherwise eligible; partial shapes, +BV32, unswizzled BV16, other NW, sequential fallback, speculative decode and +null/index corner behavior remain incumbent. + +Add a host/device constexpr mapping helper used by both product and portable +tests. For `vi=0..15`, `wk=0..7`, `j=0..15`, the direct offset from `s_head` is: + +```cpp +vi * dk + GdnDecodeRegisterLogicalColumn(wk, j) +// exact production form: vi * 128 + wk * 16 + j +``` + +The mapping is bijective over all 2,048 offsets in a full BV16×Dk128 state +tile. `s_head` already contains the resolved compact/indexed persistent row, +head and `vbase`; the helper must not repeat those terms. + +Specialize a compile-time `DIRECT_STORE` boolean with +`static_assert(!DIRECT_STORE || REGSTATE)`. For eligible valid threads, after +the existing output store, execute increasing `j=0..15`: + +```cpp +Store(s_head, GdnDecodeDirectStateOffset(vi, wk, j, dk), rr[j]); +``` + +In `DIRECT_STORE=true` only, omit the complete final `rr→r` loop, second +barrier, and cooperative `sbh→s_head` loop. `DIRECT_STORE=false` must compile +the current REGSTATE body. Preserve the initial cooperative state load and its +barrier, q/k shared staging, `rr` load, all recurrence arithmetic and shuffles, +output store, dtype conversion, state index/null resolution and global layout. +Keep gx8/bx128/smem9728 even though the final shared phase disappears. + +## Red-first correctness and mutation gates + +Before product changes, extend the portable contract tests for: + +1. exact parser semantics and invalid fallback; +2. eligibility requiring every REGSTATE production/full-tile predicate; +3. exactly one incumbent or direct-store callback and unchanged geometry; +4. representative direct offsets plus exhaustive bijection of all 2,048 full + tile elements, with no duplicate, missing or out-of-range address. + +Extend CUDA comparison of REGSTATE versus DIRECT_STORE to require byte-exact +complete output and full persistent state across production and every +ineligible/fallback surface: compact/indexed state, null indices, partial +dimensions, BF16/F32 I/O, and BF16/F16/F32 state where supported. The public +graph must execute exactly one gx8/bx128/smem9728 candidate node. Graph geometry +does not prove barrier removal; compiler evidence does. + +Scratch mutations must kill the strict parser, each eligibility predicate, +callback binding, direct logical address, direct store, output store, indexed +row and null handling. Replacing DIRECT_STORE with the incumbent shared +fallback must fail exact compiler shared-op/barrier discrimination. Forcing a +partial/ineligible path into direct storage must fail portable eligibility or +CUDA exact coverage. Restore the tree byte-for-byte after every mutation. + +## Compiler and resource gate + +Build REGSTATE and DIRECT_STORE into one immutable binary and compare the exact +`bf16,bf16,float,NW8,SWIZZLED=true,REGSTATE=true` specializations with the same +toolchain. Record compiler/binary hashes, registers, stack/local/spills, +barriers, PTX/SASS totals and shared/global load/store counts. + +DIRECT_STORE must have fewer PTX **and** SASS shared loads, fewer shared stores, +and fewer barriers than REGSTATE; zero spills, zero stack and zero local array; +unchanged gx8/bx128/smem9728; and at most **5%** total SASS instruction growth. +Record global-store count because this arm trades coalescing for static shared +work, but do not infer transaction efficiency from instruction count. Use +`cudaOccupancyMaxActiveBlocksPerMultiprocessor` on actual function pointers; +an active-warps drop greater than **12.5%** rejects the arm. Missing exact +symbols, compiler no-op, or any static/resource miss removes the candidate +before timing credit. + +NCU remains owed diagnostic when local binary and counter authority exist: +profile matched y800 nodes and record global transactions/sectors, shared +requests/wavefronts, barrier/scoreboard stalls, DRAM, achieved occupancy and +instructions for REGSTATE, DIRECT_STORE and pinned vLLM. It is attribution, +never speed credit or a substitute for static/timing gates. + +## Same-binary performance and disposition + +Under one `flock /tmp/gpu` lock, use the immutable REGSTATE Qwen3.5-4B BF16 +128-request, 128-output-token, concurrency-32 workload. Hold every other +selector fixed and run: + +`REGa -> DSTa -> DSTb -> REGb`. + +Collect graph-node `nsys`, exact call/geometry/resource evidence, token hashes, +all-fused and y800 sums/counts, total/output throughput, TTFT, TPOT/ITL, E2E, +peak GPU allocation, peak/stable PSS and available-memory drop. Run a separate +counterbalanced REG/DST memory pair without profiling. + +Retain DIRECT_STORE only if all correctness, mutation, graph and resource gates +pass; both DST legs beat their paired REG controls; counterbalanced y800 +improves by at least **1.00%**; all fused decode improves; total/output +throughput do not decrease; TTFT, TPOT/ITL and E2E do not increase; every +enclosing axis and paired memory are non-regressing; and tokens are exact. A +host-PSS increase is material only when it exceeds both 1.00% and 64 MiB, then +must reproduce once before rejection. + +If compiler structure is unchanged or any bar fails, remove selector, product +and tests in the disposition change and preserve the rejection in this spec, +benchmark record, kernel/roadmap rows, STATUS and BENCHMARKS. If every bar +passes, retain only as explicit opt-in. No default, release, pinned-vLLM, +27B or 35B acceptance follows from this local discriminator. + +## Outcome — rejected and removed + +Committed experiment `9485a551473534af9af348c10aff31b949fdd840` preserved +exact tokens and the intended gx8/gy800/bx128/smem9728 graph. REGSTATE/DIRECT +STORE used 56/55 registers and zero local storage. Static PTX improved from +848 to 694 instructions, 53 to 48 shared loads, 31 to 15 shared stores and two +to one barrier, but global stores increased from 7 to 18. + +The prescribed same-binary `REGa -> DSTa -> DSTb -> REGb` series decisively +rejects the trade: y800 **134.3005475 -> 196.3294475 ms (+46.1866%)**, +**164.584 -> 240.600 us/call**; all fused decode **264.656741 -> +388.7906645 ms (+46.904%)**. Both raw DST legs lose. Enclosing REG/DST means +also regress: total **6725.73 -> 6694.67 tok/s (-0.4618%)**, output **743.715 +-> 740.28 tok/s (-0.4619%)**, TTFT **1029.48 -> 1031.48 ms (+0.1943%)**, +TPOT/ITL **35.02 -> 35.21 ms (+0.5425%)**, and E2E **5477.21 -> 5502.83 ms +(+0.4678%)**. Every token SHA-256 remains exact. Evidence roots are +`/tmp/qwen35-gdn-direct-store-{rega,dsta,dstb,regb}-9485a5514.*`. + +The lower shared traffic cannot offset the loss of global-store coalescing. +No memory pair is owed after this hard timing rejection. The exact unpublished +tip was reverted without committing; its seven product/test/env/pending-doc +paths match parent `a9d778957` byte-for-byte. DIRECT_STORE is therefore +**REJECTED/REMOVED** with no acceptance/default claim. This disposition names +no new implementation target. diff --git a/.agents/specs/sm120-qwen35-gdn-decode-regk-2026-08-09.md b/.agents/specs/sm120-qwen35-gdn-decode-regk-2026-08-09.md new file mode 100644 index 000000000..696e5f165 --- /dev/null +++ b/.agents/specs/sm120-qwen35-gdn-decode-regk-2026-08-09.md @@ -0,0 +1,213 @@ +# sm_120 Qwen3.5 fused GDN decode register-key discriminator + +**Lifecycle:** `REJECTED/REMOVED`; structural compiler no-op, no product or +test residue at disposition head `28d3a9334d50a80e16f0f3d2ccc0e0cc1b9a4593` + +**Owner rows:** `KERNEL-SSM-MAMBA`, `ROAD-V1-C2-LOCAL-BF16` + +**Hardware/workload:** RTX 5070 Ti (`sm_120`), Qwen3.5-4B plain BF16 + +**Authority:** this disposition is records-only. No product or test code, GPU +run, download, remote operation, default, release, acceptance, or 27B/35B +claim is authorized. + +## Intake, incumbent evidence, and hypothesis + +This spike is based on immutable head +`f157ed6332ab9e5a9e78b77569eac4ae29b08eba`, with current REGSTATE product +`7476818c1b2615a07f6163c0ce192a91518e1ee2`. The regular production chain is +Qwen paged decode in `src/vllm/model_executor/models/qwen3_5.cpp:3950-3969`, +public validation in `src/vt/ops.cpp:1946-1955`, strict portable selection in +`src/vt/cuda/gdn_decode_fused.h:16-161`, and CUDA dispatch/recurrence in +`src/vt/cuda/cuda_gdn.cu:2729-3014`. Current portable coverage is +`tests/vt/test_gdn_decode_fused.cpp:17-279`; CUDA exact and graph coverage is +rooted at `tests/vt/test_ops_gdn.cpp:1867-1953,3488-3645`. + +REGSTATE is exact and provisionally faster than accepted SWZ. Its prescribed +counterbalanced means are **164.57 us/call** on grid-y 800, **-1.093499%** +versus SWZ, and **-1.263502%** over all fused calls. Both REGSTATE legs beat +both bracketing SWZ legs and every enclosing mean is non-regressing. It remains +about **1.285x** the pinned-vLLM 128.061 us/call. The production graph is +gx8/bx128/smem9728, ptxas reports 56 registers, one barrier and zero stack or +spills, and every token file is exact at SHA-256 +`83fcdc45f79ddb06a634c7d7d95eba3384543b3cd781a45a8db1fc4e2a453545`. + +The retained same-tool nvcc 12.9 artifact gives the exact +`bf16,bf16,float,NW8,SWIZZLED=true` structures: SWZ/REGSTATE cubin symbols are +27,008/26,496 bytes, PTX instruction counts are 922/894, and PTX +`ld.shared`/`st.shared` counts are 69/47 versus 53/31. Global loads/stores stay +19/7. These data establish the current baseline; PTX and ELF size do not stand +in for the required exact SASS counts. + +Current REGSTATE loads `bk[sc]` in both recurrence loops: +decay/dot at `cuda_gdn.cu:2806-2811` and update/output at `:2830-2834`. +The lane requests the identical 16 K values in the identical `j=0..15` order +both times. Thus K is the next unstacked, independently falsifiable repeated +shared-read lever. + +Pinned vLLM `5559679229bc961848b121ccdeaa8fa5d79bec98` executes vendored FLA +`vllm/third_party/flash_linear_attention/ops/fused_sigmoid_gating.py:24-178`. +Its F32 `b_k` register tile is loaded once at `:124`, optionally normalized at +`:140`, and stays live for both `b_h` reductions at `:148` and `:151`. This +grounds the K lifetime only. It does not authorize FLA's different BV32 +schedule, gate arithmetic, normalization, layout, or reduction order. + +**Hypothesis:** on the exact REGSTATE production specialization, loading only +the lane's 16 K values once into registers and reusing them in both recurrence +loops removes the repeated shared K read without changing a numerical +operation. If the compiler already performs this caching, if shared loads do +not fall, or if resource growth consumes the saving, the candidate is a no-op +or a miss and must be removed before any q+k experiment. + +## Exact selector, eligibility, and implementation + +Add strict opt-in `VT_GDN_DECODE_REGK=1`. Only the complete string `"1"` +selects the candidate. Unset, empty, whitespace, prefixes/suffixes, truthy +words and every other value preserve the currently resolved path. + +REGK is eligible only when the existing contract has already resolved the +regular, non-speculative `REGSTATE=true` specialization: exact BV16, exact +SWIZZLE, `Dv==Dk==128`, resolved NW8, and the fused kernel rather than the +sequential fallback. REGK must never promote an otherwise ineligible launch. +Invalid/unset REGK and every failed predicate retain current REGSTATE when that +path is otherwise eligible; BV32, unswizzled BV16, partial dimensions, other +NW, speculative decode and sequential fallback retain their current incumbent. + +Specialize a compile-time `REGK` boolean with +`static_assert(!REGK || REGSTATE)`. In `REGK=true`, after the existing +cooperative load and synchronization, declare `float rk[16]` and preload once, +in increasing order: + +```cpp +for (int64_t j = 0; j < 16; ++j) + rk[j] = bk[GdnDecodeRegisterSharedColumn(wk, j)]; +``` + +Use the exact `rk[j]` value in both existing REGSTATE loops: replace only the +K operand in `pdot += rr[j] * bk[sc]` and the K operand in +`rr[j] += vp * bk[sc]`; the following `po` expression continues to use +`bq[sc]`. Preserve increasing +`j=0..15`, expression spelling, accumulator sequence, shuffles, and every +other operation. `REGK=false` must compile the current REGSTATE body. + +Do **not** cache q in this candidate. Do not change `rr`, state/q/k/v global +or shared layout, the 136-float row stride, global traffic, conversions, +`expf`, beta, `vp`, FMA spelling, output/state stores, accumulator/shuffle +order, null/index semantics, graph capture, public surface, or +gx8/bx128/smem9728. Rollback is `VT_GDN_DECODE_REGK=0` or unset. + +## Red-first tests and mutation obligations + +Before product code, extend the CUDA-free contract and its focused test to +prove: + +1. the REGK parser accepts only exact `"1"`, rejecting unset, empty, + whitespace, prefix/suffix, numeric and truthy lookalikes; +2. eligibility is exactly current REGSTATE plus REGK, and failing each of BV16, + SWIZZLE, Dv128, Dk128, NW8, REGSTATE or regular fused-path eligibility + selects the current incumbent rather than REGK; +3. gx8/bx128/smem9728 is unchanged and exactly one current-REGSTATE or REGK + callback fires; +4. for every `wk=0..7,j=0..15`, `rk[j]` maps to shared column `j*8+wk` and + logical column `wk*16+j`, bijectively covering 0..127. + +Extend the CUDA exact matrix to compare REGSTATE and REGK byte-for-byte for the +complete output and persistent state. Cover the production and fallback/ +partial shapes, compact and indexed state, null indices, BF16/F32 I/O, and +BF16/F16/F32 state where the public operation supports them. The public graph +test must show exactly one gx8/bx128/smem9728 REGK node. Cached Qwen3.5 +paged-forward must retain the token SHA above. + +Correctness cannot distinguish a K operand that falls back to the same shared +value, so mutation review has two structural obligations in addition to +numerical tests. Mutating the `rk` preload mapping must fail the mapping/exact +gate. Mutating **either** REGK recurrence loop back to `bk[sc]` must fail the +same-tool exact PTX shared-load discriminator (and its SASS counterpart when +available). If the compiler makes either mutation indistinguishable, REGK is a +compiler no-op and is rejected. Also kill the parser, every eligibility +predicate, callback binding, output store and persistent-state store. Restore +the tree byte-for-byte after each mutation. + +## Build, resource, correctness, and graph gates + +Build REGSTATE and REGK into one immutable CUDA binary. Run the portable +selector/mapping test, full CUDA GDN suite, cached Qwen3.5 paged-forward, staged +and full preflight, then fresh static/targeted-mutation review and independent +operator rerun. Record the compiler/binary hashes and exact specializations. + +For exact `bf16,bf16,float,NW8,SWIZZLED=true,REGSTATE=true`, compare REGSTATE +and REGK using the same compiler/toolchain. Both PTX and SASS shared-load +counts must fall; zero stack, zero spills and zero local array are mandatory; +gx8/bx128/smem9728 must be unchanged; total SASS instruction growth must be +at most **5%**. Record registers, barriers, local/stack/spill bytes, exact +PTX/SASS totals and shared/global load/store counts. Use +`cudaOccupancyMaxActiveBlocksPerMultiprocessor` on the actual function +pointers; an active-warps drop greater than **12.5%** rejects REGK. A compiler +no-op, missing exact specialization, unchanged shared loads, or any resource +bar failure triggers cleanup before timing credit. + +NCU remains an owed diagnostic when a local binary and counter authority exist: +profile the matched 400th y800 node and record shared requests/wavefronts/bank +conflicts, DRAM, scoreboard stalls, achieved occupancy and instructions for +REGSTATE, REGK and pinned vLLM. NCU is attribution only, never speed credit and +never a substitute for the static or timing gates. + +## Same-binary measurement and disposition + +Under one `flock /tmp/gpu` lock, use the immutable cached Qwen3.5-4B BF16 +128-request, 128-output-token, concurrency-32 workload from the REGSTATE spec. +Hold every selector except REGK constant and run: + +`REGa -> KREGa -> KREGb -> REGb`. + +Collect graph-node `nsys`, exact kernel/call/geometry/resource evidence, token +hashes, total and grid-y 800 fused sums/counts, total/output throughput, TTFT, +TPOT/ITL, E2E, peak GPU allocation, peak/stable PSS and available-memory drop. +Run a separate counterbalanced REG/KREG memory pair without profiling. + +Retain REGK only if all correctness, mutation, graph and resource gates pass; +both KREG legs beat their paired REG controls; the counterbalanced y800 mean +improves by at least **1.00%**; all fused decode improves; total/output +throughput do not decrease; TTFT, TPOT/ITL and E2E do not increase; every +other enclosing axis is non-regressing; tokens are exact; and paired memory +does not materially increase. Use REGSTATE's material-host-PSS rule: an +increase must exceed both 1.00% and 64 MiB, then reproduce once before reject. + +If the compiler makes REGK a no-op or any bar fails, remove its selector, +product specialization and tests in the same disposition change and preserve +the rejection in this spec, benchmark record, kernel/roadmap rows, STATUS and +BENCHMARKS. Any later hypothesis requires a separate spike and must account for +the evidence recorded in the outcome below. +If every bar passes, retain REGK only as an explicit opt-in. No default, +release, pinned-vLLM, 27B or 35B acceptance follows from this 4B discriminator. + +## Outcome — structurally rejected and removed + +The red-first missing-contract compile was captured before candidate product +work. The candidate selector, specialization and tests existed only as +uncommitted experiment changes and were restored completely; disposition head +`28d3a9334d50a80e16f0f3d2ccc0e0cc1b9a4593` contains no REGK selector, product +code or test residue. The restored portable gate passed **13/13 · 2,005**. + +The contract mutations were effective: the parser prefix mutation killed +**0/1** with two failures, regular-path eligibility failed at `-Werror`, and +both compile-time assertions killed the shared-column mapping mutation. The +critical structural mutation did not. With the same nvcc 12.9 toolchain, exact +`bf16,bf16,float,NW8,SWIZZLED=true,REGSTATE=true,REGK=true` PTX for the baseline, +the first recurrence loop changed back to `bk`, and the second loop changed +back to `bk` is byte-identical. Each exact symbol has **848 PTX instructions** +and **53 `ld.shared`**; full PTX SHA-256 is +`1fec62b7f48dc5fa2a33bfe414ce6a10aa01979c41489375a8d67b07c88a97a0`, and +both loop comparisons are identical. Incumbent REGSTATE without the candidate +already reports the same **53 `ld.shared`**, so the compiler already performs +the intended K reuse. The mandatory shared-load reduction and loop-mutation +discriminator therefore fail: REGK is a compiler no-op and is structurally +**REJECTED/REMOVED before product commit, GPU timing, or memory measurement**. + +No performance, GPU, memory, default, release, pinned-vLLM, 27B or 35B claim +results. REGSTATE remains measured/provisional with exact SASS and NCU still +pending. The next direction requires a new spike: test direct register-to-global +REGSTATE writeback that removes the final `rr` to shared write, synchronization +and shared reread, with global-store coalescing recorded as the falsifiable +tradeoff. Do not attempt q+k caching: q is single-use and K caching is now a +proven compiler no-op. diff --git a/.agents/specs/sm120-qwen35-gdn-decode-regstate-2026-08-09.md b/.agents/specs/sm120-qwen35-gdn-decode-regstate-2026-08-09.md new file mode 100644 index 000000000..d5d5ccd20 --- /dev/null +++ b/.agents/specs/sm120-qwen35-gdn-decode-regstate-2026-08-09.md @@ -0,0 +1,277 @@ +# sm_120 Qwen3.5 fused GDN decode register-resident state discriminator + +**Lifecycle:** `MEASURED/PROVISIONAL`; exact specialization SASS-size and NCU +diagnostics `PENDING` + +**Owner rows:** `KERNEL-SSM-MAMBA`, `ROAD-V1-C2-LOCAL-BF16` + +**Hardware/workload:** RTX 5070 Ti (`sm_120`), Qwen3.5-4B plain BF16 + +**Scope:** one strict, default-OFF, numerical-order-preserving register-state +variant of the accepted BV16+NW8+shared-swizzle production kernel. The spike +commit was spec-only; the measured outcome below remains default-OFF and makes +no default, release, or 27B/35B claim. + +## Gap, evidence, and falsifiable hypothesis + +Product `824370396` accepted `VT_GDN_DECODE_SWIZZLE=1` as an exact opt-in. Its +counterbalanced means improve all fused decode 7.4555% and grid-y 800 7.5123%, +but the accepted y800 call is still **166.566 us** versus pinned vLLM +**128.061 us** (**1.301x** slower). Both local arms execute gx8/bx128/reg56; +swizzled shared memory is 9,728 bytes. The accepted token SHA-256 is +`83fcdc45f79ddb06a634c7d7d95eba3384543b3cd781a45a8db1fc4e2a453545`. + +The accepted local kernel cooperatively stages the F32 `[BV,Dk]` state into +swizzled shared memory (`cuda_gdn.cu:2758-2783`), then each lane traverses its +same 16 logical columns twice: decay/dot at `:2794-2810`, followed by +update/output at `:2813-2828`. It finally synchronizes and rereads shared for +the coalesced global writeback at `:2830-2838`. Thus each logical state element +participates in an initial shared store, two shared read-modify-writes, and a +writeback shared load. + +**Hypothesis:** on the exact production specialization, retaining the lane's +16 F32 state values in registers between the two recurrence loops removes one +shared read and one shared write per element without changing any arithmetic +or logical access. If compiler resource growth, spills, lost occupancy, or +shared-to-register moves consume the saving, the candidate will not clear the +gates and must be removed. + +## Complete execution-chain ground + +The local production path is: + +1. Qwen's paged step selects the decode segment and indexed persistent cache at + `src/vllm/model_executor/models/qwen3_5.cpp:3950-3969`. +2. Public validation and dispatch are `src/vt/ops.cpp:1946-1955`. +3. CUDA validates dtypes/state, resolves NW8 and the shared-budget fallback at + `src/vt/cuda/cuda_gdn.cu:2920-2950`; BV and swizzle environment dispatch is + at `:2968-2969`. +4. The portable strict selectors, production-only swizzle eligibility and + 9,728-byte launch contract are + `src/vt/cuda/gdn_decode_fused.h:16-78`. +5. The executing template, shared staging, recurrence and launch binding are + `src/vt/cuda/cuda_gdn.cu:2729-2871`. + +Pinned vLLM `5559679229bc961848b121ccdeaa8fa5d79bec98` calls the vendored FLA +Triton kernel in +`vllm/third_party/flash_linear_attention/ops/fused_sigmoid_gating.py:24-178`. +Its logical F32 `[BV,BK]` value `b_h` is created and loaded at `:102-120`, kept +across decay/dot/update/output at `:122-154`, and stored at `:156-170`. The +wrapper resolves `BK=128`, `BV=32`, grid `(NK,NV,N*HV)`, four warps and three +stages at `:205-212,241-277`. + +The exact pin-generated artifact is +`.triton-cache/D5PNU7TCKM4WKZZ6SN22CRUFGNRRSCWEBBGHZEFA6WQ3V7FVKNAA/`: +`fused_sigmoid_gating_delta_rule_update_kernel.source:1-28` binds the generated +IR back to upstream line 24; its JSON records sm_120, Triton 3.7.1, four warps, +three stages and 512 shared bytes. The matched production trace records +grid `1x4x800`, block 128, 96 registers and 128.061 us/call. This is evidence +that a register-heavier schedule exists, not permission to copy Triton's +different BV32 reduction schedule or its fused gate arithmetic. + +## Exact candidate and rollback contract + +Add strict opt-in `VT_GDN_DECODE_REGSTATE=1`. Only exact string `"1"` selects +the candidate. Unset, empty, whitespace, prefixes/suffixes, truthy words and +every other value retain the accepted shared-state kernel. + +Eligibility is the conjunction of: + +- resolved `VT_GDN_DECODE_BV=16`; +- resolved `VT_GDN_DECODE_SWIZZLE=1`; +- `Dv==Dk==128` and resolved `NW==8`; +- the regular, non-speculative `GdnDecodeFusedKernel` path. + +Every partial shape, BV32/default, unswizzled BV16, other NW, sequential +fallback, and speculative-decode kernel remains byte-for-byte on its incumbent +path. Specialize a compile-time `REGSTATE` boolean; do not add a runtime branch +inside either recurrence loop. + +For `REGSTATE=true`, after the existing cooperative load and synchronization: + +1. declare `float rr[16]` and load in increasing `j` order + `rr[j] = sbh[vi*136 + j*8 + wk]`; +2. execute the existing decay/dot loop in the same `j=0..15` order using + `rr[j]`, with identical expressions and accumulator; +3. retain the existing XOR shuffle sequence `off=1,2,4`; +4. execute the existing update/output loop in the same `j=0..15` order using + `rr[j]`, with identical expressions and accumulator; +5. retain the existing output conversion/store and shuffle sequence; +6. write `rr[j]` back to the same swizzled shared addresses in increasing `j`, + synchronize, then use the existing coalesced global state writeback. + +Do not change state/q/k/v global layout, shared mapping or 136-float row +stride; gx8/bx128/smem9728; input/output/state dtype conversion; `expf`, beta, +`vp`, FMA spelling; accumulator order; shuffle order; null/index semantics; or +public surface. Rollback is `VT_GDN_DECODE_REGSTATE=0` or unset. The feature +remains explicit opt-in even if accepted. + +## Red-first portable and CUDA tests + +Extend `tests/vt/test_gdn_decode_fused.cpp` (existing selector/layout coverage +at `:17-160`) before implementation: + +1. exact `"1"` parser plus unset/empty/whitespace/prefix/suffix/truthy invalids; +2. production eligibility requires BV16+swizzle+Dv128+Dk128+NW8; +3. every individual failed predicate selects the incumbent specialization; +4. the resolved contract remains gx8/bx128/smem9728 and exposes exactly one + shared-state or register-state callback; +5. for all `wk=0..7,j=0..15`, register slot `j` maps to shared `j*8+wk` and + logical column `wk*16+j`, bijectively covering 0..127. + +Extend the CUDA exact matrix rooted at `tests/vt/test_ops_gdn.cpp:1867-1953` +to compare accepted SWZ versus REGSTATE byte-for-byte for complete output and +persistent state across production and fallback/partial shapes, compact and +indexed state, null indices, BF16/F32 I/O, and BF16/F16/F32 state where the +public op supports them. Extend the public graph test at `:3488-3577` to prove +REGSTATE executes exactly one gx8/bx128/smem9728 node. Run cached Qwen3.5 +paged-forward and require the exact production token SHA above. + +Scratch mutations must make the focused test fail when they corrupt or remove +the strict parser, production eligibility, one-callback dispatch, `rr[j]` +load mapping, `rr[j]` writeback mapping, output store, or persistent-state +store. A candidate that merely compiles while silently taking SWZ is a failure. + +## Compiler/SASS and resource gate + +Build baseline and candidate into one immutable binary, then capture the two +template specializations from the same toolchain: + +```sh +nix develop .#cuda --command cmake --build build-nix-cuda-transplant-triton --target vllm-bench test_gdn_decode_fused test_ops_gdn +nix develop .#cuda --command cuobjdump --dump-resource-usage --demangle build-nix-cuda-transplant-triton/examples/vllm-bench +nix develop .#cuda --command cuobjdump --dump-sass --demangle build-nix-cuda-transplant-triton/examples/vllm-bench +nix develop .#cuda --command cuobjdump --dump-resource-usage .triton-cache/D5PNU7TCKM4WKZZ6SN22CRUFGNRRSCWEBBGHZEFA6WQ3V7FVKNAA/fused_sigmoid_gating_delta_rule_update_kernel.cubin +``` + +Record exact binary/cubin SHA-256, compiler versions, registers/thread, static +and dynamic shared, stack/local bytes, spill loads/stores, static SASS +instruction count, and shared-load/shared-store instruction counts. Candidate +must have zero spills and zero local state array, unchanged dynamic shared and +block geometry, fewer static shared-memory instructions, and no more than 5% +total static SASS instruction growth. Use +`cudaOccupancyMaxActiveBlocksPerMultiprocessor` on the two actual function +pointers; a predicted active-warps drop greater than 12.5% is material and +rejects the arm. Resource data describe compiler output only; they do not +substitute for hardware counters. + +## Fallback performance experiment without NCU + +Nsight Compute is **unavailable on this local workstation**: `ncu --version`, +`/usr/local/cuda-13.0/bin/ncu --version`, and +`nix develop .#cuda --command ncu --version` all fail because no local binary +is installed; `flake.nix` supplies Nsight Systems, not Nsight Compute. The +repository's `ncu 2025.3.1.0` evidence belongs to DGX, whose use is not +authorized for this task. Counter permission and any local installation are +external authorization gates. Therefore no bank-conflict, scoreboard, DRAM, +or achieved-occupancy conclusion may be inferred in this experiment. + +The falsifiable fallback is same-binary graph-node `nsys`, under the configured +GPU lock, in counterbalanced order `SWZa -> REGa -> REGb -> SWZb`, using the +accepted cached checkpoint and exact workload: + +```sh +GPU_LOCK=/tmp/gpu +MODEL=/home/rich/c/vllm.cpp/.hf-cache/hub/models--Qwen--Qwen3.5-4B/snapshots/851bf6e806efd8d0a36b00ddf55e13ccb7b8cd0a +BIN=/home/rich/c/vllm.cpp/build-nix-cuda-transplant-triton/examples/vllm-bench +flock "$GPU_LOCK" env LD_LIBRARY_PATH=/run/opengl-driver/lib VT_RELEASE_HOST_WEIGHTS=1 VT_DIRECT_DEVICE_LOAD=1 VT_GDN_DECODE_BV=16 VT_GDN_DECODE_SWIZZLE=1 VT_GDN_DECODE_REGSTATE=0 nsys profile --force-overwrite=true --sample=none --cpuctxsw=none --trace=cuda,nvtx --cuda-graph-trace=node --output=/tmp/qwen35-gdn-regstate-swza "$BIN" --model "$MODEL" --dataset-path /tmp/qwen35-4b-sharegpt-1024.json --num-prompts 128 --output-len 128 --concurrency 32 --temperature 0 --max-num-batched-tokens 2048 --num-blocks 1280 --output-token-ids /tmp/qwen35-gdn-regstate-swza.tokens.json +``` + +Repeat the exact command for `REGa`, `REGb`, and `SWZb`, changing only +`VT_GDN_DECODE_REGSTATE`, output names, and using one lock across the complete +series. Record all-fused and grid-y 800 sums/counts, launch geometry/resources, +token SHA, total/output throughput, TTFT, TPOT/ITL, E2E, peak GPU allocation, +peak/stable host PSS and available-memory drop. Run a separate counterbalanced +SWZ/REG memory pair without `nsys`. + +## Acceptance, rejection, and owed diagnostic + +Retain REGSTATE only if all correctness, mutation, graph and resource gates +pass; both REG samples beat their paired SWZ control; the counterbalanced y800 +mean improves by at least **1.00%**; all-fused decode improves; total/output +throughput do not decrease; TTFT, TPOT/ITL and E2E do not increase; and GPU +allocation does not materially increase. A host-PSS increase is material when +it exceeds both **1.00% and 64 MiB**; repeat such a move once, and reject if it +reproduces. Any spill/local array, >12.5% predicted active-warp loss, material +memory regression, mixed sample direction, or enclosing regression rejects the +candidate. Rejection removes the selector, specialization and tests while +preserving the refutation in this spec, benchmark record, kernel/roadmap rows, +STATUS and BENCHMARKS. Acceptance updates the same records but remains opt-in; +default/release/27B/35B gates stay open. + +When a local NCU binary and counter authority become available, the still-owed +diagnostic profiles the 400th per-launch-configuration y800 node under +`--replay-mode application --graph-profiling node --target-processes all` and +records shared requests/wavefronts/bank conflicts, DRAM bytes/throughput, +short/long-scoreboard stalls, achieved occupancy and executed instructions for +SWZ, REGSTATE and pinned vLLM. It is diagnostic attribution, not retroactive +speed credit, and must never be marked complete from SASS or timing inference. + +## Measured outcome, 2026-08-09 + +**Disposition:** `MEASURED/PROVISIONAL`, retained as an explicit opt-in at +product commit `7476818c1`. Correctness, graph geometry, observed compiler +resources, counterbalanced timing, enclosing performance and memory satisfy +their measured bars, but the exact-specialization static SASS-size <=5% check +and the NCU diagnostic remain pending external download/tool authority. +Therefore this is not a fully `ACCEPTED` candidate and receives no default, +release, or 27B/35B credit. + +The full canonical preflight is green and the fresh static plus targeted +mutation re-review reports `PASS`. Operator gates are portable selector/mapping +**10/10 · 1,962**, CUDA GDN **70/70 · 4,830**, and cached Qwen3.5 +paged-forward **4/4 · 8**, with all observed diffs zero. Every performance and +memory token file has SHA-256 +`83fcdc45f79ddb06a634c7d7d95eba3384543b3cd781a45a8db1fc4e2a453545`. +The graph contract is exactly grid x/y `8/800`, block x 128, dynamic shared +9,728 bytes, 56 registers and zero local bytes. + +The same-tool ptxas log +`/tmp/cuda_gdn-regstate-candidate-7476818c1.log` identifies the production +`GdnDecodeFusedKernel` specialization at 56 +registers, one barrier, zero-byte stack frame and zero spill loads/stores. The +accepted `SWIZZLED=true, REGSTATE=false` specialization has the same reported +resources. This proves no observed register, barrier, stack or spill regression; +it does not substitute for the still-missing exact static SASS instruction-size +and shared-instruction counts. + +The retained nvcc 12.9 intermediate build at `/tmp/cuda-gdn-keep.gXumAL/` +adds non-substitute compiler structure. `readelf` reports exact sm_120a cubin +ELF symbol sizes 27,008 bytes for SWZ and 26,496 bytes for REGSTATE +(**-1.896%**). Exact PTX function counts are 922/894 total instructions, +69/53 `ld.shared`, 47/31 `st.shared`, unchanged 19/7 global loads/stores and +two barriers for SWZ/REGSTATE. Thus the generated PTX removes exactly 16 shared +loads and 16 shared stores and the cubin symbol shrinks, supporting the intended +mechanism and <=5% byte-size direction. These are not the required SASS +instruction/shared-op counts, so the cuobjdump gate remains pending. Artifact +SHA-256 values are `dc500503...e2e6` for `cuda_gdn.ptx` and +`5169611e...d913` for `cuda_gdn.sm_120a.cubin`. + +The prescribed four-leg series is +`SWZa -> REGa -> REGb -> SWZb`. Both raw REGSTATE legs beat both bracketing SWZ +legs on grid-y 800 and all fused decode. Counterbalanced means are: + +| Axis | SWZ | REGSTATE | Change | +|---|---:|---:|---:| +| grid-y 800, 816 calls | 135.7758125 ms | 134.2911050 ms | **-1.093499%** | +| all fused decode, 1,704 calls | 268.1141915 ms | 264.7265645 ms | **-1.263502%** | +| total throughput | 6726.240 tok/s | 6739.125 tok/s | **+0.191563%** | +| output throughput | 743.770 tok/s | 745.195 tok/s | **+0.191591%** | +| mean TTFT | 1026.800 ms | 1021.745 ms | **-0.492306%** | +| mean TPOT / ITL | 35.000 ms | 35.000 ms | unchanged | +| mean E2E | 5471.960 ms | 5466.420 ms | **-0.101243%** | + +Evidence roots are +`/tmp/qwen35-gdn-regstate-{swza,rega,regb,swzb}-7476818c1.*`. The separate +non-profiled memory pair is +`/tmp/qwen35-gdn-regstate-{swz,reg}-mem-7476818c1.*`: peak GPU allocation is +13,058/13,060 MiB and peak host PSS is 2,436,964/1,989,479 KiB for SWZ/REGSTATE. +The candidate's PSS is lower, so no repeat is required; timing from this memory +pair is non-binding. + +`cuobjdump` and NCU are not locally available without the separately requested +downloads, those downloads are not authorized, and the Ordino service is +unavailable. The next disposition gate is to capture exact SWZ/REGSTATE static +SASS size and shared load/store counts, then run the specified NCU diagnostic +when both the binary and counter authority become available. Until then the +selector stays strict opt-in only. diff --git a/.agents/specs/sm120-qwen35-gdn-decode-rpt2-2026-08-09.md b/.agents/specs/sm120-qwen35-gdn-decode-rpt2-2026-08-09.md new file mode 100644 index 000000000..2bc8fc196 --- /dev/null +++ b/.agents/specs/sm120-qwen35-gdn-decode-rpt2-2026-08-09.md @@ -0,0 +1,126 @@ +# sm_120 Qwen3.5 fused GDN decode BV32 RPT2 discriminator + +**Rows:** `KERNEL-SSM-MAMBA`, feeding `ROAD-V1-C2-LOCAL-BF16`. +**Lifecycle:** `COMPLETE` — RPT2 was correctness-exact but performance-rejected +and removed at cleanup `f6a0c879141a1887f9fcf8c85a9eb3a8cf8a6275`. BV16 remains the sole opt-in and +BV32 remains the shipped default. + +## Gap, anchors and hypothesis + +Current cleanup `634ccba70` and result record `3a00d4b82` retain only exact +`VT_GDN_DECODE_BV=16`; BV8/BV24 were slower and removed. The latest same-binary +series measures BV16's dominant grid-y 800 shape at **167.858 us/call** versus +pinned vLLM **128.061 us/call** (**1.311x**). The prior 1.276x run remains +separate variance evidence. + +Local `GdnDecodeFusedCuda` -> `DispatchGdnDecodeValueTile` -> +`LaunchGdnDecodeFused{S,NW}` -> `GdnDecodeFusedKernel` is anchored at +`src/vt/cuda/gdn_decode_fused.h:11-67` and +`src/vt/cuda/cuda_gdn.cu:2730-2931`. At Dv=Dk=128, exact BV32 uses NW8 per +value row: grid-x 4, block 256/eight warps and 17,536 bytes dynamic shared. +BV16 obtains block 128/four warps, but doubles grid-x and repeated q/k loads. + +Pinned vLLM `5559679229bc961848b121ccdeaa8fa5d79bec98` vendors FLA +`fused_sigmoid_gating.py:24-277`: its recurrence is `:122-170`, +`BV=min(next_power_of_2(V),32)` is `:205-212`, grid `(NK,NV,N*HV)` is `:241`, +and `num_warps=4` is `:212,275`. The profiler proves this generated Triton +kernel executes; CUTLASS, cuBLASLt and DeepGEMM are not in this recurrence. + +The discriminator combines FLA's BV32/grid-4/four-warp schedule with the +local kernel's exact NW8 arithmetic: each eight-lane row group processes two +independent value rows sequentially. This tests whether BV16's gain came from +four-warp blocks without paying its duplicated grid/q/k traffic. + +## Exact design and selector + +Add strict opt-in `VT_GDN_DECODE_RPT=2`. Only exact string `"2"` selects RPT2; +unset, empty, whitespace, signs, zero-padding, suffixes/prefixes and every +other value select RPT1. Apply RPT2 only when the selected value tile is BV32 +and `dv>=32`; BV16 and `dv<32` remain RPT1 regardless of this environment +value. Existing `VT_GDN_DECODE_BV` semantics are unchanged. + +Specialize and unroll RPT at compile time. For RPT2, `groups=ceil(BV/RPT)=16`, +block threads are `groups*NW=128`, and each lane uses +`vi=(tid/NW)+row_iter*groups` for `row_iter=0,1`. Each row retains its same +eight consecutive lanes, Dk slice `[c0,c1)`, shuffle tree, decay, dot, update, +output reduction, output store and state writeback order. Do not change NW, +interleave two rows' recurrence, or reuse one row's accumulator for another. + +State-tile load/write use block 128 while preserving the same linear element +mapping and value semantics. The full padded BV32 shared tile remains live, so +Dv=Dk=128 resolves to **grid-x 4, block 128, dynamic shared 17,536 bytes**. +Tail rows stay shuffle-live on zero state and guarded v/output/state accesses. +The existing null indexed-state return and >48 KiB sequential fallback remain +unchanged. + +## RED-first tests and mutation review + +Extend the CUDA-free contract before product code. It must prove strict RPT +parsing; BV32 RPT1/RPT2 geometry; partial-tile ceiling/tails; `dv<32` fallback; +BV16 unchanged for RPT unset/valid/invalid; and exact callback dispatch into +one RPT1 or RPT2 specialization. Mutations deleting the RPT2 callback, +changing its production binding, row-group count or row stride must be RED. + +Extend CUDA GDN coverage to byte-compare candidate output **and full final +state** against both BV32 RPT1 and BV16. Cover production and partial Dv/Dk, +compact and indexed state including a null row, BF16/F32 state and BF16/F32 +input/output coverage already supported by the public op. Run both RPT arms in +one process. Mutations that skip the second row, corrupt its stride, remove +the RPT2 production dispatch, or suppress either output/state store must fail. +An actual public CUDA graph node must prove gx=4, bx=128 and smem=17,536. + +Focused gates are `test_gdn_decode_fused`, full CUDA `test_ops_gdn`, +`test_qwen35_paged_forward`, staged/full preflight, then fresh static and +targeted-mutation review and independent operator rerun. Every production token +file must equal SHA-256 +`83fcdc45f79ddb06a634c7d7d95eba3384543b3cd781a45a8db1fc4e2a453545`; +any mismatch stops performance work. + +## Same-binary decision gate + +Under one GPU lease and `flock /tmp/gpu`, use one rebuilt binary and the same +cached Qwen3.5-4B BF16 workload: 128 ShareGPT requests, 128 output tokens, +c32, `max_num_batched_tokens=2048`, 1,280 KV blocks and greedy sampling. Keep +all other selectors identical. Run the counterbalanced sequence +`BV16a -> RPT2a -> RPT2b -> BV16b`, collecting +`nsys --cuda-graph-trace=node`, SQLite, raw summaries, exact token hashes, +binary/head SHA, clocks/contention, peak VRAM and peak PSS. + +RPT2 is eligible only if its counterbalanced mean improves grid-y 800 by at +least **1.00%** and improves all fused-decode calls versus BV16, with both +samples moving in the winning direction. It must also be non-regressing on +every enclosing mean: total/output throughput may not decrease; TTFT, +TPOT/ITL, E2E, peak VRAM and peak PSS may not increase. Record values and +ratios for every axis. + +If any correctness, geometry, micro or enclosing bar fails, remove the RPT +selector and all product/test arms while preserving the refutation. If all +bars pass, retain RPT2 only as an opt-in; BV16 remains available, shipped +default status and Qwen3.6-27B/35B release gates remain open. No 4B result is +extrapolated to the unavailable release vehicles. + +## Result — rejected and removed + +Measured product `6ac8bf390` passed the operator gates: portable contract +**4/4 · 125**, CUDA GDN **69/69 · 5,046**, and Qwen3.5 paged-forward +**4/4 · 8**. All four production token files have SHA-256 +`83fcdc45f79ddb06a634c7d7d95eba3384543b3cd781a45a8db1fc4e2a453545`. +The intended Dv=Dk=128 launch executed: BV16 grid-x 8 / block 128 / +9,280-byte shared / 50 registers; RPT2 grid-x 4 / block 128 / 17,536-byte +shared / 56 registers. + +The counterbalanced `BV16a -> RPT2a -> RPT2b -> BV16b` series contained +1,704 total and 816 grid-y 800 calls. BV16 versus RPT2 means were +**271.577988 -> 372.4151745 ms (+37.1301%)** across all fused decode and +**136.9415415 -> 190.121047 ms (+38.8337%)** at grid-y 800. The enclosing +means also all reject RPT2: total/output throughput +**6732.58/744.47 -> 6709.16/741.88 tok/s**, TTFT +**1026.28 -> 1027.645 ms**, TPOT/ITL **35.00 -> 35.145 ms**, and E2E +**5471.38 -> 5490.87 ms**. Thus RPT2 fails both kernel bars and every enclosing +axis despite exact correctness. + +Cleanup `f6a0c879141a1887f9fcf8c85a9eb3a8cf8a6275` removes the selector, specialization and tests. +BV16 remains the sole exact opt-in. Its current-series dominant call is +**167.82 us/call** versus pinned vLLM **128.061 us/call**, about **1.310x** +slower, so default, release, 27B and 35B gates remain open. Evidence is +`/tmp/qwen35-gdn-rpt2-{bv16a,rpt2a,rpt2b,bv16b}-6ac8bf390.*`. diff --git a/.agents/specs/sm120-qwen35-gdn-decode-swizzle-2026-08-09.md b/.agents/specs/sm120-qwen35-gdn-decode-swizzle-2026-08-09.md new file mode 100644 index 000000000..65c6da4f1 --- /dev/null +++ b/.agents/specs/sm120-qwen35-gdn-decode-swizzle-2026-08-09.md @@ -0,0 +1,139 @@ +# sm_120 Qwen3.5 fused GDN decode BV16/NW8 shared-bank swizzle + +**Lifecycle:** `COMPLETE` — accepted opt-in at product `824370396`; defaults +and release gates remain open. + +**Owner rows:** `KERNEL-SSM-MAMBA`, `ROAD-V1-C2-LOCAL-BF16` + +**Hardware/workload:** RTX 5070 Ti (`sm_120`), Qwen3.5-4B plain BF16 + +**Authority:** records and docs only in this checkpoint; no product, tests, GPU, +remote operations or release-default change. + +## Gap and ground truth + +BV16 is the sole retained opt-in (`VT_GDN_DECODE_BV=16`); BV8, BV24 and BV32 +RPT2 were exact but performance-rejected and removed. The latest +counterbalanced series measures BV16's dominant grid-y 800 call at **167.82 +us/call** versus pinned vLLM **128.061 us/call** (**1.310x** slower), so the +residual remains open. The incumbent implementation is +`src/vt/cuda/cuda_gdn.cu:GdnDecodeFusedKernel` and its portable selector and +geometry contract is `src/vt/cuda/gdn_decode_fused.h`; coverage lives in +`tests/vt/test_gdn_decode_fused.cpp` and `tests/vt/test_ops_gdn.cpp`. Pinned +vLLM/FLA ground remains `fla/ops/gated_delta_rule/fused_recurrent.py` and +`fla/ops/gated_delta_rule/wy_fast.py` through vLLM's fused recurrent wrapper, +as recorded by the BV16 spike. This is a local layout discriminator, not a +claim that the candidate mirrors FLA's layout. + +The present BV16/Dk128/NW8 schedule gives lane `wk` the 16 contiguous columns +`c=wk*16+j`. For each memory instruction those lanes cluster in banks, while +one warp contains four independent eight-lane value-row groups. State shared +rows use stride `dk+1=129`, so corresponding groups do not deliberately cover +disjoint bank ranges. The hypothesis is that an exact shared-only swizzle can +reduce conflicts without changing global traffic, arithmetic or launch width. + +## Exact discriminator and layout + +Add strict opt-in `VT_GDN_DECODE_SWIZZLE=1`: only exact string `"1"` enables +it. Unset, empty, whitespace, prefixes/suffixes, numeric lookalikes and every +other value select the incumbent layout. It is eligible only when the resolved +value tile is BV16, `Dv==Dk==128`, and resolved `NW==8`; every other geometry, +including partial tiles and all BV32/default paths, uses the old layout. + +Specialize the kernel/launcher on a compile-time layout boolean. With +`ck=Dk/NW=16`, map logical column `c` to shared column +`sw=(c%ck)*NW+c/ck`. Thus lane `wk` at its existing local iteration `j` reads +and writes `sw=j*NW+wk`. Apply the same mapping to shared q and k; the four row +groups in a warp request the same q/k address and retain multicast behavior. +Use swizzled state row stride `Dk+NW=136`, shifting adjacent value-row bases by +eight banks so the four eight-lane groups cover disjoint ranges. The production +shared allocation is +`(2*128 + 16*136)*sizeof(float) = 9,728 bytes`; launch geometry remains grid-x +8, block-x 128. Keep the incumbent `dk+1` state stride and linear q/k layout in +the false specialization. + +The load, two recurrence loops, writeback and output/state stores address the +same logical elements. Per-row operation order remains `j=0..15` for each lane, +so output and persistent state must be byte-exact. Resolve all constants and +address forms outside or at compile time: the recurrence inner loops contain +no division or modulo. Do not change input/output/state dtype conversion, +shuffle reduction order, global layout, null-slot semantics, cache indexing, +grid topology, default selection, BV16 selection, or any public surface. + +## Red-first tests and mutation obligations + +Extend the CUDA-free contract tests for strict parser behavior, eligibility and +fallback, the 128-column bijection, representative indices (including +`c=0,15,16,127` and `j*8+wk`), state stride 136, shared bytes 9,728, and the +unchanged grid-x 8/block-x 128 geometry. Prove partial Dv/Dk, NW other than 8, +BV32/default and invalid selectors retain the incumbent layout and allocation. + +CUDA must compare swizzle OFF/ON byte-for-byte for both output and the complete +persistent state across the production 128/128 shape, partial/fallback shapes, +compact and indexed caches, negative/null indices, and the supported BF16/F32 +activation/output and cache-state paths. Capture the public `vt::GdnDecode` +graph and require exactly one production kernel node with grid-x 8, block-x +128 and 9,728 dynamic shared bytes when both `BV=16` and `SWIZZLE=1` are set. +Run the full CUDA GDN and Qwen3.5 paged-forward gates after the focused tests. + +Scratch mutations must make the focused gate fail when they corrupt or remove: +the swizzle mapping, state-row stride, production eligibility/binding, output +store, or state writeback. Record red-before implementation and restore every +mutation byte-for-byte. + +## Performance gate and disposition + +Build once, pin the exact executable/model/prompt/token SHA, acquire the local +GPU contention lock, and run graph-node `nsys` in counterbalanced order +`BV16a -> SWZa -> SWZb -> BV16b` on the identical cached Qwen3.5-4B BF16 +128-request/128-output-token c32 workload. Record raw reports, kernel geometry, +registers, all-call and grid-y 800 sums/counts, token SHA-256, total/output +throughput, TTFT, TPOT/ITL, E2E, peak VRAM and peak host PSS. + +Retain the swizzle only as an opt-in if both swizzled samples are byte/token +exact, the counterbalanced grid-y 800 mean improves by at least **1.00%** over +BV16, the all-fused-decode mean improves, both samples move in the winning +direction, and every enclosing axis is non-regressing (throughput may not +decrease; latency and memory may not increase). A passing 4B result does not +close default, release, Qwen3.6-27B or 35B gates. If any correctness, geometry, +microbenchmark or enclosing bar fails, remove the selector and all candidate +product/test arms while preserving the measured refutation in this spec and +the benchmark record. BV16 remains the incumbent in either disposition. + +## Measured result — accepted opt-in + +Spike `bdfaf823e` and product `824370396` implement the strict shared-swizzle +arm. Operator gates pass: portable **6/6 · 633**, CUDA GDN +**69/69 · 4,742**, and Qwen3.5 paged-forward **4/4 · 8**. Parser and mapping +negative mutations fail with 5 and 132 assertions respectively. All six +production token files have SHA-256 +`83fcdc45f79ddb06a634c7d7d95eba3384543b3cd781a45a8db1fc4e2a453545`. +Both arms execute grid-x 8 / block 128 at 56 registers; shared memory changes +from 9,280 to 9,728 bytes as designed. + +The counterbalanced `BV16a -> SWZa -> SWZb -> BV16b` series has 1,704 total +and 816 grid-y 800 calls per arm. Swizzle improves all fused decode +**289.8737275 -> 268.254130 ms (-7.4555%)** and y800 +**146.959362 -> 135.9178535 ms (-7.5123%)**. Both swizzle samples win. +Enclosing means also improve: total/output throughput +**6724.18/743.54 -> 6739.86/745.275 tok/s**, TTFT +**1027.34 -> 1023.005 ms**, TPOT/ITL **35.045 -> 34.98 ms**, and E2E +**5478.205 -> 5465.725 ms**. + +The memory pair remains positive for the user-visible and device axes: +total/output **6784.64/750.23 -> 6793.05/751.16 tok/s**, TTFT +**1020.77 -> 1018.67 ms**, TPOT **34.72 -> 34.68 ms**, E2E +**5430.24 -> 5423.44 ms**, and GPU allocation **13,060 -> 13,054 MiB**. +Peak host PSS alone rises **1,909,259 -> 2,038,631 KiB (+6.78%)**, while +available memory drops less in the swizzle arm (**1,641,476 -> 1,529,724 KiB** +available-drop basis supplied by the run). This noisy transient host-load axis +is preserved as an acceptance caveat; the arm remains explicit opt-in and +receives no default or release credit. + +Evidence is +`/tmp/qwen35-gdn-swizzle-{bv16a,swza,swzb,bv16b}-8243703.{nsys-rep,sqlite,log,tokens.json}` +and `/tmp/qwen35-gdn-swizzle-{bv16,swz}-mem-8243703.{jsonl,log,tokens.json}`. +Trace SHA-256 prefixes are `59bb6081` (BV16a), `082cd6d4` (SWZa), +`b44792c8` (SWZb), and `633fad31` (BV16b). The accepted swizzle y800 call is +**166.566 us** versus pinned vLLM **128.061 us**, leaving a **1.301x** residual. +Default, release, 27B and 35B gates remain open. diff --git a/.agents/specs/sm120-qwen35-pareto-2026-08-09.md b/.agents/specs/sm120-qwen35-pareto-2026-08-09.md new file mode 100644 index 000000000..fff40af6f --- /dev/null +++ b/.agents/specs/sm120-qwen35-pareto-2026-08-09.md @@ -0,0 +1,130 @@ +# RTX 5070 Ti Qwen3.5-4B Pareto campaign + +**Issue:** [#206 — RTX 5070 Ti: close Qwen3.5-4B TTFT, TPOT, and +VRAM gaps vs vLLM](https://github.com/mudler/vllm.cpp/issues/206) + +**Rows:** `KERNEL-SSM-MAMBA`, feeding `ROAD-V1-C2-LOCAL-BF16` + +**Lifecycle:** SPEC COMMITTED; code transplant and fresh measurement pending + +**Target:** local NVIDIA GeForce RTX 5070 Ti (`sm_120`), Qwen3.5-4B plain +BF16. This vehicle does not substitute for the hardware-unavailable 27B/35B +release gates. + +## Binding baseline and objective + +The sealed production comparison uses the pinned vLLM oracle at `555967922`, +the cached `Qwen/Qwen3.5-4B` snapshot `851bf6e...`, 128 ShareGPT requests, +128 output tokens, concurrency 32, `max_num_batched_tokens=2048`, 1,280 KV +blocks and greedy sampling. Both engines use the identical model artifacts, +prompts and production configuration; vLLM is not run with `--enforce-eager`. + +| Axis | vllm.cpp | pinned vLLM | Status | +|---|---:|---:|---| +| total throughput | **6784.743 tok/s** | 6643.593 tok/s | **PASS**, 1.021246x | +| mean TTFT | 1018.040 ms | **937.584 ms** | **OPEN**, 1.085812x | +| mean TPOT / ITL | 34.740 ms | **33.906 ms** | **OPEN**, 1.024597x | +| peak VRAM | 13053.3 MiB | **12820 MiB** | **OPEN**, 1.018199x | + +The objective is a Pareto improvement: first close the three axes where local +is behind without surrendering its throughput lead or exact correctness; once +all measured axes pass, continue only on the largest same-tool traceable cost +whose improvement preserves every passing axis. + +Binding baseline and reproduction detail: +[Qwen3.5-4B sm_120 evidence](../../docs/bench-evidence/qwen35-4b-sm120-main-20260807.md). + +## Source ground and existing evidence + +Current local exact-chunk dispatch and causal-conv kernel are grounded at +`include/vllm/v1/attention/backends/gdn_attn.h` and +`src/vt/cuda/cuda_gdn.cu:702-854`. The pinned vLLM work partition is +`${VLLM_SOURCE}/vllm/model_executor/layers/mamba/ops/causal_conv1d.py:15-28,71-79,123-124,692-742`. +The fused post-conv schedule is grounded at +`${VLLM_SOURCE}/vllm/third_party/flash_linear_attention/ops/fused_gdn_prefill_post_conv.py:57-149,208-245` +and its tests at `${VLLM_SOURCE}/tests/kernels/test_fused_gdn_post_conv.py:60-208`. +The fused decode comparison is grounded in +`${VLLM_SOURCE}/vllm/third_party/flash_linear_attention/ops/fused_sigmoid_gating.py:24-277` +and local `src/vt/cuda/cuda_gdn.cu:GdnDecodeFusedKernel`. + +The historical branch evidence imported with this spec is authoritative for +what was tried, including negative results. In particular: + +- [K=4 causal-conv specialization](sm120-qwen35-conv-channel-tile-2026-08-08.md) + was byte-exact and locally positive; its two-channel arm was slower. +- [16-token fused post-conv tile](sm120-qwen35-postconv-token-tile-2026-08-08.md) + was byte-exact and locally positive. +- Decode BV16, swizzle and provisional register-state candidates are retained + as later evidence; BV8/BV24, RPT2, REGK and scalar direct-store are recorded + failures and must not be silently repeated. +- [BF16 vector writeback](sm120-qwen35-gdn-decode-bf16-vector-writeback-2026-08-09.md) + is only a spiked later discriminator, not the first campaign step. + +## Ordered campaign + +1. Transplant the old accepted opt-in implementations and their tests onto + current `upstream/main`, without changing defaults. Resolve product changes + manually against current source; do not merge keyed records or resurrect a + state log. +2. First measure the **combined K=4 causal-conv plus 16-token post-conv tile** + in one immutable same binary. Run counterbalanced + `BASEa -> COMBINEDa -> COMBINEDb -> BASEb`, with exact token hashes and the + complete throughput, TTFT, TPOT/ITL, E2E, GPU-memory and host-PSS axes. The + earlier separate wins are not assumed additive. +3. Trace that accepted combined candidate and pinned vLLM with the same `nsys` + command and workload, both with `--cuda-graph-trace=node`. Compare the full + request interval and separate prefill, decode, scheduler/API and memory + costs. Source selects hypotheses; these matched traces select the next + executed bottleneck. +4. Optimize the largest remaining open-axis contributor. Prefer a kernel or + engine change that improves TTFT/TPOT and throughput together. If a GEMM or + GEMV is selected, first prove output dtype, compute/scale type, entry point + and algorithm policy, and resolved template dtypes with the same tool on + both engines. +5. Measure peak VRAM independently under the identical model lifetime. Profile + allocation ownership before changing it; never trade host memory, load + correctness or reusable buffers merely to lower the reported peak. +6. Attempt BF16 vector writeback only after the TTFT/Pareto work above, and + only under its own compiler, correctness and rejection gates. + +## Correctness, tests, review and acceptance + +Before product code, a fresh implementer ports or restores the smallest tests +for each selector, eligibility predicate, launch mapping and fallback, records +their intended red result, then implements the minimum change. Preserve the +upstream shape/length/dtype coverage, local packed/strided/indexed/null-state +cases, exact output and persistent-state bytes, full CUDA GDN coverage and the +cached Qwen3.5-4B gate. Production token files must remain byte-identical. + +A fresh reviewer inspects an immutable head and mutates parser, eligibility, +mapping, arithmetic-order, fallback and callback-identity guarantees in a +scratch copy. Each material mutation must make the focused test fail; the +reviewed tree is restored byte-for-byte. Findings return to a fresh +implementer. The operator reruns focused gates and the full preflight itself. + +Performance uses one GPU lock, an idle box, the 22/25 GiB user-systemd memory +limits, one loaded model per arm where supported, paired order-alternated legs +and a calibrated noise band. A result is accepted only after same-binary +reproduction. Record revisions, binary/model hashes, exact commands, raw +artifacts, contention state, values and ratios. + +Hard no-regression gates relative to the table above are: + +- exact correctness remains satisfied; +- total and output throughput do not fall below the local baseline or vLLM; +- TTFT, TPOT/ITL and E2E move toward vLLM, with no accepted-axis regression; +- peak VRAM does not exceed 13053.3 MiB and must move toward 12820 MiB; +- material host PSS, load behavior and every supported fallback do not regress. + +## Stop, removal and default gates + +Any candidate that changes tokens, misses a fallback, loses both raw paired +timing legs, regresses an enclosing axis outside noise, increases memory, or +cannot be distinguished in compiler/trace output is rejected and its product, +selector and candidate-only tests are removed. Preserve the negative evidence +in the spec and benchmark record. Do not weaken a test or checker. + +A local win stays explicit opt-in until it reproduces, passes fresh mutation +review, and meets all local axes. No local 4B result grants a shared default or +27B/35B release claim; those remain pending their named hardware gate. Never +declare an unresolved gap a ceiling: record the next same-tool discriminator. diff --git a/.agents/specs/sm120-qwen35-postconv-token-tile-2026-08-08.md b/.agents/specs/sm120-qwen35-postconv-token-tile-2026-08-08.md new file mode 100644 index 000000000..885c24c80 --- /dev/null +++ b/.agents/specs/sm120-qwen35-postconv-token-tile-2026-08-08.md @@ -0,0 +1,170 @@ +# sm_120 Qwen3.5 fused post-conv token tile — structured spike + +**Rows:** `KERNEL-SSM-MAMBA`, feeding `ROAD-V1-C2-LOCAL-BF16`. +**Hardware/workload:** RTX 5070 Ti (`sm_120`), Qwen3.5-4B plain BF16, +128 ShareGPT requests, 128 output tokens, concurrency 32, +`max_num_batched_tokens=2048`, 1,280 KV blocks, greedy. **Lifecycle:** +IMPLEMENTED, byte-exact and locally faster, but opt-in pending repeated and +release-model gates. + +## Measured selection + +The exact-chunks branch at `c3bb0f39a` was profiled under one `/tmp/gpu` lock, +the 22/25 GiB user-systemd limits, and `--cuda-graph-trace=node`. The accepted +1,280-block workload reproduces the prior post-conv baseline and emits the same +token-file SHA-256 as the exact-chunks rebenchmark: + +| Arm | Calls | Total GPU time | Mean call | Total throughput | +|---|---:|---:|---:|---:| +| fast megablock | 1,728 | **228.150171 ms** | 132.031 us | 6,742.52 tok/s | +| per-V-head split | 1,728 | **448.364941 ms** | 259.471 us | 6,677.11 tok/s | + +Both token files hash to +`83fcdc45f79ddb06a634c7d7d95eba3384543b3cd781a45a8db1fc4e2a453545`. +The split is therefore falsified on this target: it makes the selected kernel +**1.965x slower** and the enclosing run 0.97% slower. The sealed pinned-vLLM +same-tool trace has 1,923 `_fused_post_conv_kernel` calls, 108.034870 ms total, +and 56.180 us mean. The selected metric is **total GPU time for the fused GDN +post-conv family on the exact c32 workload**; the current total gap is 2.112x +and the per-call gap is 2.350x. + +An earlier diagnostic accidentally omitted `--num-blocks 1280`, producing more +scheduler waves (2,376 calls). It independently found split 427.764060 ms versus +fast 215.064822 ms, but its absolute and enclosing numbers are VOID for parity. +It is retained only as a second falsification of the split hypothesis. + +## Whole-chain cause and upstream contract + +Pinned vLLM uses one Triton program over a **16-token tile and one Q/K or V +head**: grid `(ceil(L, 16), H + HV)`, four warps and two stages +(`${VLLM_SOURCE}/vllm/third_party/flash_linear_attention/ops/fused_gdn_prefill_post_conv.py:57-74,208-245`). +Q/K values remain in the program while each token's float32 square sum is +reduced and the normalized values are stored (`:76-107`). V copy and gating use +the same token tile (`:108-149`). The upstream executable specification sweeps +the 35B, 397B and small shapes, `L={1,16,128,512,2048}`, and BF16 correctness +(`${VLLM_SOURCE}/tests/kernels/test_fused_gdn_post_conv.py:60-141`), plus sanity +and L=0 (`:144-208`). + +Our fast kernel launches one `(token, Q/K-head)` block and one whole-value +megablock per token (`src/vt/cuda/cuda_gdn.cu:1271-1333,1434-1449`). Its Q and K +paths each use a 128-thread shared-memory reduction with eight block barriers and +then **reload** the activation for the normalized store. The rejected split only +changes the V grid; it leaves those one-token Q/K reductions intact +(`src/vt/cuda/cuda_gdn.cu:1335-1406`). The trace and source therefore select the +missing 16-token Q/K tile, register reuse and warp reductions—not another V-only +grid change—as the next discriminator. + +## Port and first implementation + +Add a CUDA token-tile kernel for the production `Dk==Dv==128` shape, behind +`VT_GDN_POSTCONV_TOKEN_TILE` and default OFF while it is evaluated: + +1. Launch `grid=(ceil(T,16), Hk+Hv)` with 128 threads (four warps), matching the + upstream work partition. +2. Assign each warp tokens `warp, warp+4, warp+8, warp+12` within the tile. + For a Q/K head, each lane owns features `lane+{0,32,64,96}`. Keep Q and K in + registers, reduce their float32 square sums with warp shuffles, then store + without rereading `conv` and without shared-memory barriers. +3. For a V head, use the same warp/token mapping; each lane copies four BF16/F32 + elements and lane zero computes the existing softplus/sigmoid gate. Preserve + the current local softplus arithmetic and tensor strides so this experiment + changes scheduling and data reuse, not the operation contract. +4. Keep `VT_GDN_POSTCONV_SPLIT` and `VT_GDN_POSTCONV_FAST` as independent + rollback controls. Unsupported dimensions remain on the current dispatch. + +This is a CUDA spelling of vLLM's Triton schedule, not an unrelated sm_120 +algorithm. It deliberately does not claim bit identity in advance: the warp +reduction groups the same float32 terms differently from the current 128-lane +tree. The upstream tolerance contract must pass, and the project-level cached +model tokens must remain exact before any default flip. + +No GEMM claim is made. If a later profile selects a GEMM, it separately owes the +four-axis same-tool invocation proof (C/output dtype, compute/scale type, entry +point/algo policy, and resolved template dtype). + +## Tests and acceptance + +RED-first coverage must add: + +1. a portable flag-predicate test proving the tile is opt-in and `0` rolls back; +2. the upstream shape/length sweep for BF16, including the partial last tile, + `T=0`, `T=1`, and strided packed-BA gate inputs; +3. CUDA tile versus current fast output checks: V/g/beta exact, Q/K within the + upstream `1e-2` BF16 tolerance, finite outputs and unit norms; +4. a structural mutant that restores the one-token grid fails the tile launch + contract test; +5. full `test_ops_gdn` and cached Qwen3.5-4B 3/3·1672, followed by production + tile-OFF/ON token-file identity. + +The first performance decision is the same-binary graph-node micro-metric. +The tile must improve 228.150171 ms outside run noise and move toward the sealed +108.034870 ms total. An enclosing profile must not regress total/output +throughput, TTFT, TPOT/ITL, E2E or peak VRAM. Default ON additionally requires +token-exactness, repeated A/B evidence, and the unavailable 27B/35B gate-model +gates before claiming shared release coverage; the local 4B result is not +extrapolated. + +## Evidence and rollback + +- fast trace: `/tmp/qwen35-postconv-split-c3bb-fast-nb1280.nsys-rep`, SHA-256 + `9e430b659dd30950436b7cde57dc85248269febb059d95341b7782c8e1fa5e0a` +- split trace: `/tmp/qwen35-postconv-split-c3bb-split-nb1280.nsys-rep`, SHA-256 + `855665ea3b042f28022e5d5ccdc32a4a85ccac04d6d153dad31ff12428e1823a` +- pinned-vLLM trace: `/tmp/qwen35-async-3f35356e0-vllm.nsys-rep` +- diagnostic no-`num-blocks` traces: + `/tmp/qwen35-postconv-split-c3bb-{fast,split}.nsys-rep` + +Rollback is `VT_GDN_POSTCONV_TOKEN_TILE=0` (or unset while the experiment is +opt-in), which retains the current default fast megablock without changing any +loader, scheduler or model route. + +## Implementation outcome + +`GdnPostConvTokenTileKernel` now implements the specified four-warp, +16-token/per-head work partition for `Dk==Dv==128`. Q/K activations remain in +registers through the reduction and normalized store; V/g/beta share the same +tile. The dispatch is opt-in through `VT_GDN_POSTCONV_TOKEN_TILE=1`, the slower +explicit split retains priority, and every unsupported shape stays on the +existing fast megablock. + +The first warp reduction was a useful negative result. Sequentially summing the +four lane-owned squares before the 32-lane shuffle retained the kernel speedup +but changed production tokens: fast/tile token SHA-256 +`83fcdc45...453545`/`1d496ff0...b9756`. The accepted implementation reproduces +the existing 128-lane tree exactly: first `(i+i+64)`, then the two resident +partials `(i+i+32)`, followed by shuffle offsets `16,8,4,2,1`. This restores +byte identity while keeping the values in registers. + +Final same-binary graph-node traces have 1,728 calls per arm: + +| Arm | Total GPU time | Mean call | Total throughput | TTFT | TPOT / ITL | E2E | +|---|---:|---:|---:|---:|---:|---:| +| fast megablock | 227.887066 ms | 131.879 us | 6,734.82 tok/s | 1,024.14 ms | 35.01 ms | 5,469.87 ms | +| token tile | **122.587027 ms** | **70.942 us** | **6,770.62 tok/s** | **1,015.43 ms** | **34.85 ms** | **5,440.81 ms** | + +The tile is **1.859x faster** at the selected kernel and improves the enclosing +profile on every observed axis: total/output throughput +0.532%, TTFT -0.850%, +TPOT -0.457%, E2E -0.531%. It closes the same-tool vLLM gap from 2.112x to +**1.135x** (122.587027/108.034870 ms). Both final token files have SHA-256 +`83fcdc45f79ddb06a634c7d7d95eba3384543b3cd781a45a8db1fc4e2a453545`. + +Tests are stronger than the upstream BF16 tolerance: portable flag/grid 6/6, +CUDA GDN 67/67 and 4,384 assertions over partial/exact tiles, production and +small shapes, packed BA views with non-zero offsets/wider row strides, exact +q/k/v/g/beta bytes, finiteness and unit norms; cached Qwen3.5-4B 3/3 and 1,672 +assertions. `T=0` is pinned by the portable no-work grid contract because local +kernel tensor descriptors require positive dimensions; the runtime sweep covers +`T={1,16,17,128,512,2048}`. + +Final evidence: + +- fast trace `/tmp/qwen35-postconv-tile-exact-fast.nsys-rep`, SHA-256 + `c75e2cb27797827b3d25d40204d745b3dfa36c4be8bd17a8464229c1b70bcadc`; +- tile trace `/tmp/qwen35-postconv-tile-exact-tile.nsys-rep`, SHA-256 + `a0eb1808e216a39d1350deabe7722e1b9081940183c512cefdb10afd2706f418`; +- rejected arithmetic traces `/tmp/qwen35-postconv-tile-wip-{fast,tile}.nsys-rep`, + SHA-256 `c2d8872e...b774` / `1c42d489...d66b`. + +Default ON remains deliberately unclaimed. One local 4B profile does not close +the required repeated A/B or the unavailable Qwen3.6-27B/35B release-model +correctness/performance gates. diff --git a/docs/BENCHMARKS.md b/docs/BENCHMARKS.md index 2162fea61..54fb89ff1 100644 --- a/docs/BENCHMARKS.md +++ b/docs/BENCHMARKS.md @@ -41,7 +41,7 @@ The binding comparison. vLLM runs its **production graphed config**, never | Qwen3.6-27B | NVFP4 (`nvidia` @`0893e160`, ModelOpt `modelopt_mixed`) | 0.25.0 | 0/4 | **BEHIND, uniformly 0.85x** on decode throughput (was 0.72x before the FP8 tower fix); greedy continuation IDENTICAL to vLLM. A different model from the `unsloth` row (NVFP4 MLP + FP8 W8A8 GDN/attn tower) | | Qwen3.6-35B-A3B | NVFP4 `modelopt_mixed` | 0.25.0 | 2/18 | 3-rep grid 2026-08-05 @`1ea26427`: 0.93-1.03x (c4 wins), c16 0.93x. Both c16 levers A/B'd NEG: drain event -1.9%, mirror 0.999x. ★ probe found a prod async batch-1 greedy DEGENERATION bug the mirror fixes | | DeepSeek-V2-Lite | bf16 MLA | 0.25.0 | 4/25 | Attributed miss, row stays `ACTIVE` | -| Qwen3.5-4B | bf16 direct-load | 0.26.0.dev0 | throughput + host PSS | Exact chunks ON: total **1.021x PASS**; TTFT **1.086x**, TPOT **1.025x**, VRAM **1.018x OPEN**; local A/B **+2.152%** ([evidence](bench-evidence/qwen35-4b-sm120-main-20260807.md)) | +| Qwen3.5-4B | bf16 direct-load | 0.26.0.dev0 | throughput + host PSS | Exact chunks: total **1.021x PASS**; TTFT **1.086x**, TPOT **1.025x**, VRAM **1.018x OPEN**. #206 imports history only; no new accepted result ([evidence](bench-evidence/qwen35-4b-sm120-main-20260807.md)) | ### GDN prefill causal-convolution by GPU diff --git a/docs/STATUS.md b/docs/STATUS.md index 71a06b08e..e056a5df8 100644 --- a/docs/STATUS.md +++ b/docs/STATUS.md @@ -78,7 +78,7 @@ token-for-token correctness against the pinned oracle. | Qwen3.6-27B (NVFP4) text generation | Correctness-complete; speed is CHECKPOINT-dependent | Token-exact GB10 on both. `unsloth` @`890bdef7` beats vLLM 0.25.0 every c (1.007-1.045x), 115/124; `nvidia` @`0893e160` (ModelOpt FP8 tower) is **0.85x BEHIND**, decode ~100% GPU-busy | | Qwen3.6-35B-A3B (NVFP4, GDN MoE) | Correctness-complete; binding grid @`a0fa12c7` FLAT 0.935x-0.979x over c1-c32 (CoV <0.81%); the prior 0.87x c2 / 0.92x c8 "weak cells" were harness mismatch, not code; memory PSS 3.81x, GPU 1.40x | Token-exact SYNC+ASYNC; `VT_ASYNC_DEVICE_MIRROR` ON fixes async batch-1 token-0 degeneration; `VT_ASYNC_EXECUTOR` Option A NEUTRAL → OFF | | Qwen3 / Qwen2 dense (BF16) | Correctness-complete, speed-pending. Async-serving P0 FIXED (`ROW-SERVE-ASYNC-DENSE-MIRROR`): classic-dense `Qwen3ForCausalLM` now honors the async device token-ids mirror; CPU-only -Werror test-guard fixes x2 | Near-tie-robust token-exact vs vLLM (Qwen3-0.6B, Qwen3-4B); c1 effective parity, c8 decode residual. **Async device-mirror (`ROW-SERVE-ASYNC-DENSE-MIRROR`, `f9c969ae`): the #31 fix ported to the classic dense family, dgx-VERIFIED.** The shared dense `EmbedInto` (qwen3.cpp) raced the async combine's device input-ids write against a stale host upload → token-0 degeneration on the depth-2 AsyncLLM serving path (quant-independent). `EmbedInto` now consumes the device override published by `ForwardQwen3ForCausalLM`'s `DeviceTokenIdsScope` (27B-dense template); gate `test_qwen3_dense_async_serving` RED on `VT_ASYNC_DEVICE_MIRROR=0`, GREEN default, byte-identical mirror-off. dgx GB10: async gate RED→GREEN 0.6B+4B, SACRED 0.6B+4B 184/184 unchanged (byte-neutral sync path), memcheck 0 errors; Yi30/Qwen3-8B-MXFP4 default-config e2e coherent + 3/4 token-exact (p2 = oracle-ratified near-tie, gap 0.0000), closing the QUANT-CT-MXFP4 async-default residual. RESIDUAL: sibling InternLM2/Mistral/Llama scope one-liner; W4 bench RAN; FA2 GQA-swap default-ON, c2-c8 <1.0x. `FLASH-PTXAS` #82: codegen at PARITY (no ptxas lever); gap=engine context. **D1 (2026-07-31, `CLAIM-D1-BF16-MERGED-QKV`): the bf16 merged-QKV path (`Qwen3QkvMergeEnabled`/`VT_QWEN3_QKV_MERGE`) is now default-ON** — one `vt::MatmulBT` over the merged `[qdim+2kdim,H]` owner + a contiguous `vt::QkvSplit` (OLMo-2 exemplar), replacing three per-shard GEMMs. Bit-exact GEMM math (A/B unit `test_ops_qkv_merge` byte-identical, RED-first); the wider-N cuBLASLt K-reduction flips the 0.6B genuine bf16 near-tie so the SACRED 0.6B golden was regenerated (all tokens within the near-tie band, max 0.125 nats), while Qwen3-4B is byte-neutral (0 diffs, stays STRICT). Re-gated 0.6B 16/16 + 4B 16/16; consistency/launch-count fold (measured NEUTRAL on 4B decode), no new throughput owed | -| Qwen3.5-4B plain BF16 direct loading on discrete CUDA | Correctness-complete; throughput passes, latency/VRAM open | Exact GDN chunks default ON and byte-identical to rollback. Local A/B: total/output +2.152%, TTFT -2.945%, TPOT/ITL -1.920%; sealed-vLLM comparison 1.021x throughput, 1.086x TTFT, 1.025x TPOT, +233 MiB VRAM ([evidence](bench-evidence/qwen35-4b-sm120-main-20260807.md)) | +| Qwen3.5-4B plain BF16 direct loading on discrete CUDA | Correctness-complete; throughput passes, latency/VRAM open | Exact chunks default ON and byte-identical. vLLM ratios: throughput 1.021x PASS; TTFT 1.086x, TPOT 1.025x, VRAM +233 MiB OPEN. Issue #206 imports prior history only; no new accepted result ([evidence](bench-evidence/qwen35-4b-sm120-main-20260807.md)) | | Qwen3-Coder-30B-A3B MoE (BF16) | Correctness-complete, speed-pending | Near-tie-robust token-exact 6/6; 11 of 16 binding grid cells at or above vLLM. **D1 (2026-07-31): inherits the default-ON bf16 merged-QKV via the shared dense `AttnBlock` — byte-neutral (0 token diffs, golden UNCHANGED); re-gated 6/6** | | Llama-3.x dense (BF16) | Correctness-complete, speed-pending | Near-tie-robust token-exact 16/16 (Llama-3.2-1B); llama3 RoPE scaling | | Mistral dense (BF16) | Correctness-complete, speed-pending | Paged-engine token-exact 16/16 (Mistral-7B-v0.3) | From c7be73205d66ac63c50910deb69fb40c6bf11f4f Mon Sep 17 00:00:00 2001 From: Richard Palethorpe Date: Sun, 9 Aug 2026 18:50:25 +0100 Subject: [PATCH 02/17] perf(cuda): restore sm120 GDN experiment arms (#206) FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: Codex:gpt-5.6-sol [Codex] --- .agents/NOW.md | 4 +- .agents/kernel-matrix.md | 7 +- docs/BENCHMARKS.md | 2 +- docs/ENVIRONMENT.md | 5 + docs/STATUS.md | 3 +- src/vt/cuda/cuda_gdn.cu | 520 +++++++++++++++++++++++++---- src/vt/cuda/gdn_decode_fused.h | 171 ++++++++++ src/vt/cuda/gdn_prefill_conv.h | 90 +++++ tests/CMakeLists.txt | 5 + tests/vt/test_gdn_decode_fused.cpp | 281 ++++++++++++++++ tests/vt/test_gdn_prefill_conv.cpp | 113 ++++++- tests/vt/test_ops_gdn.cpp | 452 ++++++++++++++++++++++++- 12 files changed, 1562 insertions(+), 91 deletions(-) create mode 100644 src/vt/cuda/gdn_decode_fused.h create mode 100644 tests/vt/test_gdn_decode_fused.cpp diff --git a/.agents/NOW.md b/.agents/NOW.md index 4b13b45ee..6f7a23596 100644 --- a/.agents/NOW.md +++ b/.agents/NOW.md @@ -21,7 +21,7 @@ Work: exact-chunks on main `1ce0d662b`; sm_120 measured at `3d2581551`. | MiniMax-H3 | **PRUNED ckpts RUN (#241): Q8_0 renders, seam 0.9941** | same-binary A/B | | Kimi-Linear-48B | 122/128 held; grouped router parallelised, e2e NOT ESTABLISHED | ckpt is tiktoken-only: no warm server | | 35B binding grid | @`a0fa12c7`: **flat 0.935-0.979x, NO c2/c8 weak cell** (CoV <0.81%); mem PSS 3.81x | Attribute the flat ~5% mid-band; TTFT c2 0.872x | -| Qwen3.5-4B sm_120 | #206 spec; tput 1.021x PASS; TTFT/TPOT/VRAM OPEN | Transplant opt-ins; combined A/B | +| Qwen3.5-4B sm_120 | #206 opt-ins restored; tput 1.021x PASS; TTFT/TPOT/VRAM OPEN | Review; combined A/B | | RPi5 A76 CPU | **R5 asm GREEN; llama NOT MET**: 0.461x pf, 0.653x dec | W6: BF16 GEMM | | MXFP4 parity | c1 1.020, c2-c8 0.962-0.969. **#82 CLOSED: ptxas-lineage REFUTED** | TERMINAL: at parity | | SERVE-ASYNC-DENSE-MIRROR | **LANDED+VERIFIED** (`f9c969ae`): async mirror, dense Qwen3; SACRED 184/184 | Sibling scope one-liner | @@ -53,7 +53,7 @@ latency/memory on every axis, both gate models, reproduced 2–3x idle. See of roof. Dense-marlin +0.5%; Triton-AOT GDN a WASH. 2. **Spike the Parakeet encoder row** (vLLM: `nano_nemotron_vl.py`; the transducer half is NOT in vLLM: separate call). -3. **Qwen3.5-4B #206:** transplant opt-ins; combined K4+postconv A/B, then trace. +3. **Qwen3.5-4B #206:** review restored opt-ins; combined K4+postconv A/B. 2. **Merge the invocation-parity prevention** (CI guard + AGENTS.md checklist); CUDA build-verify the byte-exact `kGemvHeuristicAlgos` refactor on dgx. 4. **Restore `local-ai-worker`** on dgx at campaign end (`--restart=always`). diff --git a/.agents/kernel-matrix.md b/.agents/kernel-matrix.md index 55b226ef3..b16876e93 100644 --- a/.agents/kernel-matrix.md +++ b/.agents/kernel-matrix.md @@ -177,8 +177,11 @@ release gates are unchanged. [Spec and evidence](specs/sm120-qwen35-conv-chunkin Qwen3.5-4B Pareto campaign. The current sealed result passes throughput (**6784.743 vs 6643.593 tok/s**) while TTFT (**1018.040 vs 937.584 ms**), TPOT/ITL (**34.740 vs 33.906 ms**) and VRAM (**13053.3 vs 12820 MiB**) remain -open. The first ordered gate transplants the historical opt-ins and measures -combined K=4 causal-conv plus 16-token post-conv in a counterbalanced A/B; +open. The historical default-off K=4 causal-conv, 16-token post-conv and +BV16/swizzle/register-state decode arms and their focused contracts are now +restored on the current branch; fresh performance measurement remains pending. +The first ordered gate measures combined K=4 causal-conv plus 16-token +post-conv in a counterbalanced A/B; same-tool local/vLLM tracing then selects the largest residual. Lifecycle stays `INVENTORIED`; generic Mamba and 27B/35B coverage are unchanged. [Campaign spec](specs/sm120-qwen35-pareto-2026-08-09.md). diff --git a/docs/BENCHMARKS.md b/docs/BENCHMARKS.md index 54fb89ff1..7ca38eb92 100644 --- a/docs/BENCHMARKS.md +++ b/docs/BENCHMARKS.md @@ -41,7 +41,7 @@ The binding comparison. vLLM runs its **production graphed config**, never | Qwen3.6-27B | NVFP4 (`nvidia` @`0893e160`, ModelOpt `modelopt_mixed`) | 0.25.0 | 0/4 | **BEHIND, uniformly 0.85x** on decode throughput (was 0.72x before the FP8 tower fix); greedy continuation IDENTICAL to vLLM. A different model from the `unsloth` row (NVFP4 MLP + FP8 W8A8 GDN/attn tower) | | Qwen3.6-35B-A3B | NVFP4 `modelopt_mixed` | 0.25.0 | 2/18 | 3-rep grid 2026-08-05 @`1ea26427`: 0.93-1.03x (c4 wins), c16 0.93x. Both c16 levers A/B'd NEG: drain event -1.9%, mirror 0.999x. ★ probe found a prod async batch-1 greedy DEGENERATION bug the mirror fixes | | DeepSeek-V2-Lite | bf16 MLA | 0.25.0 | 4/25 | Attributed miss, row stays `ACTIVE` | -| Qwen3.5-4B | bf16 direct-load | 0.26.0.dev0 | throughput + host PSS | Exact chunks: total **1.021x PASS**; TTFT **1.086x**, TPOT **1.025x**, VRAM **1.018x OPEN**. #206 imports history only; no new accepted result ([evidence](bench-evidence/qwen35-4b-sm120-main-20260807.md)) | +| Qwen3.5-4B | bf16 direct-load | 0.26.0.dev0 | throughput + host PSS | Exact chunks: throughput **1.021x PASS**; TTFT **1.086x**, TPOT **1.025x**, VRAM **1.018x OPEN**. #206 opt-ins restored; combined A/B pending ([data](bench-evidence/qwen35-4b-sm120-main-20260807.md)) | ### GDN prefill causal-convolution by GPU diff --git a/docs/ENVIRONMENT.md b/docs/ENVIRONMENT.md index 0c58b01bf..6068401f5 100644 --- a/docs/ENVIRONMENT.md +++ b/docs/ENVIRONMENT.md @@ -87,8 +87,13 @@ portable/reference path. In normal operation leave them unset. | `VT_DEVICE_KV_CACHE` | on (CUDA) | Host-side KV cache instead of the on-device one | | `VT_GPU_SAMPLE` | on (CUDA) | Host-side sampling instead of on-GPU sampling | | `VT_GDN_PACKED_DECODE` | on (CUDA GDN) | Unpacked GDN decode path | +| `VT_GDN_DECODE_BV` | `32` (CUDA GDN decode experiment) | Exact `16` selects the byte-identical 16-value fused-recurrence tile; unset and every other spelling keep the 32-value schedule. Experimental opt-in; no release or cross-hardware default change | +| `VT_GDN_DECODE_SWIZZLE` | `0` (CUDA GDN decode experiment) | Exact `1` enables the shared-memory bank swizzle only for the `BV=16`, `Dv=Dk=128`, eight-lane production geometry; all other values and shapes keep the incumbent layout | +| `VT_GDN_DECODE_REGSTATE` | `0` (CUDA GDN decode experiment) | Exact `1` retains each lane's recurrence row in registers, only when the `BV=16` shared swizzle is eligible. Other geometries and spellings use the shared-state implementation | | `VT_CONV_REG` | on (CUDA GDN) | The non-register-tiled short causal convolution | | `VT_CONV_EXACT_CHUNKS` | on (CUDA GDN prefill) | Use `=0` for the legacy sequence-serial causal-conv mapping; default mirrors vLLM's exact `(sequence, 8-token chunk)` descriptor and is byte-identical | +| `VT_CONV_CHANNEL_TILE` | `0` (CUDA GDN prefill experiment) | Exact `1` specializes the production `K=4` causal convolution with one channel per thread; exact `2` is the retained slower two-channel control. Unset, invalid, and non-`K=4` shapes use the runtime-width schedule | +| `VT_GDN_POSTCONV_TOKEN_TILE` | `0` (CUDA GDN prefill experiment) | A non-`0` value selects the byte-identical 16-token, per-head fused post-conv schedule. It remains opt-in pending the combined K4+post-conv Pareto measurement in issue #206 | | `VT_MODELOPT_W4A4` | `0` (Qwen3.6 dense ModelOpt NVFP4) | ModelOpt NVFP4 checkpoints ship a per-tensor `input_scale` next to every projection. Consuming it sets `Nvfp4Weight::alpha`, which flips `IsTrueW4A4()` and routes the weight to the fp4-ACTIVATION GEMM; on `nvidia/Qwen3.6-27B-NVFP4` that produced incoherent text, so the default leaves `alpha` at 0 and takes the W4A16 weight-only dispatcher (verified coherent). Set `1` to consume `input_scale` and take the W4A4 path | | `VT_FA2_PREFILL` | on (CUDA) | The portable prefill attention instead of the vendored FA2 | | `VT_FA2_DECODE` | on (CUDA) | The portable decode attention instead of the vendored FA2 | diff --git a/docs/STATUS.md b/docs/STATUS.md index e056a5df8..479774aac 100644 --- a/docs/STATUS.md +++ b/docs/STATUS.md @@ -78,7 +78,7 @@ token-for-token correctness against the pinned oracle. | Qwen3.6-27B (NVFP4) text generation | Correctness-complete; speed is CHECKPOINT-dependent | Token-exact GB10 on both. `unsloth` @`890bdef7` beats vLLM 0.25.0 every c (1.007-1.045x), 115/124; `nvidia` @`0893e160` (ModelOpt FP8 tower) is **0.85x BEHIND**, decode ~100% GPU-busy | | Qwen3.6-35B-A3B (NVFP4, GDN MoE) | Correctness-complete; binding grid @`a0fa12c7` FLAT 0.935x-0.979x over c1-c32 (CoV <0.81%); the prior 0.87x c2 / 0.92x c8 "weak cells" were harness mismatch, not code; memory PSS 3.81x, GPU 1.40x | Token-exact SYNC+ASYNC; `VT_ASYNC_DEVICE_MIRROR` ON fixes async batch-1 token-0 degeneration; `VT_ASYNC_EXECUTOR` Option A NEUTRAL → OFF | | Qwen3 / Qwen2 dense (BF16) | Correctness-complete, speed-pending. Async-serving P0 FIXED (`ROW-SERVE-ASYNC-DENSE-MIRROR`): classic-dense `Qwen3ForCausalLM` now honors the async device token-ids mirror; CPU-only -Werror test-guard fixes x2 | Near-tie-robust token-exact vs vLLM (Qwen3-0.6B, Qwen3-4B); c1 effective parity, c8 decode residual. **Async device-mirror (`ROW-SERVE-ASYNC-DENSE-MIRROR`, `f9c969ae`): the #31 fix ported to the classic dense family, dgx-VERIFIED.** The shared dense `EmbedInto` (qwen3.cpp) raced the async combine's device input-ids write against a stale host upload → token-0 degeneration on the depth-2 AsyncLLM serving path (quant-independent). `EmbedInto` now consumes the device override published by `ForwardQwen3ForCausalLM`'s `DeviceTokenIdsScope` (27B-dense template); gate `test_qwen3_dense_async_serving` RED on `VT_ASYNC_DEVICE_MIRROR=0`, GREEN default, byte-identical mirror-off. dgx GB10: async gate RED→GREEN 0.6B+4B, SACRED 0.6B+4B 184/184 unchanged (byte-neutral sync path), memcheck 0 errors; Yi30/Qwen3-8B-MXFP4 default-config e2e coherent + 3/4 token-exact (p2 = oracle-ratified near-tie, gap 0.0000), closing the QUANT-CT-MXFP4 async-default residual. RESIDUAL: sibling InternLM2/Mistral/Llama scope one-liner; W4 bench RAN; FA2 GQA-swap default-ON, c2-c8 <1.0x. `FLASH-PTXAS` #82: codegen at PARITY (no ptxas lever); gap=engine context. **D1 (2026-07-31, `CLAIM-D1-BF16-MERGED-QKV`): the bf16 merged-QKV path (`Qwen3QkvMergeEnabled`/`VT_QWEN3_QKV_MERGE`) is now default-ON** — one `vt::MatmulBT` over the merged `[qdim+2kdim,H]` owner + a contiguous `vt::QkvSplit` (OLMo-2 exemplar), replacing three per-shard GEMMs. Bit-exact GEMM math (A/B unit `test_ops_qkv_merge` byte-identical, RED-first); the wider-N cuBLASLt K-reduction flips the 0.6B genuine bf16 near-tie so the SACRED 0.6B golden was regenerated (all tokens within the near-tie band, max 0.125 nats), while Qwen3-4B is byte-neutral (0 diffs, stays STRICT). Re-gated 0.6B 16/16 + 4B 16/16; consistency/launch-count fold (measured NEUTRAL on 4B decode), no new throughput owed | -| Qwen3.5-4B plain BF16 direct loading on discrete CUDA | Correctness-complete; throughput passes, latency/VRAM open | Exact chunks default ON and byte-identical. vLLM ratios: throughput 1.021x PASS; TTFT 1.086x, TPOT 1.025x, VRAM +233 MiB OPEN. Issue #206 imports prior history only; no new accepted result ([evidence](bench-evidence/qwen35-4b-sm120-main-20260807.md)) | +| Qwen3.5-4B plain BF16 direct loading on discrete CUDA | Correctness-complete; throughput passes, latency/VRAM open | Exact chunks are byte-identical. vLLM: throughput 1.021x PASS; TTFT 1.086x, TPOT 1.025x, VRAM +233 MiB OPEN. #206 opt-ins restored; combined A/B pending ([data](bench-evidence/qwen35-4b-sm120-main-20260807.md)) | | Qwen3-Coder-30B-A3B MoE (BF16) | Correctness-complete, speed-pending | Near-tie-robust token-exact 6/6; 11 of 16 binding grid cells at or above vLLM. **D1 (2026-07-31): inherits the default-ON bf16 merged-QKV via the shared dense `AttnBlock` — byte-neutral (0 token diffs, golden UNCHANGED); re-gated 6/6** | | Llama-3.x dense (BF16) | Correctness-complete, speed-pending | Near-tie-robust token-exact 16/16 (Llama-3.2-1B); llama3 RoPE scaling | | Mistral dense (BF16) | Correctness-complete, speed-pending | Paged-engine token-exact 16/16 (Mistral-7B-v0.3) | @@ -1304,6 +1304,7 @@ the sealed same-workload vLLM baseline, throughput is **1.021246x PASS**; TTFT is **1.085812x OPEN**, TPOT/ITL **1.024597x OPEN**, and mean peak VRAM 13053.3/12820 MiB OPEN. Fresh 18-leg oracle attempts were VOID JIT-environment runs and do not replace the sealed denominator. +The restored selectors remain default OFF. This local 4B diagnostic does not establish 27B/35B support. Exact evidence and reproduction: diff --git a/src/vt/cuda/cuda_gdn.cu b/src/vt/cuda/cuda_gdn.cu index ae4d83f67..fa65ba3e3 100644 --- a/src/vt/cuda/cuda_gdn.cu +++ b/src/vt/cuda/cuda_gdn.cu @@ -35,6 +35,7 @@ #include "vt/cuda/conv_update_fast.h" #include "vt/cuda/cuda_device_caps.h" #include "vt/cuda/cuda_gdn_internal.h" +#include "vt/cuda/gdn_decode_fused.h" #include "vt/cuda/gdn_packed_decode_triton.h" #include "vt/cuda/gdn_prefill_conv.h" #include "vt/cuda/gdn_packed_reg_tile.h" @@ -842,6 +843,114 @@ __global__ void CausalConv1dFwdRegKernel(Tout* out, const Tin* x, const Tin* w, } } +// Width-four experiment used by Qwen3.5. Unlike the sealed runtime-width kernel +// above, K and the number of channels owned by each lane are compile-time +// constants. ChannelsPerThread=1 isolates width specialization at the same +// 128-channel tile; ChannelsPerThread=2 covers two coalesced 128-channel stripes +// and therefore matches upstream's 256-channel feature tile. +template +__global__ void CausalConv1dFwdRegK4Kernel( + Tout* out, const Tin* x, const Tin* w, const Tin* bias, float* conv_state, + const int32_t* qsl, const THas* his, int64_t c_dim, int64_t x_row_stride, + bool silu, int chunked, const int32_t* batch_ptr, + const int32_t* token_chunk_offset_ptr, int exact) { + static_assert(ChannelsPerThread == 1 || ChannelsPerThread == 2); + constexpr int64_t k = 4; + constexpr int64_t width = k - 1; + constexpr int64_t channels_per_block = kConvRegN * ChannelsPerThread; + + const int64_t program = blockIdx.y; + const int64_t s = exact ? batch_ptr[program] : program; + const int64_t begin = qsl[s]; + const int64_t t_len = qsl[s + 1] - begin; + const int64_t chunk_m = exact ? kConvExactM : kConvRegM; + const int64_t token_offset = exact + ? static_cast(token_chunk_offset_ptr[program]) * + kConvExactM + : (chunked ? static_cast(blockIdx.z) * kConvRegM + : 0); + if (token_offset > 0 && token_offset >= t_len) return; + const int64_t token_end = + ((chunked || exact) && token_offset + chunk_m < t_len) ? token_offset + chunk_m : t_len; + const bool init = his[s] != 0; + + int64_t channels[ChannelsPerThread]; + bool active[ChannelsPerThread]; + float* state_rows[ChannelsPerThread]; + float biases[ChannelsPerThread]; + float weights[ChannelsPerThread][k]; + float windows[ChannelsPerThread][k]; + + // Load both stripes' initial history before processing either stripe. This + // matches the intended duplicated-register experiment without extending the + // baseline state-read/write exposure across an entire first-stripe token loop. +#pragma unroll + for (int lane_channel = 0; lane_channel < ChannelsPerThread; ++lane_channel) { + const int64_t c = static_cast(blockIdx.x) * channels_per_block + + threadIdx.x + static_cast(lane_channel) * kConvRegN; + channels[lane_channel] = c; + active[lane_channel] = c < c_dim; + if (!active[lane_channel]) continue; + float* srow = conv_state + (s * c_dim + c) * width; + state_rows[lane_channel] = srow; + biases[lane_channel] = bias != nullptr ? Load(bias, c) : 0.0f; +#pragma unroll + for (int j = 0; j < k; ++j) weights[lane_channel][j] = Load(w, c * k + j); +#pragma unroll + for (int j = 0; j < width; ++j) { + const int64_t ti = token_offset - width + j; + float v = 0.0f; + if (ti >= 0) { + v = Load(x, (begin + ti) * x_row_stride + c); + } else if (init) { + v = srow[width + ti]; + } + windows[lane_channel][j] = v; + } + windows[lane_channel][width] = + token_offset < t_len ? Load(x, (begin + token_offset) * x_row_stride + c) : 0.0f; + } + + for (int64_t t = token_offset; t < token_end; ++t) { +#pragma unroll + for (int lane_channel = 0; lane_channel < ChannelsPerThread; ++lane_channel) { + if (!active[lane_channel]) continue; + const int64_t c = channels[lane_channel]; + float acc = biases[lane_channel]; +#pragma unroll + for (int j = 0; j < k; ++j) + acc += weights[lane_channel][j] * windows[lane_channel][j]; + Store(out, (begin + t) * c_dim + c, silu ? Silu(acc) : acc); +#pragma unroll + for (int j = 0; j < width; ++j) + windows[lane_channel][j] = windows[lane_channel][j + 1]; + const int64_t nt = t + 1; + windows[lane_channel][width] = + nt < t_len ? Load(x, (begin + nt) * x_row_stride + c) : 0.0f; + } + } + + if (token_end == t_len) { +#pragma unroll + for (int lane_channel = 0; lane_channel < ChannelsPerThread; ++lane_channel) { + if (!active[lane_channel]) continue; + const int64_t c = channels[lane_channel]; + float* srow = state_rows[lane_channel]; +#pragma unroll + for (int j = 0; j < width; ++j) { + const int64_t tj = t_len - width + j; + float v = 0.0f; + if (tj >= 0) { + v = Load(x, (begin + tj) * x_row_stride + c); + } else if (init) { + v = srow[width + tj]; + } + srow[j] = v; + } + } + } +} + // Toggle: DEFAULT ON (VT_CONV_REG=0 restores the tiled/scalar path). Read per call // (prefill dispatch is coarse — one launch/step — so the getenv is negligible and // in-process CUDA tests can flip the selection). Predicate factored to @@ -854,6 +963,56 @@ bool ConvExactChunksEnabled() { return ConvExactChunksFlagIsOn(std::getenv("VT_CONV_EXACT_CHUNKS")); } +template +void LaunchConvFwdRegK4(cudaStream_t s, Tensor& out, const Tensor& x, + const Tensor& w, const Tensor* bias, Tensor& conv_state, + const Tensor& qsl, const Tensor& his, + const CausalConv1dArgs& args, int64_t chan_tiles) { + static_assert(ChannelsPerThread == 1 || ChannelsPerThread == 2); + const int64_t n = conv_state.shape[0], c = x.shape[1]; + const int64_t total_tokens = x.shape[0]; + const int64_t x_rs = x.stride[0]; + int64_t grid_z = 1; + int chunked = 0; + const bool exact = ConvExactChunksEnabled() && args.batch_ptr != nullptr; + int64_t grid_y = n; + if (exact) { + grid_y = args.batch_ptr->shape[0]; + VT_CHECK(grid_y <= kMaxGridY, + "cuda causal_conv1d_fwd(reg-k4): too many exact chunk programs"); + } else if (n <= kConvRegChunkMaxSeqs) { + const int64_t z = (total_tokens + kConvRegM - 1) / kConvRegM; + if (z >= 1 && z <= kMaxGridY) { + grid_z = z; + chunked = 1; + } + } + const dim3 grid(static_cast(chan_tiles), static_cast(grid_y), + static_cast(grid_z)); + const dim3 block(kConvRegN); + const int32_t* batch_ptr = exact ? args.batch_ptr->Ptr() : nullptr; + const int32_t* token_chunk_offset_ptr = + exact ? args.token_chunk_offset_ptr->Ptr() : nullptr; + if (his.dtype == DType::kI8) { + CausalConv1dFwdRegK4Kernel + <<>>( + out.Ptr(), x.Ptr(), w.Ptr(), + bias != nullptr ? bias->Ptr() : nullptr, conv_state.Ptr(), + qsl.Ptr(), his.Ptr(), c, x_rs, + args.silu_activation, chunked, batch_ptr, token_chunk_offset_ptr, + exact ? 1 : 0); + } else { + CausalConv1dFwdRegK4Kernel + <<>>( + out.Ptr(), x.Ptr(), w.Ptr(), + bias != nullptr ? bias->Ptr() : nullptr, conv_state.Ptr(), + qsl.Ptr(), his.Ptr(), c, x_rs, + args.silu_activation, chunked, batch_ptr, token_chunk_offset_ptr, + exact ? 1 : 0); + } + Check(cudaGetLastError(), "causal_conv1d_fwd(reg-k4) launch"); +} + // Register-window launcher (VT_CONV_REG=1). The default exact descriptor maps // grid.y to a flattened list of (sequence, 8-token chunk) programs, mirroring // upstream and launching neither rectangular padding nor sequence-serial work. @@ -861,14 +1020,15 @@ bool ConvExactChunksEnabled() { // chunks): it chunks grid.z only for <=4 sequences, and serially streams each // whole sequence for larger batches. template -void LaunchConvFwdReg(cudaStream_t s, Tensor& out, const Tensor& x, const Tensor& w, - const Tensor* bias, Tensor& conv_state, const Tensor& qsl, - const Tensor& his, const CausalConv1dArgs& args) { +void LaunchConvFwdRegRuntime(cudaStream_t s, Tensor& out, const Tensor& x, + const Tensor& w, const Tensor* bias, + Tensor& conv_state, const Tensor& qsl, + const Tensor& his, const CausalConv1dArgs& args, + int64_t chan_tiles) { const int64_t n = conv_state.shape[0], c = x.shape[1], k = w.shape[1]; const int64_t total_tokens = x.shape[0]; const int64_t x_rs = x.stride[0]; // padded-row (merged qkvz) x view honored VT_CHECK(k - 1 <= kConvRegMaxW, "cuda causal_conv1d_fwd(reg): conv width exceeds kConvRegMaxW"); - const int64_t chan_tiles = (c + kConvRegN - 1) / kConvRegN; int64_t gridZ = 1; int chunked = 0; const bool exact = ConvExactChunksEnabled() && args.batch_ptr != nullptr; @@ -906,6 +1066,28 @@ void LaunchConvFwdReg(cudaStream_t s, Tensor& out, const Tensor& x, const Tensor Check(cudaGetLastError(), "causal_conv1d_fwd(reg) launch"); } +template +void LaunchConvFwdReg(cudaStream_t s, Tensor& out, const Tensor& x, + const Tensor& w, const Tensor* bias, Tensor& conv_state, + const Tensor& qsl, const Tensor& his, + const CausalConv1dArgs& args) { + const int64_t c = x.shape[1], k = w.shape[1]; + DispatchConvChannelTileLaunch( + std::getenv("VT_CONV_CHANNEL_TILE"), c, k, + [&](const ConvChannelTileLaunchContract& contract) { + LaunchConvFwdRegRuntime(s, out, x, w, bias, conv_state, qsl, + his, args, contract.feature_blocks); + }, + [&](const ConvChannelTileLaunchContract& contract) { + LaunchConvFwdRegK4<1, Tin, Tout>(s, out, x, w, bias, conv_state, qsl, + his, args, contract.feature_blocks); + }, + [&](const ConvChannelTileLaunchContract& contract) { + LaunchConvFwdRegK4<2, Tin, Tout>(s, out, x, w, bias, conv_state, qsl, + his, args, contract.feature_blocks); + }); +} + // Dispatch on the toggles. reg (VT_CONV_REG, default ON) wins; else tiled // (VT_CONV_TILED); else the EXACT scalar launcher (byte-identical grid/block/kernel). template @@ -1384,7 +1566,101 @@ __global__ void GdnPostConvFastKernel(Tqkv* q_out, Tqkv* k_out, Tqkv* v_out, flo } } -// Split post-conv (VT_GDN_POSTCONV_SPLIT, default ON) — grid (T, Hk+Hv), mirroring +// Experimental 1:1 CUDA spelling of vLLM/FLA's post-conv launch schedule: +// grid=(ceil(T,16), Hk+Hv), BLOCK_T=16, four warps. Each warp owns four tokens +// in the tile. For Q/K, each lane keeps four feature values in registers across +// the float32 norm reduction and normalized store, eliminating the fast kernel's +// second conv read and all shared-memory barriers. For V, the same mapping gives +// four coalesced 32-element copy waves per token and lane zero computes gating. +// Arithmetic and strides otherwise retain the local GdnPostConv contract. +template +__global__ void GdnPostConvTokenTileKernel( + Tqkv* q_out, Tqkv* k_out, Tqkv* v_out, float* g_out, float* beta_out, + const Tconv* conv, const Tgate* araw, const Tgate* braw, const float* a_log, + const float* dt_bias, int64_t t, int64_t hk, int64_t dk, int64_t hv, int64_t dv, + int64_t a_row_stride, int64_t b_row_stride, float eps) { + constexpr int kWarps = 4; + constexpr int kFeaturesPerLane = 4; + const int lane = static_cast(threadIdx.x) & 31; + const int warp = static_cast(threadIdx.x) >> 5; + const int64_t tile_start = static_cast(blockIdx.x) * kGdnPostConvTokenTileTokens; + const int64_t head = blockIdx.y; + const int64_t key_dim = hk * dk; + const int64_t value_dim = hv * dv; + const int64_t conv_dim = 2 * key_dim + value_dim; + + if (head < hk) { +#pragma unroll + for (int tile_token = warp; tile_token < kGdnPostConvTokenTileTokens; + tile_token += kWarps) { + const int64_t tok = tile_start + tile_token; + if (tok >= t) continue; + const int64_t row = tok * conv_dim; + const Tconv* qin = conv + row + head * dk; + const Tconv* kin = conv + row + key_dim + head * dk; + Tqkv* qo = q_out + (tok * hk + head) * dk; + Tqkv* ko = k_out + (tok * hk + head) * dk; + + float qv[kFeaturesPerLane]; + float kv[kFeaturesPerLane]; +#pragma unroll + for (int item = 0; item < kFeaturesPerLane; ++item) { + const int feature = lane + item * 32; + qv[item] = Load(qin, feature); + kv[item] = Load(kin, feature); + } + // Reproduce the fast kernel's 128-lane shared reduction exactly. Its first + // two levels are (i+i+64), then (i+i+32); those four terms are resident in + // this lane. The remaining 32-lane tree is the same 16,8,4,2,1 order. + float qsum = (qv[0] * qv[0] + qv[2] * qv[2]) + + (qv[1] * qv[1] + qv[3] * qv[3]); + float ksum = (kv[0] * kv[0] + kv[2] * kv[2]) + + (kv[1] * kv[1] + kv[3] * kv[3]); +#pragma unroll + for (int offset = 16; offset > 0; offset >>= 1) { + qsum += __shfl_down_sync(0xffffffffu, qsum, offset); + ksum += __shfl_down_sync(0xffffffffu, ksum, offset); + } + qsum = __shfl_sync(0xffffffffu, qsum, 0); + ksum = __shfl_sync(0xffffffffu, ksum, 0); + const float qinv = 1.0f / sqrtf(qsum + eps); + const float kinv = 1.0f / sqrtf(ksum + eps); +#pragma unroll + for (int item = 0; item < kFeaturesPerLane; ++item) { + const int feature = lane + item * 32; + Store(qo, feature, qv[item] * qinv); + Store(ko, feature, kv[item] * kinv); + } + } + } else { + const int64_t value_head = head - hk; +#pragma unroll + for (int tile_token = warp; tile_token < kGdnPostConvTokenTileTokens; + tile_token += kWarps) { + const int64_t tok = tile_start + tile_token; + if (tok >= t) continue; + const int64_t row = tok * conv_dim; + const Tconv* vin = conv + row + 2 * key_dim + value_head * dv; + Tqkv* vo = v_out + tok * value_dim + value_head * dv; +#pragma unroll + for (int item = 0; item < kFeaturesPerLane; ++item) { + const int feature = lane + item * 32; + Store(vo, feature, Load(vin, feature)); + } + if (lane == 0) { + const int64_t idx = tok * hv + value_head; + const float av = Load(araw, tok * a_row_stride + value_head); + const float bv = Load(braw, tok * b_row_stride + value_head); + const float x = av + dt_bias[value_head]; + const float sp = x > 20.0f ? x : log1pf(expf(x)); + g_out[idx] = -expf(a_log[value_head]) * sp; + beta_out[idx] = 1.0f / (1.0f + expf(-bv)); + } + } + } +} + +// Split post-conv (VT_GDN_POSTCONV_SPLIT, opt-in) — grid (T, Hk+Hv), mirroring // vLLM's grid (ceil(L,BLOCK_T), H+HV) in _fused_post_conv_kernel:57-149 where each V // head is its own program. The shipped GdnPostConvKernel packs the ENTIRE // value_dim = Hv*Dv copy + all Hv gating scalars into ONE grid.y block per token @@ -1457,7 +1733,7 @@ __global__ void GdnPostConvSplitKernel(Tqkv* q_out, Tqkv* k_out, Tqkv* v_out, fl } } -// Toggle: DEFAULT ON (VT_GDN_POSTCONV_SPLIT=0 restores the single-megablock kernel). +// Toggle: DEFAULT OFF; a present non-'0'-leading value selects the split kernel. // Read per call (post-conv dispatch is coarse — one launch/step). bool GdnPostConvSplitEnabled() { return GdnPostConvSplitFlagIsOn(std::getenv("VT_GDN_POSTCONV_SPLIT")); @@ -1480,12 +1756,21 @@ void GdnPostConvKernelCuda(Queue& q, Tensor& q_out, Tensor& k_out, Tensor& v_out // VT_GDN_POSTCONV_SPLIT (opt-in): grid (T, Hk+Hv) — each V head its own block // (mirrors vLLM). =0 (default) uses the single-megablock grid (T, Hk+1). const bool split = GdnPostConvSplitEnabled(); + // VT_GDN_POSTCONV_TOKEN_TILE (opt-in): the full upstream 16-token, + // per-head/four-warp schedule. The explicit split control keeps priority when + // both experimental flags are set. Only the production 128-wide heads route. + const bool token_tile = + !split && GdnPostConvTokenTileFlagIsOn(std::getenv("VT_GDN_POSTCONV_TOKEN_TILE")) && + dk == 128 && dv == 128; // VT_GDN_POSTCONV_FAST: byte-identical megablock at 128 threads + 128-bit V copy. // Only for the Dk==Dv==128 gate dims (16B alignment + value_dim%8==0); mutually // exclusive with the split (both target the same megablock). See predicate. - const bool fast = !split && GdnPostConvFastFlagIsOn(std::getenv("VT_GDN_POSTCONV_FAST")) && - dk == 128 && dv == 128; - dim3 grid(static_cast(t), static_cast(split ? hk + hv : hk + 1)); + const bool fast = !split && !token_tile && + GdnPostConvFastFlagIsOn(std::getenv("VT_GDN_POSTCONV_FAST")) && dk == 128 && + dv == 128; + const unsigned grid_x = static_cast(token_tile ? GdnPostConvTokenTileGridX(t) : t); + const unsigned grid_y = static_cast((split || token_tile) ? hk + hv : hk + 1); + dim3 grid(grid_x, grid_y); cudaStream_t s = AsStream(q); // Dispatch over (q/k/v out dtype) x (conv-in dtype). conv is bf16 under the // input-side bf16 GDN path (VT_GDN_IN_BF16); the conv read upcasts to f32. @@ -1493,7 +1778,13 @@ void GdnPostConvKernelCuda(Queue& q, Tensor& q_out, Tensor& k_out, Tensor& v_out using Tqkv = decltype(qkv_tag); using Tconv = decltype(conv_tag); using Tgate = decltype(gate_tag); - if (fast) { + if (token_tile) { + GdnPostConvTokenTileKernel<<>>( + q_out.Ptr(), k_out.Ptr(), v_out.Ptr(), g_out.Ptr(), + beta_out.Ptr(), conv.Ptr(), araw.Ptr(), braw.Ptr(), + a_log.Ptr(), dt_bias.Ptr(), t, hk, dk, hv, dv, araw.stride[0], + braw.stride[0], args.eps); + } else if (fast) { GdnPostConvFastKernel<<>>( q_out.Ptr(), k_out.Ptr(), v_out.Ptr(), g_out.Ptr(), beta_out.Ptr(), conv.Ptr(), araw.Ptr(), braw.Ptr(), @@ -2487,11 +2778,14 @@ void GdnPackedDecodeKernelCuda(Queue& q, Tensor& out, // NW consecutive lanes [floor(lane/NW)*NW, +NW), aligned inside one warp (NW is // a power of two dividing 32 and the launcher only uses NW>1 when BV==32, so a // block is always a whole number of warps) — the xor butterfly stays in-warp. -template +template __global__ void GdnDecodeFusedKernel(Tout* out, const Tin* q, const Tin* k, const Tin* v, const float* g, const float* beta, TState* state, const int32_t* state_idx, int64_t hk_n, int64_t dk, int64_t hv_n, int64_t dv, int64_t bv, float scale) { + static_assert(!REGSTATE || (SWIZZLED && NW == 8), + "register state requires the production swizzled NW8 layout"); const int64_t i_v = blockIdx.x; // value-dim tile const int64_t i_nh = blockIdx.y; // fused (sequence, v-head) const int64_t i_n = i_nh / hv_n; // sequence == decode token index @@ -2516,7 +2810,7 @@ __global__ void GdnDecodeFusedKernel(Tout* out, const Tin* q, const Tin* k, cons row = si; } - const int64_t sdk = dk + 1; // padded shared row stride (kills the 32-way conflict) + const int64_t sdk = SWIZZLED ? dk + NW : dk + 1; extern __shared__ float smem[]; float* bq = smem; // [dk] q' = q*scale float* bk = bq + dk; // [dk] k @@ -2527,18 +2821,30 @@ __global__ void GdnDecodeFusedKernel(Tout* out, const Tin* q, const Tin* k, cons // q'(=q*scale) and k for this (token, head) — broadcast to every lane. const int64_t qkbase = (i_n * hk_n + hk) * dk; for (int64_t i = tid; i < dk; i += blockDim.x) { - bq[i] = Load(q, qkbase + i) * scale; - bk[i] = Load(k, qkbase + i); + const int64_t si = SWIZZLED ? (i % 16) * 8 + i / 16 : i; + bq[si] = Load(q, qkbase + i) * scale; + bk[si] = Load(k, qkbase + i); } // Coalesced load of the [BV,Dk] state slice into padded shared. The persistent // cache TState is bf16 (mirrors vLLM's default mamba_cache_dtype=auto→model // dtype; fla fused_recurrent reads bf16→f32 registers→writes bf16) or f32 // (unit test). Load() upcasts to f32; the recurrence below runs in f32. TState* s_head = state + (row * hv_n + hv) * dv * dk + vbase * dk; // [<=bv, dk] - for (int64_t e = tid; e < bv * dk; e += blockDim.x) - sbh[(e / dk) * sdk + e % dk] = e < tile ? Load(s_head, e) : 0.0f; + for (int64_t e = tid; e < bv * dk; e += blockDim.x) { + const int64_t c = e % dk; + const int64_t sc = SWIZZLED ? (c % 16) * 8 + c / 16 : c; + sbh[(e / dk) * sdk + sc] = e < tile ? Load(s_head, e) : 0.0f; + } __syncthreads(); + float* r = sbh + static_cast(vi) * sdk; + float rr[16]; + if constexpr (REGSTATE) { + for (int64_t j = 0; j < 16; ++j) { + rr[j] = r[GdnDecodeRegisterSharedColumn(wk, j)]; + } + } + // This thread's Dk column slice [c0, c1) of value-row vi (partition of [0,dk)). const int64_t ck = (dk + NW - 1) / NW; const int64_t c0 = static_cast(wk) * ck; @@ -2548,11 +2854,24 @@ __global__ void GdnDecodeFusedKernel(Tout* out, const Tin* q, const Tin* k, cons // slice); only the global v-load and o-store are guarded by vrow < dv. const float decay = expf(g[i_n * hv_n + hv]); const float beta_t = beta[i_n * hv_n + hv]; - float* r = sbh + static_cast(vi) * sdk; float pdot = 0.0f; // partial (S * exp(g)) @ k over this slice, fused w/ decay - for (int64_t c = c0; c < c1; ++c) { - r[c] *= decay; - pdot += r[c] * bk[c]; + if constexpr (REGSTATE) { + for (int64_t j = 0; j < 16; ++j) { + const int64_t sc = GdnDecodeRegisterSharedColumn(wk, j); + rr[j] *= decay; + pdot += rr[j] * bk[sc]; + } + } else if constexpr (SWIZZLED) { + for (int64_t j = 0; j < 16; ++j) { + const int64_t sc = j * 8 + wk; + r[sc] *= decay; + pdot += r[sc] * bk[sc]; + } + } else { + for (int64_t c = c0; c < c1; ++c) { + r[c] *= decay; + pdot += r[c] * bk[c]; + } } float dot = pdot; // reduce the partials across the NW lanes of the row-group #pragma unroll @@ -2560,64 +2879,116 @@ __global__ void GdnDecodeFusedKernel(Tout* out, const Tin* q, const Tin* k, cons const float vv = vrow < dv ? Load(v, (i_n * hv_n + hv) * dv + vrow) : 0.0f; const float vp = (vv - dot) * beta_t; float po = 0.0f; // partial (S + outer(v',k)) @ q' over this slice, fused w/ update - for (int64_t c = c0; c < c1; ++c) { - r[c] += vp * bk[c]; - po += r[c] * bq[c]; + if constexpr (REGSTATE) { + for (int64_t j = 0; j < 16; ++j) { + const int64_t sc = GdnDecodeRegisterSharedColumn(wk, j); + rr[j] += vp * bk[sc]; + po += rr[j] * bq[sc]; + } + } else if constexpr (SWIZZLED) { + for (int64_t j = 0; j < 16; ++j) { + const int64_t sc = j * 8 + wk; + r[sc] += vp * bk[sc]; + po += r[sc] * bq[sc]; + } + } else { + for (int64_t c = c0; c < c1; ++c) { + r[c] += vp * bk[c]; + po += r[c] * bq[c]; + } } float o = po; #pragma unroll for (int off = 1; off < NW; off <<= 1) o += __shfl_xor_sync(0xffffffffu, o, off); if (vrow < dv && wk == 0) Store(out, (i_n * hv_n + hv) * dv + vrow, o); + if constexpr (REGSTATE) { + for (int64_t j = 0; j < 16; ++j) { + r[GdnDecodeRegisterSharedColumn(wk, j)] = rr[j]; + } + } __syncthreads(); // Coalesced write-back of the updated slice from f32 registers to the // configured fp16/bf16/fp32 temporal cache. - for (int64_t e = tid; e < tile; e += blockDim.x) - Store(s_head, e, sbh[(e / dk) * sdk + e % dk]); + for (int64_t e = tid; e < tile; e += blockDim.x) { + const int64_t c = e % dk; + const int64_t sc = SWIZZLED ? (c % 16) * 8 + c / 16 : c; + Store(s_head, e, sbh[(e / dk) * sdk + sc]); + } } template void LaunchGdnDecodeFusedNW(cudaStream_t s, Tensor& out, const Tensor& q_in, const Tensor& k, const Tensor& v, const Tensor& g, const Tensor& beta, Tensor& state, - const int32_t* state_idx, int64_t n, const GdnArgs& args) { + const int32_t* state_idx, int64_t n, const GdnArgs& args, + const GdnDecodeLaunchContract& contract) { const int64_t hk_n = q_in.shape[1], dk = q_in.shape[2]; const int64_t hv_n = v.shape[1], dv = v.shape[2]; - const int64_t bv = dv < 32 ? dv : 32; // fla BV cap of 32; any dv via tail guard - const int64_t nv = (dv + bv - 1) / bv; // value-dim tiles (fla NV) - const dim3 grid(static_cast(nv), static_cast(n * hv_n)); - const size_t shmem = - (2 * static_cast(dk) + static_cast(bv) * (dk + 1)) * sizeof(float); - GdnDecodeFusedKernel<<(bv * NW), shmem, s>>>( - out.Ptr(), q_in.Ptr(), k.Ptr(), v.Ptr(), g.Ptr(), - beta.Ptr(), state.Ptr(), state_idx, hk_n, dk, hv_n, dv, bv, args.scale); + const int64_t bv = contract.value_tile; + const dim3 grid(static_cast(contract.value_tiles), + static_cast(n * hv_n)); + if constexpr (NW == 8) { + if (contract.swizzled) { + DispatchGdnDecodeStateStorage( + contract, + [&](const GdnDecodeLaunchContract&) { + GdnDecodeFusedKernel + <<(contract.block_threads), + contract.shared_bytes, s>>>( + out.Ptr(), q_in.Ptr(), k.Ptr(), + v.Ptr(), g.Ptr(), beta.Ptr(), + state.Ptr(), state_idx, hk_n, dk, hv_n, dv, bv, + args.scale); + }, + [&](const GdnDecodeLaunchContract&) { + GdnDecodeFusedKernel + <<(contract.block_threads), + contract.shared_bytes, s>>>( + out.Ptr(), q_in.Ptr(), k.Ptr(), + v.Ptr(), g.Ptr(), beta.Ptr(), + state.Ptr(), state_idx, hk_n, dk, hv_n, dv, bv, + args.scale); + }); + } else { + GdnDecodeFusedKernel + <<(contract.block_threads), contract.shared_bytes, s>>>( + out.Ptr(), q_in.Ptr(), k.Ptr(), v.Ptr(), g.Ptr(), + beta.Ptr(), state.Ptr(), state_idx, hk_n, dk, hv_n, dv, bv, + args.scale); + } + } else { + GdnDecodeFusedKernel + <<(contract.block_threads), contract.shared_bytes, s>>>( + out.Ptr(), q_in.Ptr(), k.Ptr(), v.Ptr(), g.Ptr(), + beta.Ptr(), state.Ptr(), state_idx, hk_n, dk, hv_n, dv, bv, args.scale); + } Check(cudaGetLastError(), "gdn decode(fused) launch"); } -// nw: warps-per-block for the Dk-split (occupancy lever). >1 only when BV==32 -// (dv>=32, the real gate dim) so a block is always a whole number of warps and -// the row-group shuffles stay in-warp; smaller dv (test corners) forces nw=1. +// lanes_per_row: cooperative lanes for each value row's Dk split. The launch +// contract keeps each row's NW consecutive lanes together for every value tile; +// smaller dv test corners preserve the shipped NW=1 rule. template void LaunchGdnDecodeFused(cudaStream_t s, Tensor& out, const Tensor& q_in, const Tensor& k, const Tensor& v, const Tensor& g, const Tensor& beta, Tensor& state, - const int32_t* state_idx, int64_t n, const GdnArgs& args, int nw) { - const int64_t dv = v.shape[2]; - const int nw_eff = dv >= 32 ? nw : 1; - switch (nw_eff) { + const int32_t* state_idx, int64_t n, const GdnArgs& args, + const GdnDecodeLaunchContract& contract) { + switch (contract.lanes_per_row) { case 2: LaunchGdnDecodeFusedNW(s, out, q_in, k, v, g, beta, state, state_idx, - n, args); + n, args, contract); break; case 4: LaunchGdnDecodeFusedNW(s, out, q_in, k, v, g, beta, state, state_idx, - n, args); + n, args, contract); break; case 8: LaunchGdnDecodeFusedNW(s, out, q_in, k, v, g, beta, state, state_idx, - n, args); + n, args, contract); break; default: LaunchGdnDecodeFusedNW(s, out, q_in, k, v, g, beta, state, state_idx, - n, args); + n, args, contract); break; } } @@ -2627,16 +2998,17 @@ void LaunchGdnDecodeFused(cudaStream_t s, Tensor& out, const Tensor& q_in, const template void LaunchGdnDecodeFusedS(cudaStream_t s, Tensor& out, const Tensor& q_in, const Tensor& k, const Tensor& v, const Tensor& g, const Tensor& beta, Tensor& state, - const int32_t* state_idx, int64_t n, const GdnArgs& args, int nw) { + const int32_t* state_idx, int64_t n, const GdnArgs& args, + const GdnDecodeLaunchContract& contract) { if (state.dtype == DType::kBF16) LaunchGdnDecodeFused(s, out, q_in, k, v, g, beta, state, state_idx, - n, args, nw); + n, args, contract); else if (state.dtype == DType::kF16) LaunchGdnDecodeFused(s, out, q_in, k, v, g, beta, - state, state_idx, n, args, nw); + state, state_idx, n, args, contract); else LaunchGdnDecodeFused(s, out, q_in, k, v, g, beta, state, state_idx, n, args, - nw); + contract); } // Decode dispatch. state_idx == nullptr: compact [n,Hv,Dv,Dk] state (row==i_n). @@ -2653,14 +3025,6 @@ void GdnDecodeFusedCuda(Queue& q, Tensor& out, const Tensor& q_in, const Tensor& const int64_t n = q_in.shape[0], hv_n = state.shape[1], dv = state.shape[2], dk = state.shape[3]; if (n == 0 || hv_n == 0 || dv == 0) return; VT_CHECK(n * hv_n <= kMaxGridY, "cuda gdn_decode: too many (seq×head) blocks (grid.y limit)"); - const int64_t bv = dv < 32 ? dv : 32; - const size_t shmem = - (2 * static_cast(dk) + static_cast(bv) * (dk + 1)) * sizeof(float); - if (shmem > 48 * 1024) { // corner-dim fallback (no decode test hits this; real dims are 128) - GdnScanCuda(q, out, q_in, k, v, g, beta, state, nullptr, state_idx, - args, "gdn_decode"); - return; - } // Warps-per-block for the Dk-split occupancy lever (default 8 — measured best // on GB10 sm_121: raises GdnDecodeFused theoretical occupancy 10.4%→66.7% and // cuts per-call time ~2.2× at conc-64; A/B via env, read once per call — @@ -2672,21 +3036,33 @@ void GdnDecodeFusedCuda(Queue& q, Tensor& out, const Tensor& q_in, const Tensor& if (v_nw == 1 || v_nw == 2 || v_nw == 4 || v_nw == 8) nw = v_nw; } cudaStream_t s = AsStream(q); - if (q_in.dtype == DType::kF32) { - if (out.dtype == DType::kF32) - LaunchGdnDecodeFusedS(s, out, q_in, k, v, g, beta, state, state_idx, n, args, - nw); - else - LaunchGdnDecodeFusedS(s, out, q_in, k, v, g, beta, state, state_idx, n, - args, nw); - } else { - if (out.dtype == DType::kF32) - LaunchGdnDecodeFusedS<__nv_bfloat16, float>(s, out, q_in, k, v, g, beta, state, state_idx, n, - args, nw); - else - LaunchGdnDecodeFusedS<__nv_bfloat16, __nv_bfloat16>(s, out, q_in, k, v, g, beta, state, - state_idx, n, args, nw); - } + auto launch = [&](const GdnDecodeLaunchContract& contract) { + if (contract.shared_bytes > 48 * 1024) { + // Corner-dim fallback (no decode test hits this; real dims are 128). + GdnScanCuda(q, out, q_in, k, v, g, beta, state, nullptr, state_idx, + args, "gdn_decode"); + return; + } + if (q_in.dtype == DType::kF32) { + if (out.dtype == DType::kF32) + LaunchGdnDecodeFusedS(s, out, q_in, k, v, g, beta, state, state_idx, n, + args, contract); + else + LaunchGdnDecodeFusedS(s, out, q_in, k, v, g, beta, state, state_idx, + n, args, contract); + } else { + if (out.dtype == DType::kF32) + LaunchGdnDecodeFusedS<__nv_bfloat16, float>(s, out, q_in, k, v, g, beta, state, state_idx, + n, args, contract); + else + LaunchGdnDecodeFusedS<__nv_bfloat16, __nv_bfloat16>( + s, out, q_in, k, v, g, beta, state, state_idx, n, args, contract); + } + }; + DispatchGdnDecodeValueTile(std::getenv("VT_GDN_DECODE_BV"), + std::getenv("VT_GDN_DECODE_SWIZZLE"), + std::getenv("VT_GDN_DECODE_REGSTATE"), dv, dk, + nw, launch, launch); } // Shared wrapper body: qsl_ptr == nullptr → decode (n = batch = state rows; diff --git a/src/vt/cuda/gdn_decode_fused.h b/src/vt/cuda/gdn_decode_fused.h new file mode 100644 index 000000000..c98d5a6c5 --- /dev/null +++ b/src/vt/cuda/gdn_decode_fused.h @@ -0,0 +1,171 @@ +// vllm.cpp original. CUDA-free selector and launch-geometry contract for the +// fused single-token GDN recurrence implemented in cuda_gdn.cu. +#ifndef VT_CUDA_GDN_DECODE_FUSED_H_ +#define VT_CUDA_GDN_DECODE_FUSED_H_ + +#include +#include + +namespace vt::cuda { + +enum class GdnDecodeValueTile : uint8_t { + kBv16 = 16, + kBv32 = 32, +}; + +// BV16 is an exact opt-in discriminator. Every spelling except exactly "16" +// preserves the shipped BV32 schedule, including removed sweep values. +inline constexpr GdnDecodeValueTile GdnDecodeValueTileFromEnv(const char* env_value) { + return env_value != nullptr && env_value[0] == '1' && env_value[1] == '6' && + env_value[2] == '\0' + ? GdnDecodeValueTile::kBv16 + : GdnDecodeValueTile::kBv32; +} + +// The shared-bank swizzle is an exact opt-in. In particular, numeric prefixes, +// whitespace, and truthy words must leave the incumbent layout unchanged. +inline constexpr bool GdnDecodeSwizzleFlagIsOn(const char* env_value) { + return env_value != nullptr && env_value[0] == '1' && env_value[1] == '\0'; +} + +// Register-resident recurrence state is an exact opt-in. Selection is further +// constrained by GdnDecodeLaunchContractFor to the production BV16+SWIZZLE +// geometry; the parser alone never widens eligibility. +inline constexpr bool GdnDecodeRegstateFlagIsOn(const char* env_value) { + return env_value != nullptr && env_value[0] == '1' && env_value[1] == '\0'; +} + +inline constexpr int64_t GdnDecodeStateStride(bool swizzled, int64_t dk, + int lanes_per_row) { + return dk + (swizzled ? lanes_per_row : 1); +} + +// Maps a logical Dk column to its shared-memory column. Production only selects +// the swizzled arm for Dk=128/NW=8, where ck=16. +inline constexpr int64_t GdnDecodeSharedColumn(bool swizzled, int64_t logical_column, + int64_t dk, int lanes_per_row) { + if (!swizzled) return logical_column; + const int64_t columns_per_lane = dk / lanes_per_row; + return (logical_column % columns_per_lane) * lanes_per_row + + logical_column / columns_per_lane; +} + +inline constexpr int64_t GdnDecodeRegisterLogicalColumn(int lane, int slot) { + return static_cast(lane) * 16 + slot; +} + +#if defined(__CUDACC__) +__host__ __device__ +#endif +inline constexpr int64_t GdnDecodeRegisterSharedColumn(int lane, int slot) { + return static_cast(slot) * 8 + lane; +} + +struct GdnDecodeLaunchContract { + GdnDecodeValueTile selected_tile; + int64_t value_tile; + int64_t value_tiles; + int lanes_per_row; + int64_t block_threads; + size_t shared_bytes; + bool should_launch; + bool swizzled; + bool regstate; +}; + +inline constexpr GdnDecodeLaunchContract GdnDecodeLaunchContractFor( + const char* bv_env_value, const char* swizzle_env_value, + const char* regstate_env_value, int64_t dv, int64_t dk, + int requested_nw) { + const GdnDecodeValueTile selected = GdnDecodeValueTileFromEnv(bv_env_value); + if (dv <= 0 || dk <= 0) { + return GdnDecodeLaunchContract{selected, 0, 0, 0, 0, 0, false, false, + false}; + } + const int64_t requested_bv = static_cast(selected); + const int64_t bv = dv < requested_bv ? dv : requested_bv; + const int64_t nv = (dv + bv - 1) / bv; + const int nw_eff = dv >= 32 ? requested_nw : 1; + const int64_t block_threads = bv * nw_eff; + const bool swizzled = GdnDecodeSwizzleFlagIsOn(swizzle_env_value) && + selected == GdnDecodeValueTile::kBv16 && dv == 128 && + dk == 128 && nw_eff == 8; + const bool regstate = GdnDecodeRegstateFlagIsOn(regstate_env_value) && + swizzled; + const int64_t state_stride = GdnDecodeStateStride(swizzled, dk, nw_eff); + const size_t shared_bytes = + (2 * static_cast(dk) + + static_cast(bv) * static_cast(state_stride)) * + sizeof(float); + return GdnDecodeLaunchContract{selected, bv, nv, nw_eff, block_threads, + shared_bytes, true, swizzled, regstate}; +} + +inline constexpr GdnDecodeLaunchContract GdnDecodeLaunchContractFor( + const char* bv_env_value, const char* swizzle_env_value, int64_t dv, + int64_t dk, int requested_nw) { + return GdnDecodeLaunchContractFor(bv_env_value, swizzle_env_value, nullptr, + dv, dk, requested_nw); +} + +inline constexpr GdnDecodeLaunchContract GdnDecodeLaunchContractFor( + const char* bv_env_value, int64_t dv, int64_t dk, int requested_nw) { + return GdnDecodeLaunchContractFor(bv_env_value, nullptr, nullptr, dv, dk, + requested_nw); +} + +// Storage callback dispatcher used by production and portable tests. The +// resolved contract selects exactly one compile-time kernel specialization. +template +inline decltype(auto) DispatchGdnDecodeStateStorage( + const GdnDecodeLaunchContract& contract, SharedLaunch&& shared_launch, + RegisterLaunch&& register_launch) { + if (contract.regstate) return register_launch(contract); + return shared_launch(contract); +} + +// Shared callback dispatcher used by production and portable tests. The +// callbacks receive fully resolved geometry and exactly one is invoked. +template +inline decltype(auto) DispatchGdnDecodeValueTile(const char* env_value, int64_t dv, + int64_t dk, int requested_nw, + Bv16Launch&& bv16_launch, + Bv32Launch&& bv32_launch) { + const GdnDecodeLaunchContract contract = + GdnDecodeLaunchContractFor(env_value, dv, dk, requested_nw); + if (contract.selected_tile == GdnDecodeValueTile::kBv16) { + return bv16_launch(contract); + } + return bv32_launch(contract); +} + +template +inline decltype(auto) DispatchGdnDecodeValueTile( + const char* bv_env_value, const char* swizzle_env_value, int64_t dv, + int64_t dk, int requested_nw, Bv16Launch&& bv16_launch, + Bv32Launch&& bv32_launch) { + const GdnDecodeLaunchContract contract = GdnDecodeLaunchContractFor( + bv_env_value, swizzle_env_value, dv, dk, requested_nw); + if (contract.selected_tile == GdnDecodeValueTile::kBv16) { + return bv16_launch(contract); + } + return bv32_launch(contract); +} + +template +inline decltype(auto) DispatchGdnDecodeValueTile( + const char* bv_env_value, const char* swizzle_env_value, + const char* regstate_env_value, int64_t dv, int64_t dk, int requested_nw, + Bv16Launch&& bv16_launch, Bv32Launch&& bv32_launch) { + const GdnDecodeLaunchContract contract = GdnDecodeLaunchContractFor( + bv_env_value, swizzle_env_value, regstate_env_value, dv, dk, + requested_nw); + if (contract.selected_tile == GdnDecodeValueTile::kBv16) { + return bv16_launch(contract); + } + return bv32_launch(contract); +} + +} // namespace vt::cuda + +#endif // VT_CUDA_GDN_DECODE_FUSED_H_ diff --git a/src/vt/cuda/gdn_prefill_conv.h b/src/vt/cuda/gdn_prefill_conv.h index c87277998..3510ee4f4 100644 --- a/src/vt/cuda/gdn_prefill_conv.h +++ b/src/vt/cuda/gdn_prefill_conv.h @@ -60,6 +60,8 @@ #ifndef VT_CUDA_GDN_PREFILL_CONV_H_ #define VT_CUDA_GDN_PREFILL_CONV_H_ +#include + namespace vt::cuda { // Pure predicate for the VT_CONV_REG contract: DEFAULT ON. The register-window @@ -79,6 +81,80 @@ inline bool ConvExactChunksFlagIsOn(const char* env_value) { return env_value == nullptr || env_value[0] != '0'; } +// Three-arm same-binary experiment for the remaining prefill causal-conv gap. +// Arm 0 is the sealed runtime-width kernel. Arms 1 and 2 are valid only for the +// production width K=4: respectively one and two channels per 128-thread lane. +// Unset and every spelling except the exact strings "1" and "2" preserve arm 0. +enum class ConvChannelTileArm : uint8_t { + kRuntimeWidth = 0, + kWidthFour = 1, + kWidthFourTwoChannels = 2, +}; + +inline constexpr ConvChannelTileArm ConvChannelTileArmFromEnv(const char* env_value) { + if (env_value != nullptr && env_value[0] == '1' && env_value[1] == '\0') { + return ConvChannelTileArm::kWidthFour; + } + if (env_value != nullptr && env_value[0] == '2' && env_value[1] == '\0') { + return ConvChannelTileArm::kWidthFourTwoChannels; + } + return ConvChannelTileArm::kRuntimeWidth; +} + +inline constexpr ConvChannelTileArm ResolveConvChannelTileArm(ConvChannelTileArm requested, + int64_t channels, + int64_t kernel_width) { + if (requested != ConvChannelTileArm::kRuntimeWidth && + (channels <= 0 || kernel_width != 4)) { + return ConvChannelTileArm::kRuntimeWidth; + } + return requested; +} + +struct ConvChannelTileLaunchContract { + ConvChannelTileArm arm; + int64_t feature_blocks; + int64_t threads_per_block; +}; + +inline constexpr int64_t kConvChannelTileThreads = 128; + +inline constexpr ConvChannelTileLaunchContract ConvChannelTileLaunchContractFor( + const char* env_value, int64_t channels, int64_t kernel_width) { + const ConvChannelTileArm arm = ResolveConvChannelTileArm( + ConvChannelTileArmFromEnv(env_value), channels, kernel_width); + const int64_t channels_per_block = + arm == ConvChannelTileArm::kWidthFourTwoChannels ? 256 : 128; + return ConvChannelTileLaunchContract{ + arm, + channels > 0 ? (channels + channels_per_block - 1) / channels_per_block : 0, + kConvChannelTileThreads, + }; +} + +// One shared dispatch seam for both the CUDA launcher and portable tests. Keeping +// the arm selection here means the tests exercise the exact branch logic used in +// production without adding a launch counter or other debug state to the hot path. +// The callbacks inline away at each call site and receive the already-resolved +// geometry, including runtime-width fallback for unsupported shapes. +template +inline decltype(auto) DispatchConvChannelTileLaunch( + const char* env_value, int64_t channels, int64_t kernel_width, + RuntimeWidthLaunch&& runtime_width_launch, + WidthFourLaunch&& width_four_launch, + WidthFourTwoChannelsLaunch&& width_four_two_channels_launch) { + const ConvChannelTileLaunchContract contract = + ConvChannelTileLaunchContractFor(env_value, channels, kernel_width); + if (contract.arm == ConvChannelTileArm::kWidthFour) { + return width_four_launch(contract); + } + if (contract.arm == ConvChannelTileArm::kWidthFourTwoChannels) { + return width_four_two_channels_launch(contract); + } + return runtime_width_launch(contract); +} + // Pure predicate for the VT_GDN_POSTCONV_SPLIT contract: DEFAULT OFF (OPT-IN). The // split post-conv kernel (GdnPostConvSplitKernel) is BIT-IDENTICAL (0-ulp) to the // shipped GdnPostConvKernel by construction, but the DGX nsys A/B measured it @@ -119,6 +195,20 @@ inline bool GdnPostConvFastFlagIsOn(const char* env_value) { return env_value == nullptr || env_value[0] != '0'; } +// Experimental spelling of vLLM's fused post-conv work partition: 16 tokens per +// block, one Q/K or V head per grid.y program, four warps. It remains opt-in until +// same-binary correctness and performance gates close; unset and '0' keep the +// byte-identical fast megablock default above. +inline bool GdnPostConvTokenTileFlagIsOn(const char* env_value) { + return env_value != nullptr && env_value[0] != '0'; +} + +inline constexpr int64_t kGdnPostConvTokenTileTokens = 16; + +inline constexpr int64_t GdnPostConvTokenTileGridX(int64_t tokens) { + return (tokens + kGdnPostConvTokenTileTokens - 1) / kGdnPostConvTokenTileTokens; +} + } // namespace vt::cuda #endif // VT_CUDA_GDN_PREFILL_CONV_H_ diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt index 296e66b3f..3ed1ec824 100644 --- a/tests/CMakeLists.txt +++ b/tests/CMakeLists.txt @@ -889,6 +889,11 @@ target_include_directories(test_gdn_packed_reg_tile PRIVATE ${CMAKE_SOURCE_DIR}/ # src/vt/cuda/cuda_gdn.cu (TryTritonPackedDecode -> gdn_decode_h{48,32}). vllm_cpp_add_test(test_gdn_packed_decode_triton vt/test_gdn_packed_decode_triton.cpp) target_include_directories(test_gdn_packed_decode_triton PRIVATE ${CMAKE_SOURCE_DIR}/src) +# CPU-tier contract for the opt-in fused-GDN decode BV16 value tile. Reaches +# the CUDA-free selector/geometry header under src/; the executing kernel lives +# in cuda_gdn.cu and is covered by same-process byte-exact CUDA comparisons. +vllm_cpp_add_test(test_gdn_decode_fused vt/test_gdn_decode_fused.cpp) +target_include_directories(test_gdn_decode_fused PRIVATE ${CMAKE_SOURCE_DIR}/src) # CPU-tier contract for the two NUMERICS-NEUTRAL FP4-quant vectorized-load # fast-path flags (VT_FP4_QUANT_FAST + VT_SILU_FP4_FAST, both default OFF / # '1'-opt-in). Reaches the internal header under src/ (not public API); the diff --git a/tests/vt/test_gdn_decode_fused.cpp b/tests/vt/test_gdn_decode_fused.cpp new file mode 100644 index 000000000..468618701 --- /dev/null +++ b/tests/vt/test_gdn_decode_fused.cpp @@ -0,0 +1,281 @@ +// vllm.cpp original. Portable contract for the opt-in fused GDN decode BV16 +// value tile; the CUDA recurrence remains in cuda_gdn.cu. +#include + +#include + +#include "vt/cuda/gdn_decode_fused.h" + +using vt::cuda::DispatchGdnDecodeValueTile; +using vt::cuda::DispatchGdnDecodeStateStorage; +using vt::cuda::GdnDecodeLaunchContractFor; +using vt::cuda::GdnDecodeRegisterLogicalColumn; +using vt::cuda::GdnDecodeRegisterSharedColumn; +using vt::cuda::GdnDecodeRegstateFlagIsOn; +using vt::cuda::GdnDecodeSharedColumn; +using vt::cuda::GdnDecodeStateStride; +using vt::cuda::GdnDecodeValueTile; +using vt::cuda::GdnDecodeValueTileFromEnv; +using vt::cuda::GdnDecodeSwizzleFlagIsOn; + +static_assert(GdnDecodeRegisterSharedColumn(0, 0) == 0); +static_assert(GdnDecodeRegisterSharedColumn(7, 15) == 127); + +TEST_CASE("VT_GDN_DECODE_BV selects only exact BV16") { + CHECK(GdnDecodeValueTileFromEnv("8") == GdnDecodeValueTile::kBv32); + CHECK(GdnDecodeValueTileFromEnv("16") == GdnDecodeValueTile::kBv16); + CHECK(GdnDecodeValueTileFromEnv("24") == GdnDecodeValueTile::kBv32); + CHECK(GdnDecodeValueTileFromEnv("32") == GdnDecodeValueTile::kBv32); + constexpr const char* invalid[] = { + nullptr, "", "0", "08", "016", "024", "032", "+8", + "-8", " 8", "8 ", "8x", "16x", "24x", "32x", "64", + "on", "999999999999999999999999999999999999999999999999999999999999"}; + for (const char* value : invalid) { + CAPTURE(value == nullptr ? "" : value); + CHECK(GdnDecodeValueTileFromEnv(value) == GdnDecodeValueTile::kBv32); + } +} + +TEST_CASE("fused GDN decode value-tile geometry preserves row groups") { + const auto bv32 = GdnDecodeLaunchContractFor(nullptr, 128, 128, 8); + const auto bv16 = GdnDecodeLaunchContractFor("16", 128, 128, 8); + CHECK(bv32.value_tile == 32); + CHECK(bv32.value_tiles == 4); + CHECK(bv32.lanes_per_row == 8); + CHECK(bv32.block_threads == 256); + CHECK(bv32.shared_bytes == 17536); + CHECK(bv16.value_tile == 16); + CHECK(bv16.value_tiles == 8); + CHECK(bv16.lanes_per_row == 8); + CHECK(bv16.block_threads == 128); + CHECK(bv16.shared_bytes == 9280); + + // Removed sweep values preserve the shipped BV32 launch contract. + for (const char* removed : {"8", "24"}) { + const auto fallback = GdnDecodeLaunchContractFor(removed, 128, 128, 8); + CHECK(fallback.selected_tile == GdnDecodeValueTile::kBv32); + CHECK(fallback.value_tile == 32); + CHECK(fallback.value_tiles == 4); + CHECK(fallback.block_threads == 256); + CHECK(fallback.shared_bytes == 17536); + } + + // A partial last tile is retained rather than dropped. + CHECK(GdnDecodeLaunchContractFor(nullptr, 33, 65, 8).value_tiles == 2); + CHECK(GdnDecodeLaunchContractFor("16", 33, 65, 8).value_tiles == 3); + + // Existing corner contract: Dv<32 forces one Dk lane per value row. + const auto small32 = GdnDecodeLaunchContractFor(nullptr, 17, 65, 8); + const auto small16 = GdnDecodeLaunchContractFor("16", 17, 65, 8); + CHECK(small32.lanes_per_row == 1); + CHECK(small32.value_tile == 17); + CHECK(small32.value_tiles == 1); + CHECK(small32.block_threads == 17); + CHECK(small16.lanes_per_row == 1); + CHECK(small16.value_tile == 16); + CHECK(small16.value_tiles == 2); + CHECK(small16.block_threads == 16); + for (const auto zero : {GdnDecodeLaunchContractFor(nullptr, 0, 128, 8), + GdnDecodeLaunchContractFor("16", 128, 0, 8)}) { + CHECK_FALSE(zero.should_launch); + CHECK(zero.value_tiles == 0); + CHECK(zero.block_threads == 0); + CHECK(zero.shared_bytes == 0); + } +} + +TEST_CASE("fused GDN decode shared dispatcher invokes exactly one tile arm") { + auto select = [](const char* value) { + int bv16_calls = 0; + int bv32_calls = 0; + const auto selected = DispatchGdnDecodeValueTile( + value, 128, 128, 8, + [&](const auto& contract) { + ++bv16_calls; + CHECK(contract.value_tile == 16); + return 16; + }, + [&](const auto& contract) { + ++bv32_calls; + CHECK(contract.value_tile == 32); + return 32; + }); + CHECK(bv16_calls + bv32_calls == 1); + return selected; + }; + CHECK(select("8") == 32); + CHECK(select("16") == 16); + CHECK(select("24") == 32); + CHECK(select("32") == 32); + CHECK(select(nullptr) == 32); + CHECK(select("16x") == 32); + CHECK(select("24x") == 32); +} + +TEST_CASE("VT_GDN_DECODE_SWIZZLE selects only exact one") { + CHECK(GdnDecodeSwizzleFlagIsOn("1")); + constexpr const char* invalid[] = { + nullptr, "", "0", "01", "1 ", " 1", "+1", "1x", + "2", "on", "true", "999999999999999999999999999999999999"}; + for (const char* value : invalid) { + CAPTURE(value == nullptr ? "" : value); + CHECK_FALSE(GdnDecodeSwizzleFlagIsOn(value)); + } +} + +TEST_CASE("VT_GDN_DECODE_REGSTATE selects only exact one") { + CHECK(GdnDecodeRegstateFlagIsOn("1")); + constexpr const char* invalid[] = { + nullptr, "", "0", "01", "1 ", " 1", "+1", "1x", + "2", "on", "true", "yes", "999999999999999999999999999999999999"}; + for (const char* value : invalid) { + CAPTURE(value == nullptr ? "" : value); + CHECK_FALSE(GdnDecodeRegstateFlagIsOn(value)); + } +} + +TEST_CASE("fused GDN decode swizzle contract is production-shape only") { + const auto production = + GdnDecodeLaunchContractFor("16", "1", 128, 128, 8); + CHECK(production.swizzled); + CHECK(production.value_tile == 16); + CHECK(production.value_tiles == 8); + CHECK(production.lanes_per_row == 8); + CHECK(production.block_threads == 128); + CHECK(production.shared_bytes == 9728); + CHECK(GdnDecodeStateStride(true, 128, 8) == 136); + CHECK(GdnDecodeStateStride(false, 128, 8) == 129); + + struct FallbackCase { + vt::cuda::GdnDecodeLaunchContract contract; + size_t shared_bytes; + }; + for (const auto& fallback : { + FallbackCase{GdnDecodeLaunchContractFor("16", nullptr, 128, 128, 8), 9280}, + FallbackCase{GdnDecodeLaunchContractFor("16", "0", 128, 128, 8), 9280}, + FallbackCase{GdnDecodeLaunchContractFor("16", "1", 127, 128, 8), 9280}, + FallbackCase{GdnDecodeLaunchContractFor("16", "1", 128, 127, 8), 9208}, + FallbackCase{GdnDecodeLaunchContractFor("16", "1", 128, 128, 4), 9280}, + FallbackCase{GdnDecodeLaunchContractFor(nullptr, "1", 128, 128, 8), 17536}, + FallbackCase{GdnDecodeLaunchContractFor("32", "1", 128, 128, 8), 17536}, + }) { + CHECK_FALSE(fallback.contract.swizzled); + CHECK(fallback.contract.shared_bytes == fallback.shared_bytes); + } +} + +TEST_CASE("fused GDN decode register-state contract is production-shape only") { + const auto production = + GdnDecodeLaunchContractFor("16", "1", "1", 128, 128, 8); + CHECK(production.regstate); + CHECK(production.swizzled); + CHECK(production.value_tile == 16); + CHECK(production.value_tiles == 8); + CHECK(production.lanes_per_row == 8); + CHECK(production.block_threads == 128); + CHECK(production.shared_bytes == 9728); + + struct FallbackCase { + vt::cuda::GdnDecodeLaunchContract contract; + const char* failed_predicate; + }; + for (const auto& fallback : { + FallbackCase{GdnDecodeLaunchContractFor("16", "1", nullptr, 128, 128, 8), + "regstate unset"}, + FallbackCase{GdnDecodeLaunchContractFor("16", "1", "0", 128, 128, 8), + "regstate disabled"}, + FallbackCase{GdnDecodeLaunchContractFor("32", "1", "1", 128, 128, 8), + "BV16"}, + FallbackCase{GdnDecodeLaunchContractFor("16", "0", "1", 128, 128, 8), + "swizzle"}, + FallbackCase{GdnDecodeLaunchContractFor("16", "1", "1", 127, 128, 8), + "Dv"}, + FallbackCase{GdnDecodeLaunchContractFor("16", "1", "1", 128, 127, 8), + "Dk"}, + FallbackCase{GdnDecodeLaunchContractFor("16", "1", "1", 128, 128, 4), + "NW"}, + }) { + CAPTURE(fallback.failed_predicate); + CHECK_FALSE(fallback.contract.regstate); + } +} + +TEST_CASE("fused GDN decode register-state dispatcher invokes exactly one storage arm") { + auto select = [](const char* bv, const char* swizzle, const char* regstate, + int64_t dv, int64_t dk, int nw) { + const auto contract = + GdnDecodeLaunchContractFor(bv, swizzle, regstate, dv, dk, nw); + int shared_calls = 0; + int register_calls = 0; + const auto selected = DispatchGdnDecodeStateStorage( + contract, + [&](const auto& resolved) { + ++shared_calls; + CHECK_FALSE(resolved.regstate); + return 0; + }, + [&](const auto& resolved) { + ++register_calls; + CHECK(resolved.regstate); + return 1; + }); + CHECK(shared_calls + register_calls == 1); + return selected; + }; + + CHECK(select("16", "1", "1", 128, 128, 8) == 1); + CHECK(select("16", "1", nullptr, 128, 128, 8) == 0); + CHECK(select("16", "0", "1", 128, 128, 8) == 0); + CHECK(select("32", "1", "1", 128, 128, 8) == 0); + CHECK(select("16", "1", "1", 127, 128, 8) == 0); + CHECK(select("16", "1", "1", 128, 127, 8) == 0); + CHECK(select("16", "1", "1", 128, 128, 4) == 0); +} + +TEST_CASE("fused GDN decode swizzle is a 128-column bijection") { + CHECK(GdnDecodeSharedColumn(true, 0, 128, 8) == 0); + CHECK(GdnDecodeSharedColumn(true, 15, 128, 8) == 120); + CHECK(GdnDecodeSharedColumn(true, 16, 128, 8) == 1); + CHECK(GdnDecodeSharedColumn(true, 127, 128, 8) == 127); + CHECK(GdnDecodeSharedColumn(false, 127, 128, 8) == 127); + + bool seen[128] = {}; + for (int c = 0; c < 128; ++c) { + const int sw = static_cast(GdnDecodeSharedColumn(true, c, 128, 8)); + REQUIRE(sw >= 0); + REQUIRE(sw < 128); + CHECK_FALSE(seen[sw]); + seen[sw] = true; + } + for (int wk = 0; wk < 8; ++wk) + for (int j = 0; j < 16; ++j) + CHECK(GdnDecodeSharedColumn(true, wk * 16 + j, 128, 8) == + j * 8 + wk); +} + +TEST_CASE("fused GDN decode register slots map bijectively to logical and shared columns") { + bool logical_seen[128] = {}; + bool shared_seen[128] = {}; + for (int wk = 0; wk < 8; ++wk) { + for (int j = 0; j < 16; ++j) { + const int logical = + static_cast(GdnDecodeRegisterLogicalColumn(wk, j)); + const int shared = + static_cast(GdnDecodeRegisterSharedColumn(wk, j)); + CHECK(logical == wk * 16 + j); + CHECK(shared == j * 8 + wk); + REQUIRE(logical >= 0); + REQUIRE(logical < 128); + REQUIRE(shared >= 0); + REQUIRE(shared < 128); + CHECK_FALSE(logical_seen[logical]); + CHECK_FALSE(shared_seen[shared]); + logical_seen[logical] = true; + shared_seen[shared] = true; + } + } + for (int c = 0; c < 128; ++c) { + CHECK(logical_seen[c]); + CHECK(shared_seen[c]); + } +} diff --git a/tests/vt/test_gdn_prefill_conv.cpp b/tests/vt/test_gdn_prefill_conv.cpp index eb904fdd3..47e288d55 100644 --- a/tests/vt/test_gdn_prefill_conv.cpp +++ b/tests/vt/test_gdn_prefill_conv.cpp @@ -8,19 +8,23 @@ // (src/vt/cuda/gdn_prefill_conv.h): the VT_CONV_REG and VT_GDN_POSTCONV_SPLIT flag // predicates. The kernels themselves are CUDA-only; their BIT-EXACT (0-ulp) parity // vs the shipped kernels is a DGX-gated CUDA check (tests/vt/test_ops_gdn.cpp). This -// suite pins the portable default-ON / '0'-rollback parse so the contract is -// regression-covered on every platform, not just DGX. (Both defaults are ON because -// each kernel is bit-identical to its predecessor by construction — never-slower and -// token-safe — mirroring vLLM's register-resident FLA causal_conv1d and its per-V-head -// fused post-conv grid.) +// suite pins the portable default/rollback parse so the contract is regression-covered +// on every platform, not just CUDA. The register-window and fast-megablock paths default +// ON; the slower split and not-yet-gate-validated token tile remain opt-in. #include #include "vt/cuda/gdn_prefill_conv.h" using vt::cuda::ConvRegFlagIsOn; using vt::cuda::ConvExactChunksFlagIsOn; +using vt::cuda::ConvChannelTileArm; +using vt::cuda::ConvChannelTileArmFromEnv; +using vt::cuda::DispatchConvChannelTileLaunch; +using vt::cuda::ConvChannelTileLaunchContractFor; using vt::cuda::GdnPostConvFastFlagIsOn; using vt::cuda::GdnPostConvSplitFlagIsOn; +using vt::cuda::GdnPostConvTokenTileFlagIsOn; +using vt::cuda::GdnPostConvTokenTileGridX; TEST_CASE("VT_CONV_REG defaults ON; only a '0'-leading value rolls back") { // Default (unset) is ON: CausalConv1dFwdRegKernel's output (both `out` and the @@ -50,6 +54,83 @@ TEST_CASE("VT_CONV_EXACT_CHUNKS defaults ON; only a '0'-leading value rolls back CHECK(ConvExactChunksFlagIsOn("on")); } +TEST_CASE("VT_CONV_CHANNEL_TILE selects only the three named experiment arms") { + CHECK(ConvChannelTileArmFromEnv(nullptr) == ConvChannelTileArm::kRuntimeWidth); + CHECK(ConvChannelTileArmFromEnv("0") == ConvChannelTileArm::kRuntimeWidth); + CHECK(ConvChannelTileArmFromEnv("1") == ConvChannelTileArm::kWidthFour); + CHECK(ConvChannelTileArmFromEnv("2") == ConvChannelTileArm::kWidthFourTwoChannels); + + // Invalid spellings must preserve the sealed runtime-width baseline. + CHECK(ConvChannelTileArmFromEnv("") == ConvChannelTileArm::kRuntimeWidth); + CHECK(ConvChannelTileArmFromEnv("00") == ConvChannelTileArm::kRuntimeWidth); + CHECK(ConvChannelTileArmFromEnv("10") == ConvChannelTileArm::kRuntimeWidth); + CHECK(ConvChannelTileArmFromEnv("20") == ConvChannelTileArm::kRuntimeWidth); + CHECK(ConvChannelTileArmFromEnv("2garbage") == ConvChannelTileArm::kRuntimeWidth); + CHECK(ConvChannelTileArmFromEnv("3") == ConvChannelTileArm::kRuntimeWidth); + CHECK(ConvChannelTileArmFromEnv("on") == ConvChannelTileArm::kRuntimeWidth); + CHECK(ConvChannelTileArmFromEnv(" 2") == ConvChannelTileArm::kRuntimeWidth); +} + +TEST_CASE("causal-conv channel arms keep block 128 and tile 128/128/256 channels") { + const auto baseline = ConvChannelTileLaunchContractFor("0", 8192, 4); + const auto width_four = ConvChannelTileLaunchContractFor("1", 8192, 4); + const auto two_channels = ConvChannelTileLaunchContractFor("2", 8192, 4); + CHECK(baseline.arm == ConvChannelTileArm::kRuntimeWidth); + CHECK(width_four.arm == ConvChannelTileArm::kWidthFour); + CHECK(two_channels.arm == ConvChannelTileArm::kWidthFourTwoChannels); + CHECK(baseline.threads_per_block == 128); + CHECK(width_four.threads_per_block == 128); + CHECK(two_channels.threads_per_block == 128); + CHECK(baseline.feature_blocks == 64); + CHECK(width_four.feature_blocks == 64); + CHECK(two_channels.feature_blocks == 32); + + // Partial feature tiles round up, including the second channel stripe. + CHECK(ConvChannelTileLaunchContractFor("0", 129, 4).feature_blocks == 2); + CHECK(ConvChannelTileLaunchContractFor("1", 129, 4).feature_blocks == 2); + CHECK(ConvChannelTileLaunchContractFor("2", 129, 4).feature_blocks == 1); + CHECK(ConvChannelTileLaunchContractFor("2", 257, 4).feature_blocks == 2); + + // Width-specialized arms are not valid for any other convolution width. + CHECK(ConvChannelTileLaunchContractFor("1", 8192, 3).arm == + ConvChannelTileArm::kRuntimeWidth); + const auto unsupported_two_channels = + ConvChannelTileLaunchContractFor("2", 8192, 5); + CHECK(unsupported_two_channels.arm == ConvChannelTileArm::kRuntimeWidth); + CHECK(unsupported_two_channels.feature_blocks == 64); +} + +TEST_CASE("causal-conv shared channel dispatcher invokes each arm and fallback") { + auto selected = [](const char* env_value, int64_t channels, + int64_t kernel_width) { + int runtime_calls = 0; + int width_four_calls = 0; + int two_channel_calls = 0; + const auto result = DispatchConvChannelTileLaunch( + env_value, channels, kernel_width, + [&](const auto&) { + ++runtime_calls; + return ConvChannelTileArm::kRuntimeWidth; + }, + [&](const auto&) { + ++width_four_calls; + return ConvChannelTileArm::kWidthFour; + }, + [&](const auto&) { + ++two_channel_calls; + return ConvChannelTileArm::kWidthFourTwoChannels; + }); + CHECK(runtime_calls + width_four_calls + two_channel_calls == 1); + return result; + }; + + CHECK(selected("0", 8192, 4) == ConvChannelTileArm::kRuntimeWidth); + CHECK(selected("1", 8192, 4) == ConvChannelTileArm::kWidthFour); + CHECK(selected("2", 8192, 4) == ConvChannelTileArm::kWidthFourTwoChannels); + CHECK(selected("2garbage", 8192, 4) == ConvChannelTileArm::kRuntimeWidth); + CHECK(selected("2", 8192, 5) == ConvChannelTileArm::kRuntimeWidth); +} + TEST_CASE("VT_GDN_POSTCONV_SPLIT defaults OFF (opt-in); a non-'0' value enables it") { // Default (unset) is OFF: GdnPostConvSplitKernel is BIT-IDENTICAL (0-ulp) to the // shipped GdnPostConvKernel by construction (byte-for-byte q/k L2-norm branch; same @@ -85,3 +166,25 @@ TEST_CASE("VT_GDN_POSTCONV_FAST defaults ON; only a '0'-leading value rolls back CHECK_FALSE(GdnPostConvFastFlagIsOn("0abc")); CHECK_FALSE(GdnPostConvFastFlagIsOn("00")); } + +TEST_CASE("VT_GDN_POSTCONV_TOKEN_TILE defaults OFF; a non-'0' value enables it") { + CHECK_FALSE(GdnPostConvTokenTileFlagIsOn(nullptr)); + CHECK_FALSE(GdnPostConvTokenTileFlagIsOn("0")); + CHECK_FALSE(GdnPostConvTokenTileFlagIsOn("0abc")); + CHECK_FALSE(GdnPostConvTokenTileFlagIsOn("00")); + CHECK(GdnPostConvTokenTileFlagIsOn("")); + CHECK(GdnPostConvTokenTileFlagIsOn("1")); + CHECK(GdnPostConvTokenTileFlagIsOn("on")); + CHECK(GdnPostConvTokenTileFlagIsOn(" 0")); +} + +TEST_CASE("GDN post-conv token tile covers each ceil(T/16) work item") { + CHECK(GdnPostConvTokenTileGridX(0) == 0); + CHECK(GdnPostConvTokenTileGridX(1) == 1); + CHECK(GdnPostConvTokenTileGridX(15) == 1); + CHECK(GdnPostConvTokenTileGridX(16) == 1); + CHECK(GdnPostConvTokenTileGridX(17) == 2); + CHECK(GdnPostConvTokenTileGridX(127) == 8); + CHECK(GdnPostConvTokenTileGridX(128) == 8); + CHECK(GdnPostConvTokenTileGridX(2048) == 128); +} diff --git a/tests/vt/test_ops_gdn.cpp b/tests/vt/test_ops_gdn.cpp index f952b20d5..f624887c8 100644 --- a/tests/vt/test_ops_gdn.cpp +++ b/tests/vt/test_ops_gdn.cpp @@ -27,6 +27,8 @@ #include "vt/dtype.h" #include "vt/ops.h" #ifdef VLLM_CPP_CUDA +#include + // Packed-decode debug counters (reg-tile vs legacy launch sub-counts) plus the // Triton AOT stats used by the chunked cases. Defined in cuda_gdn.cu, available // whenever the CUDA backend is built (the packed counters do not require Triton). @@ -707,6 +709,25 @@ std::vector Unpack(const std::vector& b, DType dt) { return out; } +void CheckBytesEqual(const std::vector& got, + const std::vector& want) { + REQUIRE(got.size() == want.size()); + size_t bad = 0; + size_t first_bad = 0; + for (size_t i = 0; i < got.size(); ++i) { + if (got[i] != want[i]) { + if (bad == 0) first_bad = i; + ++bad; + } + } + if (bad != 0) { + CAPTURE(first_bad); + CAPTURE(got[first_bad]); + CAPTURE(want[first_bad]); + } + CHECK(bad == 0); +} + void CheckClose(const std::vector& got, const std::vector& want, float atol, float rtol) { REQUIRE(got.size() == want.size()); @@ -1843,6 +1864,120 @@ void RunGdnCudaCase(const std::vector& qsl, int64_t batch, CheckClose(st_gpu, st_cpu, state_tol, state_tol); } +struct GdnDecodeTestEnv { + const char* bv; + const char* swizzle; + const char* regstate; + const char* nw; +}; + +void SetOrUnsetEnv(const char* name, const char* value) { + if (value == nullptr) + unsetenv(name); + else + setenv(name, value, 1); +} + +// Same-process schedule discriminator for the regular fused GDN decode path. +// Output and the complete persistent state must remain byte-identical. +void RunGdnDecodeExactCase(int64_t dv, DType io_dtype, DType state_dtype, + bool indexed, const GdnDecodeTestEnv& incumbent, + const GdnDecodeTestEnv& candidate, uint32_t seed) { + const int64_t n = indexed ? 3 : 2; + const int64_t slots = indexed ? 4 : n; + const int64_t hk = 1, hv = 2, dk = 128; + const GdnArgs args{1.0f / std::sqrt(static_cast(dk))}; + const auto q = Pack(RandomF32(static_cast(n * hk * dk), seed, + -0.5f, 0.5f), + io_dtype); + const auto k = Pack(RandomF32(static_cast(n * hk * dk), seed + 1, + -0.5f, 0.5f), + io_dtype); + const auto v = Pack(RandomF32(static_cast(n * hv * dv), seed + 2, + -0.75f, 0.75f), + io_dtype); + const auto g = RandomF32(static_cast(n * hv), seed + 3, -0.5f, + -0.01f); + const auto beta = RandomF32(static_cast(n * hv), seed + 4, 0.1f, + 0.9f); + const auto state_initial = + Pack(RandomF32(static_cast(slots * hv * dv * dk), seed + 5, + -0.25f, 0.25f), + state_dtype); + const std::vector state_idx = {2, -1, 0}; + + Backend& gpu = vt::GetBackend(DeviceType::kCUDA); + QueueGuard gq(gpu); + auto run = [&](const GdnDecodeTestEnv& env, std::vector& out_bytes, + std::vector& state_bytes) { + SetOrUnsetEnv("VT_GDN_DECODE_NW", env.nw); + SetOrUnsetEnv("VT_GDN_DECODE_BV", env.bv); + SetOrUnsetEnv("VT_GDN_DECODE_SWIZZLE", env.swizzle); + SetOrUnsetEnv("VT_GDN_DECODE_REGSTATE", env.regstate); + DeviceTensor dq(gpu, gq.q, io_dtype, {n, hk, dk}, q.data()); + DeviceTensor dkt(gpu, gq.q, io_dtype, {n, hk, dk}, k.data()); + DeviceTensor dvt(gpu, gq.q, io_dtype, {n, hv, dv}, v.data()); + DeviceTensor dg(gpu, gq.q, DType::kF32, {n, hv}, g.data()); + DeviceTensor dbeta(gpu, gq.q, DType::kF32, {n, hv}, beta.data()); + DeviceTensor dst(gpu, gq.q, state_dtype, {slots, hv, dv, dk}, + state_initial.data()); + DeviceTensor dout(gpu, gq.q, io_dtype, {n, hv, dv}); + if (indexed) { + DeviceTensor didx(gpu, gq.q, DType::kI32, {n}, state_idx.data()); + vt::GdnDecode(gq.q, dout.tensor(), dq.tensor(), dkt.tensor(), + dvt.tensor(), dg.tensor(), dbeta.tensor(), dst.tensor(), + args, &didx.tensor()); + } else { + vt::GdnDecode(gq.q, dout.tensor(), dq.tensor(), dkt.tensor(), + dvt.tensor(), dg.tensor(), dbeta.tensor(), dst.tensor(), + args); + } + out_bytes.resize(static_cast(n * hv * dv) * + vt::SizeOf(io_dtype)); + state_bytes.resize(state_initial.size()); + dout.Download(gq.q, out_bytes.data()); + dst.Download(gq.q, state_bytes.data()); + }; + + std::vector out_incumbent, state_incumbent; + std::vector out_candidate, state_candidate; + run(incumbent, out_incumbent, state_incumbent); + run(candidate, out_candidate, state_candidate); + unsetenv("VT_GDN_DECODE_REGSTATE"); + unsetenv("VT_GDN_DECODE_SWIZZLE"); + unsetenv("VT_GDN_DECODE_BV"); + unsetenv("VT_GDN_DECODE_NW"); + CheckBytesEqual(out_candidate, out_incumbent); + CheckBytesEqual(state_candidate, state_incumbent); + + if (indexed) { + // The null request must write exact zero outputs and the unreferenced cache + // slot must remain byte-for-byte untouched in both schedules. + const size_t out_row_bytes = + static_cast(hv * dv) * vt::SizeOf(io_dtype); + CHECK(std::all_of(out_candidate.begin() + static_cast(out_row_bytes), + out_candidate.begin() + static_cast(2 * out_row_bytes), + [](uint8_t byte) { return byte == 0; })); + const size_t state_row_bytes = + static_cast(hv * dv * dk) * vt::SizeOf(state_dtype); + const auto untouched_begin = + state_candidate.begin() + static_cast(state_row_bytes); + CHECK(std::equal(untouched_begin, + untouched_begin + static_cast(state_row_bytes), + state_initial.begin() + + static_cast(state_row_bytes))); + } +} + +#ifdef VLLM_CPP_CUDA +struct CudaGraphGuard { + cudaGraph_t graph = nullptr; + ~CudaGraphGuard() { + if (graph != nullptr) cudaGraphDestroy(graph); + } +}; +#endif + // Chunk-parallel prefill (VT_GDN_CHUNKED, default) vs the sequential scan // (VT_GDN_CHUNKED=0), same inputs, same binary. The chunked path is the M2 // prefill perf kernel (cuda_gdn.cu GdnChunk*); it must reproduce the sequential @@ -2392,14 +2527,28 @@ TEST_CASE("CUDA causal_conv1d_update decode-fast (VT_CONV_UPDATE_FAST) matches r // ("1" reg / "0" tiled; launcher reads getenv per call) so it is default-independent. void RunConvFwdRegByteExactCase(const std::vector& qsl, const std::vector& his, int64_t c, int64_t k, bool with_bias, bool silu, const Combo& cb, - uint32_t seed, bool i8_mask = false) { + uint32_t seed, bool i8_mask = false, + int64_t row_padding = 0) { + CAPTURE(c); + CAPTURE(k); + CAPTURE(static_cast(cb.in)); + CAPTURE(static_cast(cb.out)); + CAPTURE(row_padding); const int64_t n = static_cast(qsl.size()) - 1; const int64_t t = qsl.back(); const auto xf = RandomF32(static_cast(t * c), seed, -3.0f, 3.0f); const auto wf = RandomF32(static_cast(c * k), seed + 1, -1.0f, 1.0f); const auto bf = RandomF32(static_cast(c), seed + 2, -1.0f, 1.0f); const auto stf = RandomF32(static_cast(n * c * (k - 1)), seed + 3, -2.0f, 2.0f); - const auto xb = Pack(xf, cb.in); + const int64_t x_row_stride = c + row_padding; + std::vector x_storage(static_cast(t * x_row_stride), -12345.0f); + for (int64_t token = 0; token < t; ++token) { + for (int64_t channel = 0; channel < c; ++channel) { + x_storage[static_cast(token * x_row_stride + channel)] = + xf[static_cast(token * c + channel)]; + } + } + const auto xb = Pack(x_storage, cb.in); const auto wb = Pack(wf, cb.in); const auto bb = Pack(bf, cb.in); const auto stb = Pack(stf, DType::kF32); @@ -2408,7 +2557,8 @@ void RunConvFwdRegByteExactCase(const std::vector& qsl, const std::vect Backend& gpu = vt::GetBackend(DeviceType::kCUDA); QueueGuard gq(gpu); - DeviceTensor dx(gpu, gq.q, cb.in, {t, c}, xb.data()); + DeviceTensor dx(gpu, gq.q, cb.in, {t, x_row_stride}, xb.data()); + Tensor x_view = RowView(dx.tensor().data, cb.in, Gpu(), t, c, x_row_stride); DeviceTensor dw(gpu, gq.q, cb.in, {c, k}, wb.data()); DeviceTensor db(gpu, gq.q, cb.in, {c}, bb.data()); DeviceTensor dqsl(gpu, gq.q, DType::kI32, {n + 1}, qsl.data()); @@ -2434,10 +2584,12 @@ void RunConvFwdRegByteExactCase(const std::vector& qsl, const std::vect {static_cast(chunk_offsets.size())}, chunk_offsets.data()); - auto run = [&](bool reg, bool exact, std::vector& out_bytes, + auto run = [&](bool reg, bool exact, const char* channel_arm, + std::vector& out_bytes, std::vector& st_bytes) { ::setenv("VT_CONV_REG", reg ? "1" : "0", 1); ::setenv("VT_CONV_EXACT_CHUNKS", exact ? "1" : "0", 1); + ::setenv("VT_CONV_CHANNEL_TILE", channel_arm, 1); DeviceTensor dst(gpu, gq.q, DType::kF32, {n, c, k - 1}, stb.data()); // fresh state per arm DeviceTensor dout(gpu, gq.q, cb.out, {t, c}); gpu.Memset(gq.q, dout.tensor().data, 0x5a, static_cast(t * c) * vt::SizeOf(cb.out)); @@ -2448,7 +2600,7 @@ void RunConvFwdRegByteExactCase(const std::vector& qsl, const std::vect run_args.batch_ptr = &batch_tensor; run_args.token_chunk_offset_ptr = &offsets_tensor; } - vt::CausalConv1dFwd(gq.q, dout.tensor(), dx.tensor(), dw.tensor(), + vt::CausalConv1dFwd(gq.q, dout.tensor(), x_view, dw.tensor(), with_bias ? &db.tensor() : nullptr, dst.tensor(), dqsl.tensor(), dhis.tensor(), exact ? run_args : args); out_bytes.resize(static_cast(t * c) * vt::SizeOf(cb.out)); @@ -2457,15 +2609,24 @@ void RunConvFwdRegByteExactCase(const std::vector& qsl, const std::vect dst.Download(gq.q, st_bytes.data()); }; std::vector out_tiled, st_tiled, out_reg, st_reg, out_exact, st_exact; - run(/*reg=*/false, /*exact=*/false, out_tiled, st_tiled); - run(/*reg=*/true, /*exact=*/false, out_reg, st_reg); - run(/*reg=*/true, /*exact=*/true, out_exact, st_exact); + std::vector out_width_four, st_width_four, out_two_channels, + st_two_channels; + run(/*reg=*/false, /*exact=*/false, "0", out_tiled, st_tiled); + run(/*reg=*/true, /*exact=*/false, "0", out_reg, st_reg); + run(/*reg=*/true, /*exact=*/true, "0", out_exact, st_exact); + run(/*reg=*/true, /*exact=*/true, "1", out_width_four, st_width_four); + run(/*reg=*/true, /*exact=*/true, "2", out_two_channels, st_two_channels); ::unsetenv("VT_CONV_REG"); ::unsetenv("VT_CONV_EXACT_CHUNKS"); + ::unsetenv("VT_CONV_CHANNEL_TILE"); CHECK(out_reg == out_tiled); // out activation byte-identical CHECK(st_reg == st_tiled); // rolled conv_state byte-identical CHECK(out_exact == out_tiled); // exact descriptor changes only work assignment CHECK(st_exact == st_tiled); + CheckBytesEqual(out_width_four, out_exact); // compile-time K=4, one channel per lane + CheckBytesEqual(st_width_four, st_exact); + CheckBytesEqual(out_two_channels, out_exact); // compile-time K=4, two channels per lane + CheckBytesEqual(st_two_channels, st_exact); } TEST_CASE("CUDA causal_conv1d_fwd register kernel (VT_CONV_REG) matches tiled 0-ulp") { @@ -2495,6 +2656,11 @@ TEST_CASE("CUDA causal_conv1d_fwd register kernel (VT_CONV_REG) matches tiled 0- // i8 has_initial_state mask, no bias, silu. RunConvFwdRegByteExactCase({0, 33, 70}, {1, 0}, 1024, 4, false, true, kCudaCombos[0], seed + 30, /*i8_mask=*/true); + // Partial 256-channel tile plus the production packed-row stride. Unequal exact + // chunks cover fresh and initial state and leave the parent-row padding unread. + RunConvFwdRegByteExactCase({0, 2, 19, 28}, {0, 1, 0}, 385, 4, true, true, + kCudaCombos[2], seed + 40, /*i8_mask=*/false, + /*row_padding=*/37); } // VT_GDN_POSTCONV_SPLIT: the per-V-head split post-conv kernel (GdnPostConvSplitKernel) @@ -2526,6 +2692,7 @@ void RunGdnPostConvSplitByteExactCase(int64_t t, int64_t hk, int64_t hv, int64_t auto run = [&](bool split, std::vector& q, std::vector& k, std::vector& v, std::vector& g, std::vector& b) { ::setenv("VT_GDN_POSTCONV_SPLIT", split ? "1" : "0", 1); + ::setenv("VT_GDN_POSTCONV_TOKEN_TILE", "0", 1); DeviceTensor dq_(gpu, gq.q, qkv_dt, {t, hk, dk}); DeviceTensor dk_(gpu, gq.q, qkv_dt, {t, hk, dk}); DeviceTensor dv_(gpu, gq.q, qkv_dt, {t, hv, dv}); @@ -2549,6 +2716,7 @@ void RunGdnPostConvSplitByteExactCase(int64_t t, int64_t hk, int64_t hv, int64_t run(/*split=*/false, q0, k0, v0, g0, b0); run(/*split=*/true, q1, k1, v1, g1, b1); ::unsetenv("VT_GDN_POSTCONV_SPLIT"); + ::unsetenv("VT_GDN_POSTCONV_TOKEN_TILE"); CHECK(q1 == q0); CHECK(k1 == k0); CHECK(v1 == v0); @@ -2579,6 +2747,115 @@ TEST_CASE("CUDA gdn_post_conv split kernel (VT_GDN_POSTCONV_SPLIT) matches megab } } +// Port of vLLM tests/kernels/test_fused_gdn_post_conv.py's BLOCK_T=16 +// correctness sweep for the production 128-wide head. The experimental CUDA +// tile folds the first two levels of the 128-lane reduction into each warp lane +// in the same arithmetic order, so all five outputs must remain byte-exact +// against the shipped fast kernel (stronger than upstream's BF16 tolerance). +void RunGdnPostConvTokenTileCase(int64_t t, int64_t hk, int64_t hv, uint32_t seed) { + constexpr int64_t dk = 128; + constexpr int64_t dv = 128; + const int64_t key_dim = hk * dk; + const int64_t value_dim = hv * dv; + const int64_t conv_dim = 2 * key_dim + value_dim; + const auto convf = RandomF32(static_cast(t * conv_dim), seed, -1.5f, 1.5f); + const auto araw = RandomF32(static_cast(t * hv), seed + 1, -1.0f, 1.0f); + const auto braw = RandomF32(static_cast(t * hv), seed + 2, -1.0f, 1.0f); + const auto alog = RandomF32(static_cast(hv), seed + 3, -1.0f, 1.0f); + const auto dtb = RandomF32(static_cast(hv), seed + 4, -1.0f, 1.0f); + const auto convb = Pack(convf, DType::kBF16); + constexpr int64_t gate_prefix = 2; + constexpr int64_t gate_suffix = 3; + const int64_t gate_stride = gate_prefix + 2 * hv + gate_suffix; + std::vector packed_ba(static_cast(t * gate_stride), 7.5f); + for (int64_t row = 0; row < t; ++row) { + for (int64_t head = 0; head < hv; ++head) { + packed_ba[static_cast(row * gate_stride + gate_prefix + head)] = + braw[static_cast(row * hv + head)]; + packed_ba[static_cast(row * gate_stride + gate_prefix + hv + head)] = + araw[static_cast(row * hv + head)]; + } + } + const vt::L2NormArgs args{1e-6f}; + + Backend& gpu = vt::GetBackend(DeviceType::kCUDA); + QueueGuard gq(gpu); + DeviceTensor dconv(gpu, gq.q, DType::kBF16, {t, conv_dim}, convb.data()); + DeviceTensor dpacked_ba(gpu, gq.q, DType::kF32, {t, gate_stride}, packed_ba.data()); + Tensor dbraw = dpacked_ba.tensor().Slice(1, gate_prefix, gate_prefix + hv); + Tensor daraw = dpacked_ba.tensor().Slice(1, gate_prefix + hv, gate_prefix + 2 * hv); + DeviceTensor dalog(gpu, gq.q, DType::kF32, {hv}, alog.data()); + DeviceTensor ddtb(gpu, gq.q, DType::kF32, {hv}, dtb.data()); + + auto run = [&](bool token_tile, std::vector& q, std::vector& k, + std::vector& v, std::vector& g, + std::vector& b) { + ::setenv("VT_GDN_POSTCONV_SPLIT", "0", 1); + ::setenv("VT_GDN_POSTCONV_FAST", "1", 1); + ::setenv("VT_GDN_POSTCONV_TOKEN_TILE", token_tile ? "1" : "0", 1); + DeviceTensor dq_(gpu, gq.q, DType::kBF16, {t, hk, dk}); + DeviceTensor dk_(gpu, gq.q, DType::kBF16, {t, hk, dk}); + DeviceTensor dv_(gpu, gq.q, DType::kBF16, {t, hv, dv}); + DeviceTensor dg_(gpu, gq.q, DType::kF32, {t, hv}); + DeviceTensor db_(gpu, gq.q, DType::kF32, {t, hv}); + vt::GdnPostConv(gq.q, dq_.tensor(), dk_.tensor(), dv_.tensor(), dg_.tensor(), db_.tensor(), + dconv.tensor(), daraw, dbraw, dalog.tensor(), ddtb.tensor(), args); + q.resize(static_cast(t * key_dim) * vt::SizeOf(DType::kBF16)); + k.resize(q.size()); + v.resize(static_cast(t * value_dim) * vt::SizeOf(DType::kBF16)); + g.resize(static_cast(t * hv) * vt::SizeOf(DType::kF32)); + b.resize(g.size()); + dq_.Download(gq.q, q.data()); + dk_.Download(gq.q, k.data()); + dv_.Download(gq.q, v.data()); + dg_.Download(gq.q, g.data()); + db_.Download(gq.q, b.data()); + }; + + std::vector q0, k0, v0, g0, b0, q1, k1, v1, g1, b1; + run(/*token_tile=*/false, q0, k0, v0, g0, b0); + run(/*token_tile=*/true, q1, k1, v1, g1, b1); + ::unsetenv("VT_GDN_POSTCONV_SPLIT"); + ::unsetenv("VT_GDN_POSTCONV_FAST"); + ::unsetenv("VT_GDN_POSTCONV_TOKEN_TILE"); + + const std::vector q1f = Unpack(q1, DType::kBF16); + const std::vector k1f = Unpack(k1, DType::kBF16); + CHECK(q1 == q0); + CHECK(k1 == k0); + CHECK(v1 == v0); + CHECK(g1 == g0); + CHECK(b1 == b0); + for (const std::vector* values : {&q1f, &k1f}) { + bool all_finite = true; + float max_norm_error = 0.0f; + for (int64_t row = 0; row < t * hk; ++row) { + float sum = 0.0f; + for (int64_t j = 0; j < dk; ++j) { + const float value = (*values)[static_cast(row * dk + j)]; + all_finite = all_finite && std::isfinite(value); + sum += value * value; + } + max_norm_error = std::max(max_norm_error, std::abs(std::sqrt(sum) - 1.0f)); + } + CHECK(all_finite); + CHECK(max_norm_error < 1e-2f); + } +} + +TEST_CASE("CUDA gdn_post_conv 16-token tile matches upstream BF16 contract") { + if (!HasCuda()) { + MESSAGE("no CUDA backend registered; skipping"); + return; + } + uint32_t seed = 9900; + for (int64_t t : {1, 16, 17, 128, 512, 2048}) { + CAPTURE(t); + RunGdnPostConvTokenTileCase(t, 4, 8, seed++); + } + RunGdnPostConvTokenTileCase(17, 16, 32, seed); // upstream 35B shape +} + TEST_CASE("CUDA l2norm matches CPU (rank 2 and 3)") { if (!HasCuda()) { MESSAGE("no CUDA backend registered; skipping"); @@ -3212,6 +3489,165 @@ TEST_CASE("CUDA gdn decode matches CPU (real dims, batched multi-seq)") { RunGdnCudaCase({}, 5, 2, 2, 128, 128, kCudaCombos[0], 6320); // GQA ratio 1, odd batch } +TEST_CASE("CUDA fused GDN decode swizzle is byte-exact") { + if (!HasCuda()) { + MESSAGE("no CUDA backend registered; skipping"); + return; + } + constexpr GdnDecodeTestEnv linear{"16", "0", "0", "8"}; + constexpr GdnDecodeTestEnv swizzled{"16", "1", "0", "8"}; + RunGdnDecodeExactCase(128, DType::kBF16, DType::kBF16, false, + linear, swizzled, 6350); + RunGdnDecodeExactCase(128, DType::kF32, DType::kF32, true, + linear, swizzled, 6360); + RunGdnDecodeExactCase(37, DType::kBF16, DType::kF32, true, + linear, swizzled, 6370); +} + +TEST_CASE("CUDA fused GDN decode register-state schedule is byte-exact") { + if (!HasCuda()) { + MESSAGE("no CUDA backend registered; skipping"); + return; + } + constexpr GdnDecodeTestEnv shared{"16", "1", "0", "8"}; + constexpr GdnDecodeTestEnv regstate{"16", "1", "1", "8"}; + RunGdnDecodeExactCase(128, DType::kBF16, DType::kBF16, false, + shared, regstate, 6380); + RunGdnDecodeExactCase(128, DType::kBF16, DType::kF16, true, + shared, regstate, 6390); + RunGdnDecodeExactCase(128, DType::kF32, DType::kF32, true, + shared, regstate, 6400); + + // REGSTATE must preserve incumbent kernels whenever one production predicate + // fails. These cover partial Dv, unswizzled BV16, BV32, and NW4. + constexpr GdnDecodeTestEnv partial_shared{"16", "1", "0", "8"}; + constexpr GdnDecodeTestEnv partial_reg{"16", "1", "1", "8"}; + RunGdnDecodeExactCase(37, DType::kBF16, DType::kF32, true, + partial_shared, partial_reg, 6410); + constexpr GdnDecodeTestEnv linear_shared{"16", "0", "0", "8"}; + constexpr GdnDecodeTestEnv linear_reg{"16", "0", "1", "8"}; + RunGdnDecodeExactCase(128, DType::kBF16, DType::kBF16, false, + linear_shared, linear_reg, 6420); + constexpr GdnDecodeTestEnv bv32_shared{"32", "1", "0", "8"}; + constexpr GdnDecodeTestEnv bv32_reg{"32", "1", "1", "8"}; + RunGdnDecodeExactCase(128, DType::kF32, DType::kF32, false, + bv32_shared, bv32_reg, 6430); + constexpr GdnDecodeTestEnv nw4_shared{"16", "1", "0", "4"}; + constexpr GdnDecodeTestEnv nw4_reg{"16", "1", "1", "4"}; + RunGdnDecodeExactCase(128, DType::kBF16, DType::kF32, false, + nw4_shared, nw4_reg, 6440); +} + +TEST_CASE("CUDA public GdnDecode binds decode layout to exact production geometry") { + if (!HasCuda()) { + MESSAGE("no CUDA backend registered; skipping"); + return; + } +#ifdef VLLM_CPP_CUDA + constexpr int64_t n = 2, hk = 1, hv = 2, dk = 128, dv = 128; + const GdnArgs args{1.0f / std::sqrt(static_cast(dk))}; + const auto q = Pack(RandomF32(static_cast(n * hk * dk), 6430, + -0.5f, 0.5f), + DType::kBF16); + const auto k = Pack(RandomF32(static_cast(n * hk * dk), 6431, + -0.5f, 0.5f), + DType::kBF16); + const auto v = Pack(RandomF32(static_cast(n * hv * dv), 6432, + -0.75f, 0.75f), + DType::kBF16); + const auto g = RandomF32(static_cast(n * hv), 6433, -0.5f, -0.01f); + const auto beta = RandomF32(static_cast(n * hv), 6434, 0.1f, 0.9f); + const auto state = RandomF32(static_cast(n * hv * dv * dk), 6435, + -0.25f, 0.25f); + + Backend& gpu = vt::GetBackend(DeviceType::kCUDA); + QueueGuard gq(gpu); + // All device and host allocations precede capture. The captured work is the + // actual public op call on the queue's production stream. + DeviceTensor dq(gpu, gq.q, DType::kBF16, {n, hk, dk}, q.data()); + DeviceTensor dkt(gpu, gq.q, DType::kBF16, {n, hk, dk}, k.data()); + DeviceTensor dvt(gpu, gq.q, DType::kBF16, {n, hv, dv}, v.data()); + DeviceTensor dg(gpu, gq.q, DType::kF32, {n, hv}, g.data()); + DeviceTensor dbeta(gpu, gq.q, DType::kF32, {n, hv}, beta.data()); + DeviceTensor dst(gpu, gq.q, DType::kF32, {n, hv, dv, dk}, state.data()); + DeviceTensor dout(gpu, gq.q, DType::kBF16, {n, hv, dv}); + gpu.Synchronize(gq.q); + + struct ExpectedGeometry { + const char* bv_env; + const char* swizzle_env; + const char* regstate_env; + bool register_state; + unsigned grid_x; + unsigned block_x; + unsigned shared_bytes; + }; + constexpr ExpectedGeometry cases[] = { + {nullptr, "1", "1", false, 4, 256, 17536}, + {"16", "0", "1", false, 8, 128, 9280}, + {"16", "1", "0", false, 8, 128, 9728}, + {"16", "1", "1", true, 8, 128, 9728}, + {"16", "1", "1x", false, 8, 128, 9728}, + {"16", "1x", "1", false, 8, 128, 9280}, + {"32", "1", "1", false, 4, 256, 17536}, + }; + const cudaStream_t stream = static_cast(gq.q.handle); + setenv("VT_GDN_DECODE_NW", "8", 1); + void* shared_production_function = nullptr; + for (const auto& expected : cases) { + CAPTURE(expected.bv_env == nullptr ? "" : expected.bv_env); + CAPTURE(expected.swizzle_env); + CAPTURE(expected.regstate_env); + if (expected.bv_env == nullptr) + unsetenv("VT_GDN_DECODE_BV"); + else + setenv("VT_GDN_DECODE_BV", expected.bv_env, 1); + setenv("VT_GDN_DECODE_SWIZZLE", expected.swizzle_env, 1); + setenv("VT_GDN_DECODE_REGSTATE", expected.regstate_env, 1); + CudaGraphGuard captured; + REQUIRE(cudaStreamBeginCapture(stream, cudaStreamCaptureModeThreadLocal) == + cudaSuccess); + vt::GdnDecode(gq.q, dout.tensor(), dq.tensor(), dkt.tensor(), dvt.tensor(), + dg.tensor(), dbeta.tensor(), dst.tensor(), args); + REQUIRE(cudaStreamEndCapture(stream, &captured.graph) == cudaSuccess); + REQUIRE(captured.graph != nullptr); + + size_t node_count = 0; + REQUIRE(cudaGraphGetNodes(captured.graph, nullptr, &node_count) == cudaSuccess); + REQUIRE(node_count == 1); + cudaGraphNode_t node = nullptr; + REQUIRE(cudaGraphGetNodes(captured.graph, &node, &node_count) == cudaSuccess); + REQUIRE(node_count == 1); + cudaGraphNodeType node_type = cudaGraphNodeTypeEmpty; + REQUIRE(cudaGraphNodeGetType(node, &node_type) == cudaSuccess); + REQUIRE(node_type == cudaGraphNodeTypeKernel); + cudaKernelNodeParams params{}; + REQUIRE(cudaGraphKernelNodeGetParams(node, ¶ms) == cudaSuccess); + if (expected.bv_env != nullptr && + std::string(expected.bv_env) == "16" && + std::string(expected.swizzle_env) == "1" && + std::string(expected.regstate_env) == "0") { + shared_production_function = params.func; + } + if (expected.register_state) { + REQUIRE(shared_production_function != nullptr); + CHECK(params.func != shared_production_function); + } + CHECK(params.gridDim.x == expected.grid_x); + CHECK(params.gridDim.y == static_cast(n * hv)); + CHECK(params.gridDim.z == 1); + CHECK(params.blockDim.x == expected.block_x); + CHECK(params.blockDim.y == 1); + CHECK(params.blockDim.z == 1); + CHECK(params.sharedMemBytes == expected.shared_bytes); + } + unsetenv("VT_GDN_DECODE_REGSTATE"); + unsetenv("VT_GDN_DECODE_SWIZZLE"); + unsetenv("VT_GDN_DECODE_BV"); + unsetenv("VT_GDN_DECODE_NW"); +#endif +} + // vLLM's public mamba_ssm_cache_dtype accepts float16 independently of the // model activation dtype. Preserve the decomposed rollback path for that // configuration as well as the packed path; otherwise disabling packed decode From 41b0095efe9ac1244bed7f33a7636a3a3d58d41e Mon Sep 17 00:00:00 2001 From: Richard Palethorpe Date: Sun, 9 Aug 2026 19:18:44 +0100 Subject: [PATCH 03/17] test(cuda): guard post-conv token-tile fallback (#206) FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: Codex:gpt-5.6-sol [Codex] --- src/vt/cuda/cuda_gdn.cu | 5 ++--- src/vt/cuda/gdn_prefill_conv.h | 11 +++++++++++ tests/vt/test_gdn_prefill_conv.cpp | 12 ++++++++++++ 3 files changed, 25 insertions(+), 3 deletions(-) diff --git a/src/vt/cuda/cuda_gdn.cu b/src/vt/cuda/cuda_gdn.cu index fa65ba3e3..0ad4eff66 100644 --- a/src/vt/cuda/cuda_gdn.cu +++ b/src/vt/cuda/cuda_gdn.cu @@ -1759,9 +1759,8 @@ void GdnPostConvKernelCuda(Queue& q, Tensor& q_out, Tensor& k_out, Tensor& v_out // VT_GDN_POSTCONV_TOKEN_TILE (opt-in): the full upstream 16-token, // per-head/four-warp schedule. The explicit split control keeps priority when // both experimental flags are set. Only the production 128-wide heads route. - const bool token_tile = - !split && GdnPostConvTokenTileFlagIsOn(std::getenv("VT_GDN_POSTCONV_TOKEN_TILE")) && - dk == 128 && dv == 128; + const bool token_tile = GdnPostConvTokenTileEligible( + split, std::getenv("VT_GDN_POSTCONV_TOKEN_TILE"), dk, dv); // VT_GDN_POSTCONV_FAST: byte-identical megablock at 128 threads + 128-bit V copy. // Only for the Dk==Dv==128 gate dims (16B alignment + value_dim%8==0); mutually // exclusive with the split (both target the same megablock). See predicate. diff --git a/src/vt/cuda/gdn_prefill_conv.h b/src/vt/cuda/gdn_prefill_conv.h index 3510ee4f4..b0df230e9 100644 --- a/src/vt/cuda/gdn_prefill_conv.h +++ b/src/vt/cuda/gdn_prefill_conv.h @@ -203,6 +203,17 @@ inline bool GdnPostConvTokenTileFlagIsOn(const char* env_value) { return env_value != nullptr && env_value[0] != '0'; } +// Shared production eligibility contract for the token-tile kernel. The +// explicit split experiment retains priority, and the token-tile indexing owns +// exactly four 32-element feature groups per lane, so both Q/K and V head +// widths must independently be 128. Unsupported shapes keep the existing +// megablock/split dispatch. +inline bool GdnPostConvTokenTileEligible(bool split, const char* env_value, + int64_t dk, int64_t dv) { + return !split && GdnPostConvTokenTileFlagIsOn(env_value) && dk == 128 && + dv == 128; +} + inline constexpr int64_t kGdnPostConvTokenTileTokens = 16; inline constexpr int64_t GdnPostConvTokenTileGridX(int64_t tokens) { diff --git a/tests/vt/test_gdn_prefill_conv.cpp b/tests/vt/test_gdn_prefill_conv.cpp index 47e288d55..c7fcd3ac9 100644 --- a/tests/vt/test_gdn_prefill_conv.cpp +++ b/tests/vt/test_gdn_prefill_conv.cpp @@ -23,6 +23,7 @@ using vt::cuda::DispatchConvChannelTileLaunch; using vt::cuda::ConvChannelTileLaunchContractFor; using vt::cuda::GdnPostConvFastFlagIsOn; using vt::cuda::GdnPostConvSplitFlagIsOn; +using vt::cuda::GdnPostConvTokenTileEligible; using vt::cuda::GdnPostConvTokenTileFlagIsOn; using vt::cuda::GdnPostConvTokenTileGridX; @@ -178,6 +179,17 @@ TEST_CASE("VT_GDN_POSTCONV_TOKEN_TILE defaults OFF; a non-'0' value enables it") CHECK(GdnPostConvTokenTileFlagIsOn(" 0")); } +TEST_CASE("GDN post-conv token tile requires both 128-wide heads") { + CHECK(GdnPostConvTokenTileEligible(false, "1", 128, 128)); + CHECK_FALSE(GdnPostConvTokenTileEligible(false, "1", 128, 64)); + CHECK_FALSE(GdnPostConvTokenTileEligible(false, "1", 64, 128)); + CHECK_FALSE(GdnPostConvTokenTileEligible(false, "1", 64, 64)); + + CHECK_FALSE(GdnPostConvTokenTileEligible(true, "1", 128, 128)); + CHECK_FALSE(GdnPostConvTokenTileEligible(false, nullptr, 128, 128)); + CHECK_FALSE(GdnPostConvTokenTileEligible(false, "0", 128, 128)); +} + TEST_CASE("GDN post-conv token tile covers each ceil(T/16) work item") { CHECK(GdnPostConvTokenTileGridX(0) == 0); CHECK(GdnPostConvTokenTileGridX(1) == 1); From 9f8dfd3eae5d1fad68562945ac4e7c41fc53c580 Mon Sep 17 00:00:00 2001 From: Richard Palethorpe Date: Sun, 9 Aug 2026 19:40:17 +0100 Subject: [PATCH 04/17] spec(bench): correct Qwen3.5 frontend timing gate (#206) Record the accepted combined local A/B, void mismatched cross-engine timing ratios, and commit the pretokenized-admission contract before implementation. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: Codex:gpt-5.6-sol [Codex] --- .agents/NOW.md | 4 +- .agents/benchmark-record.md | 86 +++++++++++++++++++ .agents/engine-matrix.md | 2 +- .agents/kernel-matrix.md | 20 +++-- .agents/roadmap_v1.md | 2 +- .../specs/sm120-qwen35-pareto-2026-08-09.md | 52 ++++++++++- docs/BENCHMARKS.md | 2 +- docs/STATUS.md | 2 +- 8 files changed, 151 insertions(+), 19 deletions(-) diff --git a/.agents/NOW.md b/.agents/NOW.md index 6f7a23596..78463d0ba 100644 --- a/.agents/NOW.md +++ b/.agents/NOW.md @@ -21,7 +21,7 @@ Work: exact-chunks on main `1ce0d662b`; sm_120 measured at `3d2581551`. | MiniMax-H3 | **PRUNED ckpts RUN (#241): Q8_0 renders, seam 0.9941** | same-binary A/B | | Kimi-Linear-48B | 122/128 held; grouped router parallelised, e2e NOT ESTABLISHED | ckpt is tiktoken-only: no warm server | | 35B binding grid | @`a0fa12c7`: **flat 0.935-0.979x, NO c2/c8 weak cell** (CoV <0.81%); mem PSS 3.81x | Attribute the flat ~5% mid-band; TTFT c2 0.872x | -| Qwen3.5-4B sm_120 | #206 opt-ins restored; tput 1.021x PASS; TTFT/TPOT/VRAM OPEN | Review; combined A/B | +| Qwen3.5-4B sm_120 | #206 K4+postconv Pareto win; IDs exact. Cross-engine tput/TTFT VOID; TPOT/VRAM OPEN | Pretokenize; rerun | | RPi5 A76 CPU | **R5 asm GREEN; llama NOT MET**: 0.461x pf, 0.653x dec | W6: BF16 GEMM | | MXFP4 parity | c1 1.020, c2-c8 0.962-0.969. **#82 CLOSED: ptxas-lineage REFUTED** | TERMINAL: at parity | | SERVE-ASYNC-DENSE-MIRROR | **LANDED+VERIFIED** (`f9c969ae`): async mirror, dense Qwen3; SACRED 184/184 | Sibling scope one-liner | @@ -53,7 +53,7 @@ latency/memory on every axis, both gate models, reproduced 2–3x idle. See of roof. Dense-marlin +0.5%; Triton-AOT GDN a WASH. 2. **Spike the Parakeet encoder row** (vLLM: `nano_nemotron_vl.py`; the transducer half is NOT in vLLM: separate call). -3. **Qwen3.5-4B #206:** review restored opt-ins; combined K4+postconv A/B. +3. **Qwen3.5-4B #206:** pretokenize benchmark; rerun cross-engine. 2. **Merge the invocation-parity prevention** (CI guard + AGENTS.md checklist); CUDA build-verify the byte-exact `kGemvHeuristicAlgos` refactor on dgx. 4. **Restore `local-ai-worker`** on dgx at campaign end (`--restart=always`). diff --git a/.agents/benchmark-record.md b/.agents/benchmark-record.md index c1b9663f5..2d5dbe0cc 100644 --- a/.agents/benchmark-record.md +++ b/.agents/benchmark-record.md @@ -18031,3 +18031,89 @@ loads 53 -> 48, shared stores 31 -> 15, barriers 2 -> 1—but global stores rose 7 -> 18 and the timing proves lost coalescing dominates. No memory pair is needed after the hard timing rejection. No acceptance, default, release, pinned-vLLM, 27B or 35B claim changes. + +## 2026-08-09 — sm_120 Qwen3.5-4B combined prefill arms accepted locally; cross-engine frontend ratios void + +**Issue:** [#206](https://github.com/mudler/vllm.cpp/issues/206). +The immutable measured binary was commit `84be99763bfcd8de94c7e6e2dd871e7a793348ea` +(SHA-256 `581ea963303b37a92f838c96e5f8f828f79b0c29081c0d4d1c240606ed3fd041`). +The fresh mutation review found that the post-conv token-tile fallback for the +independent Q/K and V head widths was not owned by a focused test. Commit +`23bc978f0e2c43265d858fc154d484be6e995f42` extracted the production +eligibility predicate and added split/env/Dk/Dv fallback coverage; the fresh +scoped mutation re-review then passed. This was a test-contract finding, not a +token or measured-product change. + +The reviewed same-binary series used the standard cached Qwen3.5-4B BF16 +workload: ShareGPT SHA-256 +`9ea13603767c62c267e3f381fbccf42d0c9ca0c393655c37533eadca7aefca0c`, +128 requests, 128 output tokens, c32, greedy, 2,048 batched-token cap and 1,280 +KV blocks. BASE set `VT_CONV_CHANNEL_TILE=0` and +`VT_GDN_POSTCONV_TOKEN_TILE=0`; COMBINED set both to `1`. The order was +`BASE-r1 -> COMBINED-r1 -> COMBINED-r2 -> BASE-r2 -> BASE-r3 -> COMBINED-r3`. +The GPU was idle before the series (RTX 5070 Ti, driver 595.71.05, P8, 0% +utilization, no compute application); the whole series held `/tmp/gpu`. + +| Arm / leg | total tok/s | output tok/s | mean TTFT | mean TPOT / ITL | mean E2E | +|---|---:|---:|---:|---:|---:| +| BASE-r1 | 6762.97 | 747.83 | 1031.42 ms | 34.77 ms | 5447.71 ms | +| BASE-r2 | 6775.87 | 749.26 | 1019.15 ms | 34.79 ms | 5437.30 ms | +| BASE-r3 | 6780.66 | 749.79 | 1018.40 ms | 34.76 ms | 5433.45 ms | +| **BASE mean** | **6773.1667** | **748.9600** | **1022.9900 ms** | **34.7733 ms** | **5439.4867 ms** | +| COMBINED-r1 | 6820.31 | 754.17 | 1010.46 ms | 34.58 ms | 5401.67 ms | +| COMBINED-r2 | 6823.24 | 754.49 | 1009.11 ms | 34.57 ms | 5399.32 ms | +| COMBINED-r3 | 6821.75 | 754.33 | 1008.28 ms | 34.59 ms | 5400.59 ms | +| **COMBINED mean** | **6821.7667** | **754.3300** | **1009.2833 ms** | **34.5800 ms** | **5400.5267 ms** | +| **COMBINED change** | **+0.7175%** | **+0.7170%** | **-1.340%** | **-0.556%** | **-0.716%** | + +Every one of the six token files has SHA-256 +`83fcdc45f79ddb06a634c7d7d95eba3384543b3cd781a45a8db1fc4e2a453545`. +Both candidate legs beat every base leg on every enclosing timing axis, so the +combined K=4 causal-conv plus 16-token post-conv arm is **ACCEPTED as a local, +default-OFF opt-in**. Raw evidence is +`/tmp/qwen35-ab-combined-84be99763/`; the six raw log hashes in arm order are +`a776a34e...dd02`, `15858c44...5c6`, `09b42ddd...e7f`, +`413713fb...c201`, `c1a664df...3f19`, and `6b0b1780...f214`. + +### TTFT attribution and the frontend mismatch + +A local COMBINED diagnostic produced **6816.79 tok/s**, **1010.16 ms TTFT** +and **34.60 ms TPOT**. Across all 128 request rows, mean +intake/queue/prefill was **196.252 / 344.389 / 469.181 ms**, summing to +**1009.822 ms**. The pinned vLLM `555967922` diagnostic produced +**6633.514 tok/s**, client/core TTFT **942.785 / 942.261 ms**, and TPOT +**33.91494 ms**; its mean queue/prefill components were +**338.960 / 449.291 ms**. vLLM's request arrival is a wall-clock timestamp +while its event timestamps are monotonic, so the raw recorded intake delta is +invalid. The only comparable intake inference is +`core_ttft - queue - prefill` = **154.011 ms**. The resulting TTFT-gap +decomposition is **42.241 ms intake + 5.429 ms queue + 19.890 ms prefill = +67.560 ms**. + +Request-id waves further localize the mismatch: + +| IDs | local mean TTFT | pinned vLLM mean TTFT | +|---|---:|---:| +| 0-31 | 2104.505 ms | 1752.132 ms | +| 32-63 | 656.106 ms | 657.784 ms | +| 64-95 | 643.766 ms | 670.053 ms | +| 96-127 | 634.913 ms | 689.077 ms | + +The source proves the frontends are not timing the same work. Local starts +`t0` before admission and submits prompt **strings** through `AsyncLLM`, so +tokenization is inside its measured interval +(`examples/bench/bench_core.h:508-529`). The oracle tokenizes every prompt +before `run_closed_loop` starts its timer and submits `TokensPrompt` IDs +(`tools/bench/vllm_closed_loop_metrics.py:59-82,137-167`). Therefore every +prior cross-engine **total-throughput and TTFT ratio on this harness is VOID**; +both axes are `PENDING` a pretokenized local rerun. The local same-binary +COMBINED A/B above remains accepted because both arms timed the same frontend. +TPOT remains comparable and open because it begins after the first token. + +Diagnostic evidence hashes: local log/token +`f7359b58...f41b` / `83fcdc45...3545`; vLLM log/metrics/splits/tokens +`9980b7b4...7148` / `61739a24...12a` / `a46753d9...3c4` / +`eebbb644...a38`. Roots are `/tmp/qwen35-combined-ttft-split-84be99763.*` +and `/tmp/qwen35-vllm-ttft-split-pin.*`. The next gate pre-tokenizes every +local prompt before `t0`, submits the existing token-ID overload, then repeats +counterbalanced same-binary rollback and fresh local/vLLM comparisons. diff --git a/.agents/engine-matrix.md b/.agents/engine-matrix.md index 5930fc695..159698900 100644 --- a/.agents/engine-matrix.md +++ b/.agents/engine-matrix.md @@ -204,7 +204,7 @@ claims it. | `SERVE-HTTP-TRANSPORT` | Serving-socket transport parity: mirror vLLM's uvicorn/asyncio default `TCP_NODELAY` on every accepted SSE socket so per-token stream frames are not held by Nagle against the peer's delayed ACK. Implemented + CPU-tested; the non-binding localhost A/B sizing is COMPLETE and NEUTRAL within noise on c1/c2 ITL/TPOT/throughput (loopback ACKs are instant, so Nagle never coalesces ~100 ms-cadence token frames) — no gate-axis credit expected; the mirror stays for real-network parity. Future keep-alive / read-write-timeout / listening-socket option parity noted, not done | T0 | vLLM serves via uvicorn over asyncio `vllm/entrypoints/launcher.py:71,76`, `vllm/entrypoints/openai/api_server.py:591,630`; asyncio disables Nagle per accepted TCP stream socket `asyncio/base_events.py:192-197` (`_set_nodelay`) called from `asyncio/selector_events.py:950`; cpp-httplib default-off `third_party/httplib/httplib.h:142`, applied on accept only when set `third_party/httplib/httplib.h:12083` | `src/vllm/entrypoints/openai/api_server.cpp:69` (`set_tcp_nodelay(true)` in the ApiServer setup) | behavioral accepted-socket `getsockopt(TCP_NODELAY)` case `tests/vllm/entrypoints/openai/test_api_server.cpp:1076` (helper `:380`); RED accepted `TCP_NODELAY` 0 → GREEN 1, full `test_openai_api_server` **22/22 cases / 242 assertions**; non-binding sizing root `~/work/vllm.cpp-tcpnodelay-sizing/ff915e8…` (raw-set SHA `f5b52900…2128`) neutral within noise; closure [ledger](parity-ledger.md#L451) | [serve-tcp-nodelay.md](specs/serve-tcp-nodelay.md) | `DONE` | `ff915e8` | | `SERVE-C-ABI` | Stable LocalAI-style C FFI (**19** exported `VLLM_API` symbols at `VLLM_ABI_VERSION 10`; blocking and nonblocking request handles. Count corrected 2026-07-24 from a stale `17`, which predated ABI v4/v5 adding `tool_parser`/`reasoning_parser` and the chat entry points; `include/vllm.h` is the source of truth and README:231 already said 19). **ABI v9 2026-07-28 (`CLAIM-CAPI-ENGINE-CONFIG-V9`): the ABI carried strictly LESS engine config than `EngineParams` does** - `max_num_batched_tokens`, the scheduler `scheduling_policy` (`fcfs` / `priority` / `lpm`), and `kv_transfer_config` (the external KV connector / LMCache JSON) were reachable from the bundled server's flags and from NO embedder. All three added, inert at their defaults (zero-filled v8 growth == byte-identical pre-v9 engine); the connector NAME is validated against `KVConnectorFactory` at load, mirroring the server's startup check. `tokenizer_config_path` stopped being a declared-since-v1 no-op and now selects the chat template's source file. Malformed `speculative_config`/`kv_transfer_config` documents now report `VLLM_ERR_INVALID_ARGUMENT` (the contract vllm.h documented since v6) instead of `VLLM_ERR_MODEL_LOAD`, via a catch scoped to the parse block so a real `FromModelDir` failure still reports MODEL_LOAD. Driver: the LocalAI vllm-cpp backend could not expose LMCache or the prefill budget in a model config) | T0 | Original project ABI; pinned vLLM has no C ABI | `include/vllm.h:143,181,207`; `src/capi/vllm_c.cpp:229,264,327,391` | `tests/capi/test_capi.cpp:320,428,505,574,606,640`; `tests/capi/test_dlopen.cpp:77,86`; `tests/capi/c_header_compile.c:1` | [c-api-library.md](specs/c-api-library.md) | `ANCHOR-BACKFILL` | `CLAIM-SERVE-C-ABI-SPIKE` | | `SERVE-CPP-API` | Rich `LLM` and `AsyncLLM` C++ API | T1 | `vllm/entrypoints/llm.py:66,422`; `vllm/v1/engine/async_llm.py:70` | - | - | `planned: specs/cpp-api.md` | `INVENTORIED` | - | -| `SERVE-CLI-BENCH` | Serve and latency/throughput/serve benchmark modes | T0 | `vllm/entrypoints/cli/serve.py:44`; `vllm/entrypoints/cli/benchmark/main.py:29`; production queue `vllm/v1/engine/core.py:200-231,622-669` | separate binaries + explicit scheduler-capacity flags `examples/server/main.cpp:63,96,116,170`; production `AsyncLLM` benchmark frontend + auditable scheduler depth `examples/bench/bench_core.h:426,495,595`; `examples/bench/main.cpp:51` | server help contract `examples/CMakeLists.txt:34`; production-frontend and metric assertions `tests/examples/test_bench.cpp:18,29-32,61,81,97` | [CLI/serve/benchmark spike](specs/cli-serve-bench.md) | `PARTIAL` | - | +| `SERVE-CLI-BENCH` | Serve and latency/throughput/serve benchmark modes | T0 | `vllm/entrypoints/cli/serve.py:44`; `vllm/entrypoints/cli/benchmark/main.py:29`; production queue `vllm/v1/engine/core.py:200-231,622-669`; pinned comparison pretokenizes before timing `tools/bench/vllm_closed_loop_metrics.py:137-167` | separate binaries + explicit scheduler-capacity flags `examples/server/main.cpp:63,96,116,170`; production `AsyncLLM` benchmark frontend + auditable scheduler depth `examples/bench/bench_core.h:426,495,595`; current string admission starts after `t0` at `examples/bench/bench_core.h:508-529`, so #206 cross-engine throughput/TTFT is VOID pending the committed pretokenized-admission correction; `examples/bench/main.cpp:51` | existing server-help/production-frontend/metric assertions `examples/CMakeLists.txt:34`, `tests/examples/test_bench.cpp:18,29-32,61,81,97`; owed #206 default/rollback callback-identity and exact prompt/output-token A/B tests | [CLI/serve/benchmark spike](specs/cli-serve-bench.md); [#206 campaign contract](specs/sm120-qwen35-pareto-2026-08-09.md) | `PARTIAL` | - | | `SERVE-GATE-ONLINE` | Same-corpus online correctness, TTFT/TPOT/ITL, throughput and peak-memory gate vs vLLM v0.25.0 | T0 | `vllm/benchmarks/serve.py:1,581-615`; [v0.25 audit](sync/2026-07-12-702f481.md); `tests/benchmarks/test_serve_cli.py:1` | Schema-v5 harness plus [trace controller](../include/vt/cuda/cuda_profiler_control.h#L13), [production component driver](../scripts/dgx-gdn-packed-component.sh), and fail-closed [component finalizer](../tools/bench/gdn_packed_component.py) | **BINDING `9ecd9d0`: 114/124** (async default ON; mem 4/4, c1 20/20, c2 20/20, c16 19/20, c4 & c32 18/20, c8 15/20; `benchmark_binding` refers here, superseding `3f256ab` 55/124 and `246a23c` 49/124, both retained immutable). Two-grid totality with `f0fb727` (111/124) is 115/124 effective parity vs vLLM 0.25.0 (27B). Async CLOSED the c16/c32 ITL tails (ours now BEATS vLLM: c16 p99 1.055, c32 p90 1.034/p99 1.078) and leaves a stable c8 `p99_itl` ~0.86 residual, ROOT-CAUSED (2026-07-18, `CLAIM-C8-P99-TAIL-1`, [spec](specs/c8-p99-itl-tail-2026-07-18.md)) as IRREDUCIBLE-AS-MIRRORED: our deterministic synchronous forward keeps co-admitted c8 requests in byte-identical lockstep where vLLM's async-future jitter de-phases them; the c16/c32 INVERSION proves this is the trailing edge of the per-step determinism that wins c16/c32 + throughput, not a capability gap (scheduler + async placeholder byte-identical, `tests/vllm/v1/test_scheduler_wave.cpp:265`, [tail spec](specs/tail-stall-analysis-2026-07-16.md)). Full grid + per-binding forensics: roadmap_v1.md + parity ledger; no packed speed credit | [online serving gate](specs/cuda-online-serving-gate.md); [merged GDN projections](specs/gdn-merged-input-projections.md); [packed decode](specs/gdn-packed-decode.md) | `ANCHOR-BACKFILL` | CLAIM-SERVE-GATE-1 | | `SERVE-E2E-NIGHTLY` | Server conformance and real-model nightly suites for all release gates | T0 | `tests/entrypoints/openai/`; `tests/v1/e2e/`; `.buildkite/test-pipeline.yaml` | current unit/conformance tests only; no scheduled DGX suite | `tests/vllm/entrypoints/openai/test_conformance.cpp:1`; `tests/parity/test_qwen36_paged_engine.cpp:78`; `tests/parity/test_qwen27_paged_engine.cpp:110` | `planned: specs/server-e2e-nightly.md` | `INVENTORIED` | - | | `ENG-RELEASE-BINARIES` | Downloadable host-ABI-specific `vllm-server` bundles: adaptive CPU and fat CUDA primary artifacts, optional per-SM diagnostics, and literal-static feasibility boundary | T0 | vLLM release lanes `.buildkite/release-pipeline.yaml:1-18,34-170` @ `555967922`; release-image dependency boundary `docker/Dockerfile.cpu:262-290` | Required W1-W11/W13 implementation is complete in draft PR #196: ten-SM gencode/AOT, adaptive CPU tiers, extracted-archive validation and supply chain, least-privilege immutable handoff, eight primary bundles, byte-derived indexes, attestation, and exact-file publication; W12 remains optional/non-primary | Local manifest/archive/workflow/mutation gates, full adaptive x86 tier execution, clean CPU archive, and Vulkan 35/35 backend + 11/11 cross-device archive run are green; hosted ten-SM completion, full eight-tuple dry run, matching-hardware gates, and tagged publication remain pending, so no published-binary claim exists | [release-binary-matrix.md](specs/release-binary-matrix.md) | `ACTIVE` | `CLAIM-ENG-RELEASE-BINARIES-W1-W13` | diff --git a/.agents/kernel-matrix.md b/.agents/kernel-matrix.md index b16876e93..290f26474 100644 --- a/.agents/kernel-matrix.md +++ b/.agents/kernel-matrix.md @@ -174,15 +174,17 @@ release gates are unchanged. [Spec and evidence](specs/sm120-qwen35-conv-chunkin **2026-08-09 `KERNEL-SSM-MAMBA` sm_120 campaign anchor.** [#206](https://github.com/mudler/vllm.cpp/issues/206) tracks the RTX 5070 Ti -Qwen3.5-4B Pareto campaign. The current sealed result passes throughput -(**6784.743 vs 6643.593 tok/s**) while TTFT (**1018.040 vs 937.584 ms**), -TPOT/ITL (**34.740 vs 33.906 ms**) and VRAM (**13053.3 vs 12820 MiB**) remain -open. The historical default-off K=4 causal-conv, 16-token post-conv and -BV16/swizzle/register-state decode arms and their focused contracts are now -restored on the current branch; fresh performance measurement remains pending. -The first ordered gate measures combined K=4 causal-conv plus 16-token -post-conv in a counterbalanced A/B; -same-tool local/vLLM tracing then selects the largest residual. Lifecycle +Qwen3.5-4B Pareto campaign. The historical sealed measurement reported a +throughput pass (**6784.743 vs 6643.593 tok/s**) while TTFT +(**1018.040 vs 937.584 ms**), TPOT/ITL (**34.740 vs 33.906 ms**) and VRAM +(**13053.3 vs 12820 MiB**) remain open. The reviewed combined default-off K=4 +causal-conv plus 16-token post-conv A/B is accepted locally: **+0.7175%** total +throughput, **-1.340%** TTFT, **-0.556%** TPOT and **-0.716%** E2E, with six +identical token hashes. The cross-engine throughput and TTFT ratios are now +**VOID/PENDING** because local timed string tokenization while pinned vLLM +pre-tokenized before its timer; TPOT remains comparable and open. The committed +next gate pre-tokenizes local prompts before `t0`, retains +`VT_BENCH_PRETOKENIZE=0` same-binary rollback, and reruns both engines. Lifecycle stays `INVENTORIED`; generic Mamba and 27B/35B coverage are unchanged. [Campaign spec](specs/sm120-qwen35-pareto-2026-08-09.md). diff --git a/.agents/roadmap_v1.md b/.agents/roadmap_v1.md index 3269e645f..54feca634 100644 --- a/.agents/roadmap_v1.md +++ b/.agents/roadmap_v1.md @@ -101,7 +101,7 @@ models we already ship + benchmark. Full seam map + M0–M5 W-plan: | MEM | `ROAD-V1-MEM` | **Memory budgeting: auto-size to the declared workload by default, optional total-footprint cap, pre-flight error instead of an OOM (user-directed 2026-08-06, [#83](https://github.com/mudler/vllm.cpp/issues/83))** — the user-facing wart that every operator hits before they hit any perf question: vLLM makes you compute your own VRAM budget, express it as a PERCENT, and nail it or OOM | [coverage view §2](feature-matrix.md#2-kv-cache--memory), [porting inventory](porting-inventory.md) | — **M1+M2 LANDED 2026-08-08** ([`specs/kv-sizing.md`](specs/kv-sizing.md)): the absolute `--kv-cache-memory` knob sizes the pool via a group-aware `KVBytesPerBlock` divisor, `--num-blocks` is the override, both mirrored on the C ABI at v16; `ResolveNumBlocks` precedence `num_blocks > bytes > 256`, CPU-gated (`KVBytesPerBlock` 5/5 + `test_capi` v16). M3 (the `gpu_memory_utilization` profile run) stays dgx-gated. (M0 design grounded in vLLM `config/cache.py` + `gpu_worker.py:497-599`; GB10 unified-pool caveat carried) | `M1+M2 DONE` | **Source-verified 2026-08-06 (records-only, NO code).** WE ARE CURRENTLY BEHIND vLLM ON THIS AXIS, not ahead: there is NO memory profiling at all and the KV pool is a RAW BLOCK COUNT the user types by hand — `EngineParams::num_blocks = 256` (`include/vllm/entrypoints/model_loader.h:58`, beside `block_size = 32` `:57` / `max_model_len` `:59` / `max_num_seqs = 8` `:60`), exposed verbatim as `--num-blocks N` (`examples/server/main.cpp:100,203-204,370`), carried on the C ABI as `vllm_model_params.num_blocks` at the same 256 default (`src/capi/vllm_c.cpp:429,485`), landing as `BlockPool(num_gpu_blocks, ...)` which asserts `> 0` and otherwise TRUSTS it (`include/vllm/v1/core/block_pool.h:96,223`; `src/vllm/v1/core/block_pool.cpp:51`). So a user must convert "40 GB free, 32k context, concurrency 8" into a block count themselves — strictly worse ergonomics than a percentage. UPSTREAM HAS THREE KNOBS, all `config/cache.py`, all T0, all rowed at [porting-inventory.md](porting-inventory.md) §T0: `gpu_memory_utilization` (fraction of TOTAL, default 0.9), `kv_cache_memory_bytes` (absolute KV pool) and `num_gpu_blocks_override` (exact block pin), sized off a profile run as `total x utilization - non-torch - peak activation`. **Mirroring that is necessary but does NOT solve the complaint**, for three source-grounded reasons: (1) the fraction is of TOTAL not FREE, so on any shared card the right fraction is a function of what someone else already holds — exactly the arithmetic the user is being asked to do; (2) WEIGHTS LOAD BEFORE THE KNOB ENGAGES (utilization sizes the KV pool AFTER the model is resident), so an oversized model OOMs during load and never reaches the check — which is the failure operators actually hit; (3) 0.9 is taken whether or not it is needed (a 4B model on an 80 GB card reserves 72 GB it will never touch and blocks everything else on the device). THE DESIGN (user-ratified 2026-08-06) is three modes: **Mode 1 default = SIZE TO THE DECLARED WORKLOAD** — per-allocation-class accounting BEFORE allocating anything (weights from checkpoint metadata / safetensors header / GGUF manifest, known before reading a byte of tensor data; CUDA context measured at creation; peak activation from a profile run at `max_num_batched_tokens`; KV for `max_model_len x max_num_seqs` at the resolved `block_size`/KV dtype; CUDA-graph capture-set footprint) → allocate exactly that and LEAVE THE REMAINDER OF THE DEVICE FREE. This is the SURPASS over vLLM, which takes its 90% regardless of whether the workload needs 8 GiB or 80. **Mode 2 = a cap on the TOTAL ENGINE FOOTPRINT** (weights + activations + KV + graph pools + context), NOT on the KV pool alone — load-bearing, because a KV-only cap cannot prevent the weight-load OOM; three spellings of the same cap: `--memory-limit 40GiB` (absolute, the primary form), `--gpu-memory-utilization 0.85` (vLLM's exact flag name and fraction semantics so existing vLLM launch lines port unchanged, per [[mirror-vllm-always-no-asking]]) and `--num-gpu-blocks-override N` (upstream's reproducibility escape hatch — this is where today's `--num-blocks` GOES, demoted from primary knob to explicit override), with precedence spelled out and TESTED, not left to argument order. **Mode 3 = REFUSE BEFORE ALLOCATING** with the full per-class breakdown and remedies COMPUTED FROM THE ACTUAL BUDGET (`--max-model-len N` / `--max-num-seqs M` / `--kv-cache-dtype fp8` / smaller quant, each with the GiB it recovers) — "you are 43.9 GiB over and here are the three levers that close it" is the difference between a usable error and a stack trace. **UNIFIED-MEMORY HAZARD (not hypothetical):** on GB10 the ~119 GiB pool is UNIFIED, a fraction-of-total setting reserves HOST RAM too, and `gpu_memory_utilization=0.85` has HARD-REBOOTED our DGX three separate times ([[gb10-unified-memory-oom-reboots-box]]) — hence absolute bytes is the PRIMARY form with the percentage kept only for vLLM compatibility, and hence the accounting must know whether the pool is unified, which makes free/total + an is-unified predicate a PLATFORM-SEAM question belonging behind `ROAD-V1-C1`'s abstraction (note `Platform::needs_weight_staging()` was deliberately NOT `is_unified_memory()` because the latter FLIPS GB10 — the distinction matters here) rather than a CUDA-specific branch, since discrete and unified devices need different safety margins. CORRECTNESS: pool size changes preemption/scheduling TIMING but not emitted tokens, so the SACRED token-exact gates are unaffected — and M2's gate makes that explicit by re-running them with no block flag at all. **Next gate = M0 the `specs/kv-sizing.md` spike (accounting model + precedence rules + upstream `file:line`); then M1 a `MemoryBudget` computing required bytes per class WITHOUT allocating plus the Platform free/total + is-unified seam, unit-gated predicted-vs-actual weight bytes; M2 auto-sizing as the default with `--num-blocks` demoted to `--num-gpu-blocks-override`, gated by every existing model gate running with NO block flag and staying token-exact; M3 the three caps + precedence through the server flags and the C ABI, gated by our KV pool matching vLLM's own at a matched `--gpu-memory-utilization`; M4 the pre-flight refusal, gated by a deliberately over-subscribed config exiting cleanly on GB10 (non-zero exit, no OOM, NO BOX REBOOT) covering the unified-pool path specifically; M5 (optional) a runtime guard failing the REQUEST rather than the engine.** Docs (README, [STATUS](../docs/STATUS.md)) update in the SAME change as whichever milestone shifts externally-visible behaviour, per [[keep-readme-current]] | | 1 | `ROAD-V1-C1` | **Extensibility-first (USER PRIORITY 2026-07-18):** drop-in kernel ABI + the MISSING Platform seam + model self-registration — make new GPUs/archs/models ADDITIVE (plan: [extensibility-platform-seam-2026-07-18.md](specs/extensibility-platform-seam-2026-07-18.md)) | [`BACKEND-ABI-VT`](backend-matrix.md), [kernel matrix](kernel-matrix.md) | exhaustive kernel/dependency inventory and [raw-pointer adapter ABI](specs/dropin-kernel-abi.md) accepted; additive W0 implemented and CPU 94/94. `CLAIM-BACKEND-ABI-W0-GPU-1` repaired the GCC13/doctest blocker without runtime changes; exact sm_121a all-target build, focused CUDA/ABI sanitizer, and both gate-model tests pass at `1141b79`. Cross-arch/trace/A-B and scalar-forwarder/backend-shim debts remain explicit | `PARTIAL` | **★ NEW ORDER-1 HEAD (user-directed 2026-07-19): the PORTABLE AUTOMATIC OP-FUSION FRAMEWORK (`KERNEL-FUSION-FRAMEWORK`, spike [portable-fusion-framework.md](specs/portable-fusion-framework.md), `SPIKE`).** The extensibility cornerstone: fusions DECLARED ONCE (backend-agnostic `constexpr FusedRecipe` catalog above `vt::`, transcribing vLLM's finite pattern-pass set `passes/fusion/*` @ `pass_manager.py:138-200`, mirroring the `CustomOp` `forward_native`/`forward_cuda` seam `custom_op.py:103`) and REALIZED PER-BACKEND through the existing `vt::` op table (Tier-0 composite = the CPU oracle inherited free by any backend; Tier-1 interpreter = one kernel port per backend lights up every recipe). Makes a new vLLM fusion PR a ONE-declaration port, a new GPU an additive catalog realization, a new model an additive pattern declaration — the PR-#4 remedy, composed with the Platform/attn-registry/model-registry seams below. The TDR Phase-0 skeleton is already LANDED (`fused_recipe.h`/`recipes.h` one recipe + `OpId::kFusedChain` Tier-0/1 on CPU+CUDA + byte-exact `test_ops_fused_chain.cpp`). **W0 ADOPTED 2026-07-19 (`CLAIM-FUSION-FRAMEWORK-W0`):** the seam is now used in production at ONE real site — the 35B `RunLayerPaged` post-attention layernorm routes its plain add+residual+gemma-RMSNorm through `vt::FusedChain(kFusedAddRmsNorm)` (`VT_FUSED_CHAIN_ADOPT` default-ON / `=0` rollback), behaviour-preserving + byte-identical to the prior hand-call (DGX: clean CUDA `-Werror` 0 warn, byte-exact composite==interp==golden incl. H=2048, 35B 315/315 + 27B 235/235 token-exact BOTH arms, memcheck 0 errors) — proving the declare-once/realize-per-backend seam end-to-end; the current 3-opcode/4-role POD sufficed byte-identically for the plain add+rmsnorm class, so W0 needed NO generalization. **W1 GENERALIZED the POD 2026-07-20 (`CLAIM-FUSION-FRAMEWORK-W1`, `1115648`):** full activation/norm/quant/rope opcode set + indexed operand table; all 5 quant-fused chains declared byte-exact; Tier-0 composite = ONE device-agnostic walker (kills CPU/CUDA oracle drift); infrastructure only, no call site changed (DGX: `-Werror` 0-warn, byte-exact CPU 196 + CUDA 361, memcheck 0, 27B 235/235 + 35B 315/315 both arms). **W2 MIGRATED the hand-fusions 2026-07-20 (`CLAIM-FUSION-FRAMEWORK-W2`):** the framework now OWNS the fusion dispatch — each recipe binds (new backend-agnostic `FusedRecipe.fast_op`) to its EXISTING single-launch fused kernel, so `FusedChain(recipe)` dispatches to the SAME fast kernel the model called directly pre-migration (byte-identical + perf-neutral by construction; composite is the graceful fallback + oracle). SIX call sites migrated behind `VT_FUSED_CHAIN_ADOPT` (`kSiluMulFp4Quant`/`kSigmoidGateFp4Quant`/`kRmsNormGatedQuantFp8`×2/`kRmsNormQuantFp8`/`kAttnQkNormRopeGate`×2). DGX: `-Werror` 0-warn, byte-exact CPU 228 + CUDA 420, memcheck 0, 27B 235/235 + 35B 315/315 BOTH arms. **W3 MECHANICAL-SYNC PROOF LANDED 2026-07-20 (`CLAIM-FUSION-FRAMEWORK-W3`):** ported a NEW, previously-unported vLLM fusion pass — `SiluMulFp8StaticQuantPattern` (`act_quant_fusion.py:81` → `_C.silu_and_mul_quant`, the static-per-tensor-FP8 sibling of `kSiluMulFp4Quant`) — as ONE `constexpr FusedRecipe kSiluMulQuantFp8` + its byte-exact test, touching EXACTLY 2 shared files (`recipes.h` + `test_ops_fused_chain.cpp`), NO kernel/dispatch/composite-walker/model-site edit and NO new primitive (composite = existing `vt::MoeSiluMul` + `vt::QuantFp8Static`; `fast_op=kNoFastOp`). The PR-#4 additivity test made concrete: a whole new fusion pattern = one declaration. DGX: `-Werror` 0-warn, byte-exact CUDA 432, memcheck 0, no token regression (recipe declared-only) 27B 235/235 + 35B 315/315. **W4 BACKEND-ADDITIVITY PROOF LANDED 2026-07-20 (`CLAIM-FUSION-FRAMEWORK-W4`) — the W-series proof milestone is DONE:** made the additivity claim EXECUTABLE — new test `test_fused_chain_additivity.cpp` treats the EXISTING CPU backend AS the 'second backend' relative to CUDA (no mock `DeviceType` — that would edit the core enum + every switch, ironically non-additive) and, in ONE generic loop over the WHOLE catalog (all 7 recipes), asserts each runs byte-exact on the CPU backend via the Tier-0 composite — 4 CPU-full end-to-end + 3 fp8-prefix (byte-exact prefix + the FULL composite asserted to THROW on CPU, documenting the CUDA-only static-fp8 backend-negotiated tail, §3b/§6). Additivity evidence: catalog `recipes.h` grew 1→6→7 while the composite walker stayed ONE per-OPCODE function (12 `FOp::` cases) + the CPU/CUDA `kFusedChain` registration ONE line each + `cpu_ops.cpp` never `#include`s `recipes.h` — W3's whole new recipe `kSiluMulQuantFp8` is in ZERO backend TUs, inherited free. CPU `-Werror` 0-warn, `test_fused_chain_additivity` 17/17 + `test_ops_fused_chain` 228/228; engine byte-identical (no `src/`/`include/` change) ⇒ 27B 235/235 + 35B 315/315 structurally unchanged; memcheck N/A. Honest deferred (named, non-blocking the ORDER-1 milestone): Tier-1 perf interpreter for the quant chains (composite-only today), a REAL Metal/Vulkan catalog realization (M4 HW-blocked), and per-recipe fast single-launch kernels. Honest payoff: perf ceiling ~3.5%/step compute-bound on 35B (NOT a perf lever — tasks #61/#62; W0 is perf-neutral by construction); primary value = extensibility + mechanical upstream-sync + CPU/CUDA oracle-drift elimination. Incremental W0 adopt-one **(DONE)** → W1 generalize POD **(DONE)** → W2 migrate hand-fusions **(DONE)** → W3 mechanical-sync proof **(DONE)** → W4 backend-additivity proof **(DONE)** → Wn honest re-measure (optional, off the extensibility critical path). **W-SERIES ORDER-1 PROOF MILESTONE DONE 2026-07-20 (`CLAIM-FUSION-FRAMEWORK-W4`).** **PRIOR extensibility items (all LANDED, the seams this composes with):** **#1 extensibility item — extract the Platform seam — LANDED 2026-07-18** (`BACKEND-PLATFORM` `ACTIVE`, `CLAIM-BACKEND-PLATFORM-1`): `include/vllm/platforms/interface.h` + `src/vllm/platforms/{platform,cpu,cuda}.cpp` mirror `vllm/platforms/interface.py:134-229` 1:1; `CurrentPlatform()` self-registered per `DeviceType`; the 7 memory-model/residency `device.type == kCUDA` sites (of PR #4's ~37) now route through it → new-GPU memory model is ONE additive `platforms/.cpp`. Behavior-preserving (clean CPU build + `test_platform` + full CPU CTest green; DGX 235/235 + 315/315 pending). **Item 2 residency-as-Platform-capability LANDED 2026-07-19** (`CLAIM-BACKEND-PLATFORM-2`): the host-free / load-stream / DevicePool-cap decisions in `qwen3_5.cpp` now READ `GetPlatform(.device.type).residency_policy()` (per-device) via the pure `ShouldReleaseHostWeights`/`ShouldInterleaveLoadStream` helpers + `device_pool_cap_bytes`, not an inline `device.type`/env gate; `CudaPlatform.release_host_weights_after_upload` flipped false→true (now CONSUMED ⇒ reproduces today's GB10 host-free-after-Marlin + ~4 GiB load peak EXACTLY); `MarlinMoeEnabled()` stays the orthogonal kernel-path gate. **A new (discrete) GPU sets `residency_policy()` values ⇒ ZERO model edit.** Behavior-preserving (clean CPU build + `test_platform` consumption cases 7/43 + full CPU CTest + tools 164/164 green; **DGX-CONFIRMED @ `62fc0e0`: clean CUDA `-Werror`, 27B 235/235 + 35B 315/315 token-exact, 35B VmHWM ≈ 4.0 GiB load-stream win preserved, memcheck 0 errors**). Then item 3 drop-in ABI family migration. **Item 4 attn-backend registry LANDED 2026-07-19** (`CLAIM-ATTN-REGISTRY-1`, `BACKEND-ATTN-REGISTRY`): NEW `include/vllm/v1/attention/registry.{h,cpp}` `(DeviceType,name)` registry + `SelectAttentionBackendName` selector (mirror `registry.py` self-registration + `cuda.py:361-470` `get_attn_backend_cls`/`_get_backend_priorities`); `Platform::get_attn_backend_priority()` filled (was the item-1 STUB) → capability-ordered name lists on `CudaPlatform` (major-10 vs else) + `CpuPlatform`; FLASH_ATTN/GDN self-register. **Adding a backend's attention = 1 self-registering TU + 1 priority slot, ZERO selector/model/runner edit.** Behavior-preserving — the walk returns FLASH_ATTN on CUDA+CPU (the same FA2 attention runs); clean CPU build + `test_attn_backend_registry` (8/25) + full CPU CTest, tools 164/164, checkers green; **DGX-CONFIRMED @ `2c732e7`: 27B 235/235 + 35B 315/315 token-exact (FA2 sm_121a), memcheck 0/315**. **Item 5 model self-registration LANDED 2026-07-19** (`CLAIM-MODEL-SELFREG-1`, `MODEL-FACTORY-registry`): the fixed `kRegistrations` array → `REGISTER_VLLM_MODEL(...)` static-`Registrar` idiom (`model_registry.h:167-189`) + Qwen dense/MoE arch entry points split into per-variant TUs (`qwen3_5_dense.cpp`/`qwen3_5_moe.cpp`) over shared `qwen3_5_common.{h,cpp}`, so **adding a model = 1 new TU + 1 REGISTER line, ZERO shared-array edit**; behavior-preserving (clean CPU build + `test_model_registry` extension + full CPU CTest, tools 164/164, checkers green; DGX 27B/35B token-exact pending). Deep `qwen3_5.cpp` machinery factoring deferred. Metal/MLX bring-up proves the seams (needs M4). **★ THE ARCH HALF OF THIS ITEM IS NOW PROVEN BY MEASUREMENT, NOT ARGUED (2026-07-22, `CLAIM-CUDA-SM120-BRINGUP`, [spec §W8](specs/cuda-arch-additivity.md)):** consumer-Blackwell `sm_120a` was brought up as a BUILD-supported target and required **ZERO kernel, model, runner, sampler or feature-table edits** — the additive seams (per-arch FEATURE TABLE, capability-keyed tactic registry keyed on `major == 12`, Platform auto-probe, `pageable && integrated` residency classification) already covered it, so the entire diff is build configuration, a configure-tier test and records. That is the PR-#4 additivity test passed on a real second architecture. It is deliberately NOT a runtime-support claim: no sm_120 board exists here. **★ THE MODEL/QUANT HALF NOW ADVANCES TOO (2026-07-23, `CLAIM-BACKEND-SEAM-S4-1`): the `model_executor/layers/` `LinearMethod`/`QuantizationConfig` seam the [accelerator-seam audit](specs/accelerator-seam-audit.md) §9 named ABSENT now EXISTS in part.** `S4` landed byte-identical: the dense model's projections route through a `method.Apply()` chosen ONCE from the checkpoint (retiring the per-call `IsNvfp4()` tensor-name probe), and 18 shared-layer `device==kCUDA` availability gates became `vt::OpRegistered` op-table queries — the policy(scheme)/implementation(kernel) split the audit's binding rule requires. **DSR 86 → 67**; all six SACRED gates byte-identical on dgx (27B/35B/Coder/dense/OPT/DeepSeek-V2); the fragile 27B-W4A4/fp8-recipe gates are correctly deferred to `S6` behind `S5`'s reference tier. **★ `S6` ASSESSED 2026-07-23 (`CLAIM-BACKEND-SEAM-S6-1`) → NO-OP / BLOCKED, DSR stays 67 (§11):** the deferred fp4/fp8 gates convert ZERO sites byte-identically — every one bottoms out at a **dual-registered** (CPU+CUDA) bespoke op (none CUDA-only, unlike S4's convertible gates), so `OpRegistered(op,dev)` is TRUE on `kCPU` ⇒ the class-A swap is bit-changing on the CPU reference/emulation path (two numerics per device); S5's reference tier does not change this (those CPU kernels are present natively, never a miss). No `src/`/`include/`/test byte changed, no baseline moved. The genuine byte-identical unlock is re-scoped to **`S3`** (Platform capability fields mirroring `supports_fp8`/`cutlass_fp4_supported` — the audit's own class-D fix) and **`S7`** (layer extraction); the plan's `~37` S6 target assumed the class-A `OpRegistered` swap was byte-identical, which holds only for CUDA-only ops (all taken by S4). **★ `S3` LANDED 2026-07-23 (`CLAIM-BACKEND-SEAM-S3-1`) — the byte-identical unlock S6 re-scoped:** mirrors vLLM's `Platform` capability surface (`supports_fp8`/`cutlass_fp4_supported`/`opaque_attention_op`/`is_integrated_gpu`/`support_static_graph_mode`/`is_device_capability_family`, base false in `interface.h`; CudaPlatform answers GB10 values in `cuda.cpp`/`platform.cpp`) and converts **12** deferred `qwen3_5.cpp` gates onto it (7 fp4-act `cutlass_fp4_supported`, 3 fp8-fused `supports_fp8`, 2 decode-graph `support_static_graph_mode`) — byte-identical because a capability answers the base false off CUDA, exactly what `device==kCUDA` did (where S6's `OpRegistered` was TRUE on `kCPU`), and it DECOUPLES (a future accelerator answers for itself). **DSR 67 → 55** (`kcuda` 25→13), baseline lowered same commit, ratchet + 24-case suite green; all seven SACRED gates byte-identical on dgx (27B 235/235 · 35B 315/315 · Coder · dense-32B · OPT · DeepSeek-V2 · Llama), `test_platform` CUDA-leg proves each predicate == former `device==kCUDA` on GB10, memcheck 0 errors, clean CUDA+CPU `-Werror`. Residency/stream/FA2-dtype/merged-layout sites LEFT for `S7`. **★ `S7` LANDED 2026-07-23 (`CLAIM-BACKEND-SEAM-S7-1`) — the seam campaign's TERMINAL runtime-decoupling state, the closest this extensibility work comes to a finish line:** ALL 23 remaining runtime `kCUDA`/`is_cuda()` sites in the shared model layer hoisted onto capabilities — new `Platform::needs_weight_staging()` (the CUDA device-resident staging policy, NOT `is_unified_memory()` which would FLIP GB10; covers residency/merged-GDN/packed-decode/direct-load), `Platform::supports_fa2_attention()` (FA2 dtype), `Backend::SupportsAuxStream()` (MoE aux-stream), reusing S3's `supports_fp8`/`cutlass_fp4_supported`/`support_static_graph_mode`/`is_integrated_gpu` (runner combine/scatter) and `vt::OpRegistered(kMoeGroupedGemmBf16)` (a CUDA-only op). Each returns the former `device==kCUDA` value on GB10 → byte-identical. **DSR 55 → 32 — the IRREDUCIBLE build-gate floor:** the shared model layer holds ZERO runtime device tests; the 32 residual are all `#ifdef VT_*` compile-time gates for kernels that only build on one GPU family (a kernel that only compiles on one arch is legitimately irreducible), so the audit's `<10` is NOT reachable and this is the honest answer to "how additive can the shared layer get" — every runtime device coupling is gone. baseline lowered same commit, ratchet + 24-case suite green; all seven SACRED gates byte-identical on dgx (27B 235/235 · 35B 315/315 · Coder 6/6 · dense-32B 16/16 · OPT 6/6 · DeepSeek-V2 8/8 · Llama 16/16), new `test_platform`/`test_backend`/`test_cuda_backend` cases green, memcheck 0, clean CUDA+CPU `-Werror`. The `layers/`-library physical relocation (shrinking `qwen3_5.cpp` toward `qwen3_next.py`'s 802-line shape) is a follow-on refactor; the device coupling it was to remove is already gone | | 2 | `ROAD-V1-C2` | Model families: Llama/Qwen3/Mistral, MoE, Qwen3-Next | [model matrix](model-matrix.md) | current pin has 353 static IDs; v0.25.0 adds three sync-target rows (MOSS-Transcribe-Diarize, Laguna DFlash, Bailing hybrid MTP), yielding 356 after pin advance. **FIRST ADDITIVE-MODEL BRING-UP W0-W4 LANDED 2026-07-20 — CORRECTNESS COMPLETE (0.6B + 4B gates PASS 16/16; SPEED pending)** ([first-additive-model-qwen3-dense.md](specs/first-additive-model-qwen3-dense.md), `MODEL-TEXT-qwen3-qwen3-for-causal-lm` `ACTIVE`(correctness COMPLETE, speed pending), runner generalization `ENG-RUNNER-MODELSHAPE`, `CLAIM-MODEL-QWEN3-DENSE`) **MLA CAMPAIGN SPIKED 2026-07-21** ([mla-deepseek-campaign](specs/mla-deepseek-campaign.md), `CLAIM-MLA-DEEPSEEK`): five rows `INVENTORIED` -> `SPIKE` (DeepSeek V2 / V3+V3.2 / v1-MHA, Kimi-Linear, MiniMax-M2). **KIMI-LINEAR-48B W0 DEDICATED SPIKE 2026-08-05** ([kimi-linear.md](specs/kimi-linear.md), `CLAIM-KIMI-LINEAR-W0`): full dedicated W0 spike for `MODEL-TEXT-kimi-linear-*` (stays `SPIKE` — actively claimed) — the ONE Kimi text model that FITS one GB10 (91.5 GiB, 0.77x pool) with a real e2e SACRED gate; HEAVY reuse (MLA + sigmoid/`noaux_tc` MoE + GDN family + KDA host refs landed), NET-NEW = the KDA device kernel + NoPE-MLA branch + hybrid schedule/loader; W1 implementation can start. Answers the Tier-3 "MLA = new attention, new campaign" item in [breadth-sweep-plan](specs/breadth-sweep-plan.md) §B.3. Key determinations: GB10/sm_121 selects **`TRITON_MLA`** for dense MLA decode and **`FLASH_ATTN`** for MLA prefill (`vllm/platforms/cuda.py:129-133`, `mla/prefill/selector.py:74-77`), so the sm90/sm100-only MLA kernels are out of reach and out of scope; the cross-cutting cost is the **compressed-latent KV cache** (one 576-wide latent per token, `num_kv_heads=1`, no separate V), which our allocator and `vt::ReshapeAndCache`/`vt::PagedAttention` cannot express; and **only DeepSeek-V2-Lite (~29.3 GiB bf16) fits GB10** — V3/V3.2, Kimi-K2.5, MiniMax-M2/M3 are HW-BLOCKED e2e, Kimi-Linear-48B is HW-MARGINAL. W0-W10 plan recorded; nothing implemented. **GLM + DSA + LATEST-DEEPSEEK SPIKED 2026-07-21** ([glm-dsa-latest-deepseek](specs/glm-dsa-latest-deepseek.md), `CLAIM-GLM-DSA-LATEST-DEEPSEEK`): seven rows `INVENTORIED` -> `SPIKE` (ChatGLM, Glm, Glm4, Glm4Moe, Glm4MoeLite, GlmMoeDsa, DeepSeek-V4). Answers the user's "also glm, and deepseek latest versions" priority. Headline: **`Glm4MoeLiteForCausalLM` / `zai-org/GLM-4.7-Flash` (31.2B, 58.2 GiB bf16) FITS GB10 and is a SECOND MLA gate vehicle that closes BOTH coverage gaps the MLA campaign named as unit-gated-only** (it has `q_lora_rank=768` and `noaux_tc`/`e_score_correction_bias`, which DeepSeek-V2-Lite lacks). **DSA is DOUBLY BLOCKED on GB10:** the sparse XOR filter eliminates `TRITON_MLA` for sparse models leaving `FLASHINFER_MLA_SPARSE_SM120` as the sole candidate, and that path is non-functional on flashinfer 0.6.12 (XQA backend is dense-only, discards `sparse_mla_top_k`); GLM-5 is 1404 GiB and V3.2 642 GiB regardless. DeepSeek-V4 is a NEW architecture (Sinkhorn-normalized Manifold Hyper-Connections, CSA/HCA compressor with recurrent state, hash-routed MoE) and HW-BLOCKED at 148.7 GiB — but its TOKENIZER risk is LOW (standard HF fast BPE; only the chat template needs porting, with upstream golden fixtures). Glm4/Glm need two primitives we have NONE of: partial rotary factor and sandwich norms. Nothing implemented. **NVFP4A16 (W4A16)** on the already-done dense `Qwen3ForCausalLM` (`RedHatAI/Qwen3-32B-NVFP4A16`, 64L) — the QUANT-SCHEME additivity experiment, serving user priorities #2 (models) and #4 (quants) at once. KERNEL LAYER FULLY ADDITIVE (ZERO new kernel code: vLLM FORCES Marlin for `use_a16`, OBSERVED `Using MarlinNvFp4LinearKernel`, and that is the GEMM we already vendored for the 35B). **CORRECTNESS CLOSED 2026-07-21 (W4b):** the strict gate's 4/6 was diagnosed by the ratified TEACHER-FORCING isolation — all 29 divergent positions gap <= 0.0625 nats with 28/29 EXACTLY 0.0, one root flip an EXACT bf16 tie at which vLLM's teacher-forced argmax is OURS and vLLM contradicts its own greedy. **NOT a W4A16 defect: it is the pre-existing dense-forward bf16 near-tie drift, recorded against `MODEL-TEXT-qwen3-qwen3-for-causal-lm`.** Gate closes **6/6** under the ratified near-tie-robust bar with the nats evidence committed. SPEED still pending ⇒ `ACTIVE`, not `DONE`. **GEMMA FAMILY SPIKED 2026-07-24** ([sweep-gemma](specs/sweep-gemma.md), `CLAIM-SWEEP-GEMMA`): four rows `INVENTORIED` → `SPIKE` (Gemma 1/2/3/4). Answers the user's "and then we do gemma" ("gemma 4") next-target. **The newest registered Gemma is Gemma 4** (real, public, but ALL checkpoints multimodal-wrapped `Gemma4*ForConditionalGeneration`, ≥12B, 0.25.0 oracle-support unverified, needs a PLE/YOCO/MoE/k_eq_v stack) — it leads the characterization but is gate-BLOCKED as a first vehicle. **The recent-first gate vehicle that FITS + is oracle-certain is Gemma 3** (`Gemma3ForCausalLM` on `google/gemma-3-1b-it`). Headline: Gemma reduces MOSTLY to landed infra — gemma-RMSNorm, sandwich norms (glm4 `b568d20`), SentencePiece (names "Gemma"), sliding-window (FA-2 + SlidingWindow/ChunkedLocalAttention specs), the `kAttnQkNormRopeGate` QK-norm+rope recipe, tied embeddings are ALL REUSE; the one genuinely-new compute kernel is GeGLU (`gelu_pytorch_tanh`+mul; we have only SiLU), plus the final logit soft-cap + qpas/embed-scale scalars + dual-rope routing. Per-version delta: Gemma-2 has an attn logit soft-cap, Gemma-3 removed it for QK-norm. **GEMMA-3 W0-W2 LANDED 2026-07-24 — CORRECTNESS COMPLETE, the FIRST Gemma family** (`MODEL-TEXT-gemma3-gemma3-for-causal-lm` `ACTIVE`, speed pending): `Gemma3ForCausalLM` on `google/gemma-3-1b-it`. W1 = two additive default-inert vt ops `kGeluAndMul` (GeGLU `gelu_pytorch_tanh`) + `kMulScalar` (bf16 embed-scale), CUDA+CPU, unit 12/12. W2 = `gemma3.{h,cpp}`/`gemma3_weights.cpp`/`gemma3_registry.cpp` reusing the GLM-4 sandwich-norm layout + `dense_attn_block.h` + FA-only KV: GemmaRMSNorm `(1+w)`, per-head Gemma q/k norm, dual per-layer RoPE theta, qpas scale, per-layer sliding window, GeGLU MLP, `sqrt(hidden)` embed-scale, tied lm_head. **SACRED gate STRICT token-exact 48/48** greedy vs vLLM 0.25.0 (K=5 ALL-DETERMINISTIC → STRICT; BOS-verified; tokenizer-free like Mistral's `LOAD-SENTENCEPIECE` path). Loader 340 tensors, registry 23/23, clean `-Werror` 0 warn. **GEMMA-2 + GEMMA-1 W3-W6 LANDED 2026-07-24 — CORRECTNESS COMPLETE** (`MODEL-TEXT-gemma2-gemma2-for-causal-lm` + `MODEL-TEXT-gemma-gemma-for-causal-lm` `ACTIVE`, speed pending): W3 = the logit soft-cap primitives (`vt::SoftCap` final cap + `PagedAttentionArgs.logits_soft_cap` attention cap threaded into the native/CPU/FA-2 attention, default-inert). W4 `Gemma2ForCausalLM` (gemma-2-2b-it) = the inverse of Gemma-3 (BOTH soft-caps, NO QK-norm, single rope) — **near-tie-band SACRED 48/48** (44/48 strict + 4/48 at 0.0-nat vLLM-own ties, 0 forward-divergent; soft-cap PROVEN applied by a cap-on≠cap-off A/B). W5 `GemmaForCausalLM` (gemma-2b) = the original Gemma (two fused norms, head_dim scale) — **STRICT 48/48**. W6 = Gemma-4 honesty pass (HW/DEP-BLOCKED, not registered). Regressions byte-identical (Gemma-3 48/48, Qwen3-dense 184/184, OPT 63/63, Llama 92/92, Mistral 92/92) + DeepSeek-V2 asserts-on 223/223; compute-sanitizer 0; clean `-Werror` 0 warn. Gemma-4 stays `BLOCKED`. | `PARTIAL` | **ACTIVE: the first additive-model bring-up = Qwen3 dense (`Qwen3ForCausalLM`) on `Qwen3-0.6B` BF16 — W0+W1 landed 2026-07-20.** W0 (config+registry stub: new TU `qwen3_dense.cpp`+`qwen3.h`, one `REGISTER_VLLM_MODEL`, full-attention-only KV spec, forward stub) + W1 (the RUNNER GENERALIZATION `ENG-RUNNER-MODELSHAPE`) are DONE and gated: dgx CUDA `-Werror` 0-warn, **27B 235/235 + 35B 315/315 token-exact UNCHANGED** (behaviour-preserving), new CPU runner tests RED(SIGSEGV)→GREEN, registry resolves `Qwen3ForCausalLM`, ASan/UBSan clean. The runner is now MODEL-SHAPE-AGNOSTIC (a full-attention-only KV config allocates+steps without the hybrid GDN path) → every future dense/non-hybrid arch adds new-files-only. Qwen3-0.6B is the only standard-dense arch with a checkpoint + runnable vLLM 0.25.0 oracle on dgx TODAY (no Llama/Mistral checkpoint present → Llama-first needs a download, sequenced as W-next for genuine cross-family additivity). **W2 loader + W3 forward LANDED 2026-07-20:** dense forward `qwen3.cpp` (`Qwen3DenseModel::Forward/ForwardDevice`) composed from vt:: ops + the fusion catalog (2 new byte-exact recipes: `kFusedAddRmsNormStd`, `kAttnQkNormRope`); bf16 attention numerics mirror vLLM. The first pure-dense bf16 model forced out + FIXED 2 genuine latent bugs: tokenizer `kQwen2Classic` (classic Qwen2/Qwen3 pre-tokenizer was hard-rejected) and `cuda_paged_attn.cu` WMMA prefill mistokenizing at head_dim≠256 (now gated to the validated d=256). **W4 CORRECTNESS COMPLETE 2026-07-20 — near-tie-robust gate PASSES on 0.6B AND a bigger 4B.** The 2026-07-20 razor's "vLLM greedy non-deterministic" premise was a BATCHING artifact: per-prompt (batch=1, the gate regime) vLLM 0.25.0 greedy is DETERMINISTIC (0.6B 0-multi/K=10, 4B 0-multi/K=5). Forward correctness is PROVEN by teacher-forcing vLLM on OUR exact prefix (`scripts/qwen3-neartie-gap.py`): at all-but-2 positions vLLM's own argmax given our prefix IS our token (gap 0.0000, bit-identical logprobs — our forward matches vLLM's prefill logits); residual flips are bf16 near-ties (0.6B ≤0.125 nats, 4B ≤0.25) where vLLM's own prefill argmax disagrees with its decode. Gate = our token within 0.5 nats of vLLM's teacher-forced argmax (strict where equal): **Qwen3-0.6B 16/16** (strict 12 + near-tie 4) and the **bigger-model complete-correctness proof Qwen3-4B (36L, GQA 32/8, hidden 2560, same forward code) 16/16** (strict 10 + near-tie 6). Regression 27B 235/235 + 35B 315/315 UNCHANGED, CUDA `-Werror` 0-warn, memcheck 0. Correctness-complete. **SPEED — d128 FA2 PREFILL + DECODE DEFAULT-ON 2026-07-20 (`Qwen3-4B` vs vLLM 0.25.0 production/graphed, in1024/out128) — big gap-close, still below vLLM, `MODEL-TEXT-qwen3-qwen3-for-causal-lm` stays `ACTIVE` NOT `DONE`:** implemented the dominant prefill lever (a d128 FlashAttention-2 varlen prefill — generalized the vendored FA2 launcher to head_dim 128, `VT_FA2_PREFILL_QWEN3` default-ON) and flipped the FA2 varlen d128 decode default ON (near-tie gate re-passes 16/16 on 0.6B + 4B). Total tput now 0.90× (c1)/0.62× (c8) (up from 0.80×/0.48×), c1 decode at parity (TPOT 1.04×, ITL P99 0.98× win); prefill A/B = +7%/+41% total, −55%/−48% TTFT. STILL failing TTFT median 5.85×/10.2× + total <1×: the full prefill STEP (not the attention kernel, now vLLM's FA2 family) is ~6× vLLM = non-attention glue (GEMM/MLP fusion) + host-side launch overhead (un-graphed prefill); plus c8 decode batch efficiency (TPOT 1.38×). Dominant residual lever = portable prefill-step fusion + graphed prefill (secondary = c8 split-KV decode occupancy). **RoPE cos/sin cache flipped DEFAULT-ON 2026-07-20** (`VT_QWEN3_ROPE_CACHE`): the opt-in blocker (an alleged FA2-split-KV-combine run-to-run nondeterminism) was GROUNDED + DISPROVEN — the paged engine is byte-deterministic run-to-run and the combine never launches on the gate (`num_splits==1`); goldens regenerated on the canonical `$HOME/cutlass-4.5.0` build (the flashinfer cutlass copy tips the 27B tok6 razor to 234/235; cutlass-4.5.0 = 235/235), gate 16/16 both, 27B 235/235 + 35B 315/315 unchanged. RoPE-ON closes total tput 0.90×→0.97× (c1) / 0.62×→0.82× (c8) and c1 TTFT ratio 5.85×→2.27×, still `ACTIVE`. **SPEED RE-BOUND 2026-07-21 (same-session, matching-recipe) — TTFT residual RESOLVED, cutlass claim CORRECTED:** the "TTFT 2.27×/5.85×" + "c8 ITL 4.3×" were BAD-DENOMINATOR/num-prompts artifacts — a fresh same-session vLLM capture gives c1 TTFT ~152 ms & c8 ITL P99 ~130 ms, and OURS WINS TTFT at both concurrencies (c1 0.90×, c8 0.38×). **c1 = effective every-axis parity** (tput 0.98× / TPOT 1.01× / TTFT+ITL wins); **c8 residual = decode** (tput 0.93× / TPOT 1.10× / ITL P99 1.12×), which nsys shows is 93% GPU-busy/compute-bound (small-M=8 `cutlass_80_wmma` projections). The **qkv-merge** (new GQA `QkvSplit` op mirroring vLLM `QKVParallelLinear`) was implemented + MEASURED NEUTRAL (doesn't cut decode FLOPs) ⇒ default-OFF. **CUTLASS CLAIM CORRECTED: 27B `test_qwen27_paged_engine` = 235/235 on the FLASHINFER cutlass build** (the "flashinfer ⇒ 234/235" was a build artifact). Stays `ACTIVE`; named residual = c8 decode-GEMM efficiency (a decode-fusion sub-campaign). **SWEEP MODEL #1 — Qwen3-Coder-30B-A3B (`Qwen3MoeForCausalLM`) W0+W1 LANDED 2026-07-21** ([sweep-qwen3-coder-30b.md](specs/sweep-qwen3-coder-30b.md), `MODEL-TEXT-qwen3-moe-qwen3-moe-for-causal-lm` `ACTIVE`, `CLAIM-MODEL-QWEN3-CODER`): the first full-attention BF16 MoE, composed from the done dense attention + the done 35B MoE experts (ZERO runner change). W0 = registry stub (`qwen3_moe_registry.cpp` + `qwen3_moe.h`, one `REGISTER_VLLM_MODEL`, full-attn-only KV, `is_dense_model=false`, W2/W3 throwing stubs). W1 = three behaviour-preserving refactors making the two done pieces reusable: (#1) dense `AttnBlock` + glue extracted to `dense_attn_block.h` (Qwen3-dense byte-identical), (#2) bf16 `MoeBlock` exposed cross-TU via `RunMoeBlock` (`qwen3_5_moe_block.h`; 35B untouched), (#3) no-shared-expert guard (inert for the 35B). Gated: dgx CUDA `-Werror` 0-warn; Qwen3-dense 0.6B+4B 16/16 + 27B 235/235 + 35B 315/315 UNCHANGED; registry resolves `Qwen3MoeForCausalLM`; memcheck 0. W2 bf16 loader → W3 forward → W4 near-tie token-exact → W5 fast bf16 grouped-MoE GEMM remain. Then Llama dense (download), Mistral, MoE families **SWEEP MODEL — GLM-4-9B-0414 (`Glm4ForCausalLM`) G2 LANDED 2026-07-24 — CORRECTNESS COMPLETE** ([glm-dsa-latest-deepseek](specs/glm-dsa-latest-deepseek.md), `MODEL-TEXT-glm4-glm4-for-causal-lm` now `READY` per the [live-state audit](specs/live-state-audit-2026-08-06.md), `CLAIM-GLM-DSA-LATEST-DEEPSEEK` amended) **GLM-4.7-Flash (`Glm4MoeLiteForCausalLM`, 31.2B MLA+MoE) G1 LANDED 2026-07-24 — SACRED gate 8/8, CORRECTNESS COMPLETE** (`MODEL-TEXT-glm4-moe-lite-glm4-moe-lite-for-causal-lm` `ACTIVE`, speed pending): reuses the DeepSeek-V2 MLA stack; first e2e coverage of the q_lora branch + noaux_tc sigmoid router, closing the MLA campaign's two C2 coverage gaps: the FIRST GLM-family model. SACRED gate 16/16 vs vLLM 0.25.0 (STRICT 13/16 + near-tie 3/16, max gap 0 nats; vLLM K=5 self-deterministic ⇒ STRICT bar), speed PENDING. The two "new primitives" the spike named reduced to EXISTING infra: partial + interleaved `RopeFromCache` (`is_neox_style=false`, the DeepSeek-V2 decoupled-rope path) over `rotary_dim=64`; standalone `vt::RmsNorm` sandwich norms. Biased qkv, no QK-norm, GQA 32/2, untied lm_head. New files + one REGISTER, reusing the shared dense glue. **SWEEP MODEL — Llama-3.2 (`LlamaForCausalLM`) W0-W4 LANDED 2026-07-23 — CORRECTNESS COMPLETE** ([sweep-llama-3.2](specs/sweep-llama-3.2.md), `MODEL-TEXT-llama-llama-for-causal-lm` `ACTIVE`, `CLAIM-MODEL-LLAMA-3.2`): the roadmap's explicit "Llama-first" increment and the first mainstream non-Qwen/non-OPT dense arch. `LlamaForCausalLM` (Llama-3.2-1B) = the Qwen3-dense forward with exactly two ADDITIVE deltas — NO qk-norm (shared `AttnBlock` skips it when q_norm/k_norm empty) + llama3 rope-scaling (4 default-0 `RopeArgs` fields + a `Llama3ScaleFreq` kernel helper, no-op elsewhere; formula verified 2e-7 rel vs vLLM) — reusing the shared dense forward VERBATIM (`LlamaModel == Qwen3DenseModel`). 3 new Llama files, ZERO edit to runner/scheduler/platforms/attn-registry/`hf_config`/any qwen3-opt model. vLLM 0.25.0 greedy MEASURED DETERMINISTIC (K=6, 0 multi-valued cells) ⇒ STRICT token-exact bar, PASS **16/16 (12 strict + 4 near-tie band, max gap 0.0000 nats, 0 divergent)** — at all 13 divergent positions vLLM's own teacher-forced argmax given our prefix IS our token. A correctness-fatal tokenizer bug (Llama's `Sequence` post_processor wrapping `TemplateProcessing` ⇒ BOS 128000 never prepended, silently 1/16) was isolated by a CUDA prefill-argmax diagnostic (forward proven 4/4 correct given vLLM's exact tokens) and fixed byte-preservingly (Qwen/OPT/DeepSeek unaffected — all ByteLevel/top-level-TemplateProcessing). Regressions 27B 235/235 · 35B 315/315 · Coder 6/6 · Qwen3-dense 16/16 · OPT 6/6 · DeepSeek-V2 8/8 UNCHANGED; `-Werror` 0-warn; memcheck 0; DSR 67. SPEED pending (head_dim 64 → generic paged path; Llama-3.2-3B head_dim-128 is the FA2-toggle W-next). **MLA CAMPAIGN W0+W1 LANDED 2026-07-21** (`CLAIM-MLA-DEEPSEEK`; rows STAY `SPIKE` — W0/W1 make no model supported). **W0 grounded every fact the spike flagged as an unverified source read; ALL CONFIRMED, none contradicted:** DeepSeek-V2-Lite fetched to dgx (30 GB, 4 shards) and loading in the vLLM 0.25.0 oracle; the oracle's own DEBUG startup on sm_121 prints `Using TRITON_MLA attention backend out of potential backends: ['TRITON_MLA']` and `Using FLASH_ATTN MLA prefill backend` — so the dense-MLA decode + MLA-prefill targets are OBSERVED, not inferred, and the sm90/sm100-only MLA kernel class stays out of scope; the real `config.json` confirms every §5.1 number (`kv_lora_rank=512`, `qk_nope=128`, `qk_rope=64` -> the **576-wide latent**, `v_head_dim=128`, `q_lora_rank=null`, `n_group=topk_group=1`, softmax/greedy, 64+2 experts, 27 layers) plus `is_neox_style=False` and the mscale2 scale correction; and BOTH recorded coverage gaps (no `fused_qkv_a_proj` branch, no `e_score_correction_bias`) are confirmed real, so those pieces stay unit-gated only. **W1 = the behaviour-preserving spec-driven KV allocation, ZERO MLA math:** the attention cache is now sized `num_blocks * spec->page_size_bytes()` and viewed from the spec's own `block_size`/`num_kv_heads`/`head_size`/`dtype` instead of the hardcoded `2 * block * Hkv * Dh` with shape reconstructed from the HF config (`runner.cpp`), plus `MLAAttentionSpec` with upstream's factor-1 single-tensor page formula (`kv_cache_interface.py:397-398`) registered against the ORDINARY `FullAttentionManager` (`single_type_kv_cache_manager.py:1539`) — the spike's key finding, which is why block manager/prefix caching/eviction need no change. Gated: dgx clean CUDA `-Werror` 0 warnings/0 errors; **27B 235/235 + 35B 315/315 + Qwen3-Coder 6/6 + Qwen3-dense 16/16 ALL UNCHANGED** (behaviour-preserving proven, not assumed); `test_runner` 15/15, `test_kv_cache_interface` 21/21 (4 new MLA-spec cases), `test_llm_engine` 5/5; the new path is proven EXERCISED (not merely compiled) by `fa_page_size_bytes()` + a `page_size_padded` case no HF-config formula can produce. **W2 + W3 LANDED 2026-07-21** (base `a05437f`; rows STAY `SPIKE` — still no MLA attention math, no MLA model, no forward). **W2 = the MLA branch of `_get_backend_priorities` the pre-W2 comment deferred, ported as DATA:** the whole of `cuda.py:84-176` (BOTH branches — MLA sm_100 including the `:96-115` adaptive sparse tail, MLA sm_12x, MLA `else`, and the two pre-existing non-MLA arms) is now a TABLE in the new header `include/vllm/platforms/cuda_attn_priority.h`, one row per upstream arch arm keyed on `(use_mla, major)`, so a future arch is a ROW rather than a code path; putting it in a header (not the CUDA-only TU) also let the CPU test tier assert the REAL table and DELETED the hand-copied `FakeCudaPlatform` duplicate. On sm_121 a `use_mla=true` request now RESOLVES to `TRITON_MLA`, matching the W0 oracle observation. **The sparse/DSA seam is left OPEN and unit-proven:** GB10's row keeps both upstream entries and the sparse one loses to a real FILTER — `AttentionBackend::is_mla()`/`is_sparse()` checked against the request (`backend.py:307-360 validate_configuration`) — so a future DSA backend is selected purely by declaring `is_sparse() == true`, with ZERO edit to the table or the selector. `TritonMLABackend` lands the NAME plus upstream's 3-D `get_kv_cache_shape` (no K/V axis; `num_kv_heads != 1` REFUSED), `get_impl_cls()` deliberately still `nullptr`. MLA prefill priority ported too (GB10 -> `[FLASH_ATTN]` alone). **W3 = the two new `vt::` ops, both CPU-reference-gated.** `vt::ConcatAndCacheMla` mirrors `csrc/libtorch_stable/cache_kernels.cu:401-442` — and per the whole-chain rule this was VERIFIED, not assumed, to be vLLM's OWN csrc kernel (`_custom_ops.py:2532` -> `torch.ops._C_cache_ops`), with no flashinfer/cutlass variant in the dense-bf16 path; it concatenates the latent + rope part into ONE 576-wide entry, the write `ReshapeAndCache`'s K/V-pair signature cannot express, stride-driven so a per-layer cache slice and the two column halves of `kv_a_proj_with_mqa` both work copy-free. The **grouped-topk (`noaux_tc`) router** extension — flagged in `coordination.md` as SHARED with `CLAIM-GLM-DSA-LATEST-DEEPSEEK` and "must not be implemented twice" — is landed HERE and that claim now consumes it: additive `MoeRouterTopKArgs` fields + an optional `e_score_correction_bias`, with `num_expert_group == 0` still dispatching the ORIGINAL kernel so the 27B/35B/Coder/dense routers are byte-identical BY CONSTRUCTION. **Stated plainly: the `noaux_tc` correctness evidence is UNIT-ONLY.** V2-Lite has `n_group=topk_group=1` and no bias, so the e2e vehicle exercises none of it; the gate is `tests/vt/test_ops_moe_router_grouped.cpp` at DeepSeek-V3's REAL dimensions (256 experts, n_group=8, topk_group=4, sigmoid, scaling 2.5, WITH the bias) against an INDEPENDENT sort-based transcription of the upstream formula. **W4 LANDED 2026-07-22** (base `ed2c342`; rows STAY `SPIKE` — W4 adds a kernel and fills a `nullptr`, it makes no model supported). **`vt::MlaDecodeAttention` — the MQA decode over the compressed latent (QK 576 / V 512, `num_kv_heads=1`), a structure port of the two-stage split-KV pair W0 OBSERVED EXECUTING:** `MlaDecodeStage1` <- `_fwd_grouped_kernel_stage1` (`triton_decode_attention.py:278-458`, the `IS_MLA` branch whose `v = tl.trans(k)` at `:424-431` is the whole MLA trick — V is the leading 512 columns of the SAME latent row already loaded as K, so one shared-memory tile serves as both), `MlaDecodeStage2` <- `_fwd_kernel_stage2` (`:575-639`), `ComputeNumKvSplits` <- `_compute_num_kv_splits` (`triton_mla.py:40-47`), split workspace <- `_reserve_attn_logits_workspace` (`:57-78`) realized as the house grow-only per-stream scratch. **Honest reuse verdict:** our FA-2 split+combine machinery fit at the ALGORITHM level (the split schedule, the LSE merge algebra, the fixed-ascending no-atomicAdd determinism rule) and NOT at the code level — the vendored FA-2 launcher takes separate 4-D k/v caches and is instantiated for symmetric head_dim {128,256}, which cannot express a 3-D single-buffer cache with QK 576 / V 512; that is recorded in the TU header rather than forced. **Evidence is unit-level and deliberately strong** (there is no e2e model until W7): [`tests/vt/test_ops_mla_attn.cpp`](../tests/vt/test_ops_mla_attn.cpp), a port of `tests/kernels/attention/test_mla_decode_cpu.py` whose `ref_mla` becomes an INDEPENDENT TWO-PASS oracle (a different algorithm from the streaming online-softmax both impls use) plus its NaN-padding out-of-bounds detector, run at the REAL V2-Lite geometry (576/512/64, block 16, mscale^2 scale) over ragged / multi-block / single-block / every split boundary (`num_kv_splits` ∈ {1..512} incl. splits > seq_len) / 128-head V3 / non-BLOCK_H head counts / a 288-256 block-32 geometry / bf16 + f32, with run-to-run BIT-exactness. dgx sm_121: 11/11 cases, 2,303,193 assertions; `compute-sanitizer` memcheck **0 errors**, racecheck **0 hazards**, synccheck **0 errors**; clean CUDA build **0 warnings / 0 errors**; **27B 235/235 + 35B 315/315 + Coder 6/6 + Qwen3-dense 16/16 + OPT 6/6 ALL UNCHANGED**. `TritonMLABackend::get_impl_cls()` is no longer `nullptr` — it returns a real `TritonMLAImpl` whose `forward_mqa` is the 1:1 counterpart of `triton_mla.py:189-260`; PREFILL is W5 and `forward()` refuses a prefill-shaped batch BY NAME rather than producing wrong numbers. NO speed number — decode perf is W9. **W5 LANDED 2026-07-22** (base `5395203`; rows STAY `SPIKE`). **MLA PREFILL + the workspace-bounded CHUNKED-CONTEXT loop.** Three new ops — `vt::MlaPrefillAttention` (<- `mla/prefill/flash_attn.py:153-248`, the ONLY MLA prefill backend reachable on sm_121 and the one W0 OBSERVED the oracle logging), `vt::GatherMlaCache` (<- `cache_kernels.cu:992-1064`) and `vt::MergeAttnStates` (<- `merge_attn_states.cu:18-192`, both `-inf` edge cases verbatim) — plus the loop itself in the new `mla_chunked_context.h` (<- `mla_attention.py:1422-1451,1667-1745,2094-2199,2344-2425`), which is what keeps a long-context prefill inside a bounded workspace instead of materializing a 3 GB up-projected context. **The vendored FA-2 launcher WAS generalized, and W4's prediction that it would be tractable held for a reason worth recording: upstream does not ask FA-2 for asymmetric head dims either.** `requires_v_padding` is TRUE on GB10, so upstream ZERO-PADS V from 128 to 192 and slices the output back — the kernel stays a plain SYMMETRIC head_dim-192 instantiation. The whole change is two new explicit instantiations of an UNCHANGED generic template, one new launcher entry for the contiguous-varlen mode, and the pad/slice pair; the paged launcher every non-MLA prefill calls is textually untouched (211 insertions / **0 deletions**), and 27B 235/235 + 35B 315/315 + Coder 138/138 + Qwen3-dense 664/664 + OPT 36/36 are all UNCHANGED. Evidence is UNIT-ONLY (there is still no model): 4/4 cases / **2,377,052 assertions** and 5/5 / **306,037 assertions** on dgx sm_121 at the real QK 192 / V 128 geometry, against an INDEPENDENT double-precision two-pass oracle and — for the loop — a SINGLE-SHOT whole-sequence oracle that never chunks, over exact / +1 / -1 chunk boundaries, zero-context and zero-key-in-chunk requests, ADVERSARIAL reverse-interleaved block tables, NaN-poisoned outputs and run-to-run bit-exactness; memcheck **0**, racecheck **0 hazards**, synccheck **0**. A genuine upstream FA-2 quirk was found on the way and worked around rather than papered over: its EMPTY-K early exit ignores the unpadded-LSE flag, which a zero-key chunk request would turn into an out-of-bounds LSE write. **W6 LANDED 2026-07-22** (base `2846467`; rows STAY `SPIKE` — W6 adds an attention LAYER, not a model). **The MLA attention BLOCK + LOAD-TIME WEIGHT ABSORPTION — the piece that finally COMPOSES W3's cache write, W4's MQA decode and W5's MHA prefill into one layer:** the projections with BOTH `q_lora_rank` branches (`fused_qkv_a_proj` -> `q_a_layernorm` -> `q_b_proj`, or the direct `q_proj`), the two RMSNorms (the rope part deliberately NOT normed), the DECOUPLED RoPE (`is_neox_style=False`, only the trailing 64-dim slice rotates) with its YaRN cos/sin cache and the SEPARATE mscale^2 softmax-scale correction, the `kv_b_proj -> W_UK/W_UV` split, the prefill-MHA / decode-MQA dispatch with decode tokens packed FIRST, and the `kv_b_proj` up-projection callback W5 left open. **The spike's most useful prediction held: absorption needed NO new attention kernel** — it is a LOAD-TIME weight transform plus TWO batched GEMMs, so the entire new-kernel surface is two general primitives, `vt::BatchedMatmul` (<- `torch.bmm` at `mla_attention.py:789,1034`; on CUDA torch resolves that to cuBLAS `gemmStridedBatchedEx`, and ours is the cuBLASLt strided-batched form of the same GEMM) and `vt::ConcatMlaNopeRope` (<- `concat_mla_q`, generalized so one op also serves `_concat_k_nope_k_pe`). **The absorbed-vs-unabsorbed equivalence — the heart of W6 — is PROVEN NUMERICALLY, three independent ways, rather than argued:** an INDEPENDENT double-precision block oracle computes the attention BOTH ways and agrees to < 1e-11 (the identity itself); our absorbed decode reproduces the UNABSORBED oracle to < 2e-4 in f32; and — the strongest — the SAME batch is driven once through our ABSORBED MQA decode kernel (QK 576 / V 512, one KV head, K/V never materialized) and once through our UNABSORBED materialized-MHA prefill path (QK 192 / V 128 plus the chunked-context loop), agreeing to < 3e-4 (CPU f32) / < 4e-2 (CUDA bf16) with nothing but the weights shared between them. Evidence on dgx sm_121: `test_mla_attention_block.cpp` 10/10 cases / 2,372,644 assertions and `test_ops_mla_absorb.cpp` 9/9 / 1,644,807 (CUDA cases proven to EXECUTE; NaN-poisoned outputs; run-to-run BIT-exact), porting `tests/kernels/test_concat_mla_q.py` in both arms. memcheck 0, racecheck 0 hazards, synccheck 0 (the last needing `--num-cuda-barriers 65536` — the default table overflows on a binary driving this many kernel families and the tool then reports a bogus launch failure, a trap worth knowing). Clean CUDA build 0 warn / 0 err; **27B 235/235 + 35B 315/315 + Coder 138/138 + Qwen3-dense 664/664 + OPT 36/36 ALL UNCHANGED**. **Coverage stated plainly: the `q_lora` query branch has NO e2e coverage and cannot get any on GB10** — DeepSeek-V2-Lite has `q_lora_rank=null`, so it is unit-gated at DeepSeek-V3's real dimensions only; GLM-4.7-Flash (`q_lora_rank=768`, 58.2 GiB, fits) is what would close it. **W7 LANDED 2026-07-22** (base `ce43c51`; the row STILL stays `SPIKE`). **THE DEEPSEEK-V2 MODEL — registry + config parse + loader + forward: the first MLA model in the tree, and the first one that runs a real MLA checkpoint end to end.** Four new files plus ONE shared-code edit (a two-line additive condition in `runner.cpp` recognising a `kMlaAttention` KV group as the model's attention group — upstream maps MLA onto the ordinary `FullAttentionManager`, so block tables/prefix caching/eviction are untouched). **LOADER GATE PASSED on the real 4-shard DeepSeek-V2-Lite: 5291/5291 checkpoint tensors accounted for, ZERO unmapped and ZERO leftover** (4/4 cases / 37,331 assertions), every shape asserted including the LOAD-TIME `kv_b_proj -> W_UK_T [16,128,512]` / `W_UV [16,512,128]` absorption split — the same transform, at the same lifecycle point, as upstream's `process_weights_after_loading`. **V2-Lite takes the DIRECT `q_proj` query branch** (`q_lora_rank: null`), asserted with the fused branch EMPTY on every layer. **FORWARD GATE PASSED and obviously right, not merely finite: the real checkpoint prefill of `The capital of France is` -> argmax ` Paris`** (top-5 ` Paris`/` the`/` a`/` one`/` also`, run-to-run bit-exact) — the direct analogue of the Qwen3-Coder W3 sanity case. **BATCH-ORDERING GATE:** the ordering invariant W6 measured 0.86 relative error from is now VALIDATED, not assumed — `BuildMlaBatchSplit` throws (naming the request and citing the upstream line) if a decode follows a prefill or a with-context prefill follows a context-free one. **SHARED EXPERTS — new for this family and UNGATED unlike Qwen3.6's sigmoid-gated one — gated two ways:** a MoE layer with every routed expert zeroed is BIT-IDENTICAL to a dense layer holding the same MLP, and turning the shared expert off CHANGES the logits. **The CUDA path is EXERCISED, not merely compiled:** a case at the real MLA head geometry drives the CUDA MLA kernels and the CUDA-only grouped bf16 MoE GEMM, bit-exact on device and within 0.0061 worst relative logit error of the CPU reference path. 11/11 forward cases; memcheck/racecheck/synccheck all **0**; clean CUDA build **0 warn / 0 err**; **regression set UNCHANGED**. **Only `DeepseekV2ForCausalLM` is REGISTERED** — `DeepseekForCausalLM` (plain MHA), V3 (fp8/671B) and V3.2 (DSA indexer) are REFUSED BY NAME in the config parse rather than falsely claimed. A pre-existing tree-wide hazard was found on the way and recorded: the shared `DevicePool` is a process-wide singleton keyed only on a byte size class, so a single process driving BOTH a CPU and a CUDA forward hands the second backend the first's recycled pointers. **NEXT: W8 — the SACRED token-exact gate on DeepSeek-V2-Lite** (wire the paged engine to produce the MLA batch order the model already validates, capture oracle goldens, run the STRICT form W0 determined). A loading, forwarding model is NOT a supported model, so no model row moves until that gate passes. **W8 LANDED 2026-07-22 — THE SACRED CORRECTNESS GATE PASSES 8/8, and `MODEL-TEXT-deepseek-v2-deepseek-v2-for-causal-lm` moves `SPIKE` -> `ACTIVE` (correctness COMPLETE, speed PENDING). NOT `DONE` — that additionally requires vLLM-speed parity on every axis, which is W9 and has NO number yet; the other four campaign rows stay `SPIKE`.** An 8-prompt battery is driven through the FULL paged `LLMEngine` and compared to the pinned vLLM 0.25.0 oracle: **8/8 PASS — STRICT token-exact 5/8, near-tie band 3/8, 92/128 tokens strictly exact, max teacher-forced gap 0.25 nats, 0 forward-divergent** (223 assertions). **The bar was ARRIVED AT by measurement, not chosen:** vLLM is DETERMINISTIC on this model at batch=1 (W0's K=5 8/8, re-confirmed by W8's own capture at T=16 with 0 multi-valued cells), so the STRICT form ran FIRST and came out 5/8; the ratified TEACHER-FORCING diagnostic then showed **36 divergent positions with 35 at gap EXACTLY 0.0000 nats** — vLLM's own argmax GIVEN OUR PREFIX is our token, so they are the downstream tail of one earlier flip — **exactly ONE root flip with any gap at all (prompt[3] tok 9, 0.2500 nats, inside the ratified 0.5-nat band and equal to the landed Qwen3-dense 4B gate's worst)**, and **ZERO tokens outside vLLM's top-20**, with the per-position nats COMMITTED as goldens and anything beyond the band still FAILING. **W8's first job — the scheduler/runner wiring — turned out to need NO new code, for a non-accidental reason:** `runner.cpp:671` already reorders with `decode_threshold = 1`, exactly MLA's `reorder_batch_threshold` (`mla_attention.py:1420`), and its `decode -> short_extend -> long_extend -> pure_prefill` ordering satisfies BOTH MLA invariants (decodes form a batch prefix; with-context prefills lead the prefill tail). W8 PROVES that end to end rather than duplicating it, with new DIAGNOSTIC `MlaBatchSplitStats` counters and a non-vacuity bar: the battery is admitted CONCURRENTLY with staggered arrival, producing **7 genuinely MIXED decode+prefill steps at up to 8 concurrent requests** with `BuildMlaBatchSplit` (which throws naming the request) never firing, plus a prefix-cache-driven **with-context prefill**, and a phase-0 check that the engine really allocated the compact MLA cache (`fa_page_size_bytes = 36864`, no factor 2). **THE REAL BLOCKER WAS THE TOKENIZER, NOT THE MODEL:** the first run REFUSED to load (`unsupported normalizer "Sequence"`), and behind it sat a whole NEW pre-tokenizer family — DeepSeek's is a HF `Sequence` PIPELINE of SEVEN stages (five `Split(Isolated)` over ENUMERATED codepoint ranges, then `Digits(individual_digits=true)`, then `ByteLevel(use_regex=false)`), whose stage ORDER is load-bearing because stage 2's punctuation class spans 0x3A-0x7E and CONTAINS A-Z/a-z. Landed as `SplitPattern::kDeepSeek` with the five patterns compared VERBATIM at load, and MEASURED token-for-token against the REAL HF `tokenizers` library over a stage-stress corpus (**6/6 cases / 2461 assertions**). **The TOKENIZATION goldens earned their keep by REFUTING a fix that was already written:** `tokenizer_config.json` declares `add_bos_token: true`, which reads as exactly the OPT missing-BOS bug — but vLLM's resolved tokenizer (`TokenizersBackend`) adds NO BOS, our loader already matched bit-for-bit, and the "fix" would have BROKEN a passing gate; it was reverted and the measured behaviour PINNED by a guard case ([[ground-premises-before-dispatching]]). Regression set UNCHANGED (27B 235/235, 35B 315/315, Coder 6/6, Qwen3-dense 16/16, OPT 6/6, plus every tokenizer test — W8 touches SHARED tokenizer code, so that was proved, not assumed); clean CUDA rebuild 0 warn/0 err; local CPU suite 151/151; memcheck/racecheck/synccheck 0. Batch invariance is REPORTED (6/8) and deliberately NOT a bar, because the ORACLE itself changed on 3/8 of this battery under batched generation (W0). One W9 input recorded: the oracle must run `moe_backend='triton'` — vLLM's auto-selected FlashInfer CUTLASS unquantized MoE REBOOTED dgx three times on GB10's unified memory. **W9 SPEED CLOSE LANDED 2026-07-22 — the track has its FIRST binding speed number, and it is an ATTRIBUTED MISS: `MODEL-TEXT-deepseek-v2-deepseek-v2-for-causal-lm` STAYS `ACTIVE` (correctness COMPLETE, speed SHORT), NOT `DONE`** ([grid](../docs/BENCHMARKS.md), [spike §W9](specs/mla-deepseek-campaign.md)). Denominator SETTLED with evidence — CUTLASS MoE has now rebooted dgx **five times** (two more at W9, the second on a pristine box with a 0 GiB page cache and every mitigation applied, both deaths at the identical post-`torch.compile` phase), so `--moe-backend triton` IS vLLM's best STABLE GRAPHED configuration here and is the bar; the substitution does not flatter us, we lose to it. `nsys` (both sides, `--cuda-graph-trace=node`) overrode the plan: the lever was not the planned MLA fusion recipes but `MlaDecodeStage1` sitting at **44.7% of all GPU time and ~180x off its own memory-bound floor** on a **2-CTA grid at batch 1**; applying upstream's own occupancy target made it **18.3x faster** (837 -> 45.8 us) for **+69.5%/+53.3%/+32.0%/+19.5%** end-to-end at c1/c2/c4/c8, while the planned decode-graph sibling is worth only ~+2% (this decode is GPU-bound). Grid vs vLLM: output throughput **0.87/0.95/0.86/0.88** (was 0.50 at c1), TTFT **1.06/1.14/0.96/0.88** (we WIN at c4/c8), TPOT **1.11/0.97/1.16/1.17**. SACRED gate **8/8 UNCHANGED** with both levers default-ON; a real latent CUDA-graph use-after-free in the MLA metadata upload was found and fixed (its whole class now guarded); regression set UNCHANGED; clean rebuild 0 warn/0 err; sanitizers 0. **NEXT LEVER, NAMED: route the batch-1 dense projections off cuBLAS `gemvx` (31.8% of our GPU time) onto a tensor-core GEMM — vLLM splits the same work `gemvx` 12.7% + `nvjet_sm121_tst_mma_*` 6.6%.** **W10 BLOCKED-ROW HONESTY PASS LANDED 2026-07-22 — the campaign's W-plan is COMPLETE; records only (no code, no build, no GPU work, nothing downloaded, no number claimed).** Rows set to their final honest state: `MODEL-TEXT-deepseek-v2-deepseek-v3-for-causal-lm` (V3 + V3.2, and Kimi-K2/K2.5's text backbone by config composition) and `MODEL-TEXT-minimax-m2-mini-max-m2-for-causal-lm` move `SPIKE` -> `BLOCKED`, joined cross-claim by `MODEL-TEXT-deepseek-v2-glm-moe-dsa-for-causal-lm` (GLM-5) under `CLAIM-GLM-DSA-LATEST-DEEPSEEK`; each is HW-BLOCKED on 119 GiB (~642 GiB fp8 / ~428 GiB / 1404 GiB) and the two DSA models are additionally DEP-BLOCKED — for a SPARSE model the XOR filter eliminates `TRITON_MLA`, leaving `FLASHINFER_MLA_SPARSE_SM120` alone, whose sm12x dispatch goes to flashinfer's DENSE-ONLY XQA backend that discards `sparse_mla_top_k` (upstream's own test monkeypatches the probe and asserts nothing numerical). `MODEL-TEXT-deepseek-v2-deepseek-for-causal-lm` stays `SPIKE` with the record repaired to say it is plain MHA and needs NO MLA; Kimi-Linear stays `SPIKE` (MLA half unlocked, KDA a separate kernel campaign, HW-MARGINAL). Each blocked row states what CAN still be gated (config resolution, weight-map on a slice, unit parity at real dimensions) versus what CANNOT (anything e2e). **Two PERMANENT coverage gaps now stated in the rows:** the `noaux_tc` grouped router and the `q_lora` query branch have NO e2e coverage and are unit-gated only, because V2-Lite is `n_group=topk_group=1`/softmax with no `e_score_correction_bias` and `q_lora_rank=null`. **NAMED NEXT VEHICLE: GLM-4.7-Flash** (`MODEL-TEXT-glm4-moe-lite-glm4-moe-lite-for-causal-lm`, 31.2B / 58.2 GiB, FITS GB10) — the only reachable checkpoint that closes BOTH gaps. **BLOCK NOT CLOSEABLE, nothing archived:** the DeepSeek-V2 row is `ACTIVE`, not `DONE`, so the plan/spec stay LIVE; the single open item is the named `gemvx` -> tensor-core dispatch lever. **MISTRAL FIFTH FAMILY W0-W3 LANDED 2026-07-23** ([sweep-mistral](specs/sweep-mistral.md), `MODEL-TEXT-mistral-mistral-for-causal-lm` `ACTIVE`, `CLAIM-MODEL-MISTRAL`): the closest-to-Llama dense arch (vLLM `mistral.py` = "Mistral adaptation of the LLaMA architecture") — plain rope θ1e6 (no rope_scaling) + qk-norm-optional + untied lm_head + null sliding_window, all PRE-EXISTING ⇒ NO new primitive, 3 new files + additive CMake/registry-test rows only, ZERO shared-code edit. **MODEL forward gate 30/30 greedy tokens vs vLLM 0.25.0** (tokenizer-free: fed vLLM's exact prompt tokens through our CUDA prefill; 29 STRICT token-exact + 1 near-tie, 0 forward-divergent; vLLM greedy det 4/5 K=3). W2 loader real-weights 1541 assertions. **REAL FINDING:** Mistral's SentencePiece/Metaspace tokenizer is unsupported by our ByteLevel-BPE tokenizer → the FULL paged-engine SACRED gate is BLOCKED, the pre-inventoried `LOAD-SENTENCEPIECE` row (SentencePiece tokenizer family). `-Werror` 0-warn, DSR 32, regressions UNCHANGED (Llama paged 16/16, Qwen3-dense forward 1031, registry 299; MoE/GDN gates unaffected by construction). SPEED + full paged gate both PENDING (row `ACTIVE`, not `DONE`). **OLMo-2 SPIKED 2026-07-24** ([sweep-olmo2](specs/sweep-olmo2.md), `CLAIM-SWEEP-OLMO2`): one row `INVENTORIED` → `SPIKE` (`MODEL-TEXT-olmo2-olmo2-for-causal-lm`, covering `Olmo2ForCausalLM` + its `Olmo3ForCausalLM` alias). Answers the breadth-sweep §B.3 Tier-2 rank-8 "GLM4 / Olmo2-3" item (GLM-4 + Gemma landed; OLMo-2 next). **HEADLINE: OLMo-2 is the cleanest dense bring-up yet — ZERO new compute kernels.** The two distinctive facts both reduce to WIRING over landed ops: (1) the **pure post-norm (`norm_after`) placement** is a strict SUBSET of the GLM-4/Gemma sandwich (keeps ONLY the standalone-output-norm op `glm4.cpp:174-178` — the exact primitive flagged — DROPS the pre-norms, plain residual add); (2) the **QK-norm is FULL-WIDTH not per-head** → reuses `vt::RmsNorm` at a `[T,q_size]`/`[T,kv_size]` shape but CANNOT use the fused per-head `kAttnQkNormRopeGate`. Everything else REUSES (plain RMSNorm, SiLU SwiGLU, NeoX rope, GQA paged glue, Gemma-3 sliding-window for Olmo-3, tied embeddings, packed loader, ByteLevel BPE). Gate vehicle `allenai/OLMo-2-0425-1B` (1.485B, ~2.77 GiB, fits GB10 tight ~30 GiB free); Olmo-3 rides the same row (0.25.0 oracle-support UNVERIFIED). OLMo-1 (non-parametric LayerNorm), OLMoE/FlexOlmo (MoE), OlmoHybrid (SSM) stay `INVENTORIED`. Nothing implemented. | -| 2a | `ROAD-V1-C2-LOCAL-BF16` | Close local discrete-Blackwell Qwen3.5 plain-BF16 production parity | [`MODEL-MM-qwen3-5-qwen3-5-for-conditional-generation`](model-matrix.md), [`LOAD-SAFETENSORS-DIRECT-DENSE`](engine-matrix.md), [`SERVE-CLI-BENCH`](engine-matrix.md), [`KERNEL-SSM-MAMBA`](kernel-matrix.md) | Exact `(sequence, 8-token chunk)` GDN conv dispatch is byte-identical and default ON. Rebased-main same-binary reprofile: **718.704→233.955 ms (3.072x)** and whole run **+2.272%**; sealed vLLM remains 145.421 ms (**1.609x residual**). The binding three-pair A/B improved total/output **2.152%**, TTFT **2.945%**, TPOT/ITL **1.920%**, no VRAM regression. Against sealed vLLM: throughput **1.021246x PASS**; TTFT, TPOT and VRAM OPEN | `GATING` | Spike/profile the residual causal-conv gap. Latency, VRAM and 27B/35B correctness remain gates; no 4B-to-gate-model extrapolation. Evidence: [production baseline and exact-chunk outcome](../docs/bench-evidence/qwen35-4b-sm120-main-20260807.md), [conv spike/result](specs/sm120-qwen35-conv-chunking-2026-08-07.md) | +| 2a | `ROAD-V1-C2-LOCAL-BF16` | Close local discrete-Blackwell Qwen3.5 plain-BF16 production parity | [`MODEL-MM-qwen3-5-qwen3-5-for-conditional-generation`](model-matrix.md), [`LOAD-SAFETENSORS-DIRECT-DENSE`](engine-matrix.md), [`SERVE-CLI-BENCH`](engine-matrix.md), [`KERNEL-SSM-MAMBA`](kernel-matrix.md) | Exact `(sequence, 8-token chunk)` GDN conv dispatch is byte-identical and default ON. Rebased-main same-binary reprofile: **718.704→233.955 ms (3.072x)** and whole run **+2.272%**; sealed vLLM remains 145.421 ms (**1.609x residual**). The reviewed combined default-off K4 causal-conv + 16-token post-conv local A/B is exact and improves total throughput **0.7175%**, TTFT **1.340%**, TPOT **0.556%** and E2E **0.716%**. Prior cross-engine throughput and TTFT ratios are **VOID/PENDING** because only local timed string tokenization; TPOT and VRAM remain OPEN | `GATING` | Implement the committed pretokenized benchmark frontend with `VT_BENCH_PRETOKENIZE=0` rollback, review, then rerun local/vLLM. Latency, VRAM and 27B/35B correctness remain gates; no 4B-to-gate-model extrapolation. Evidence: [campaign spec](specs/sm120-qwen35-pareto-2026-08-09.md), [production baseline](../docs/bench-evidence/qwen35-4b-sm120-main-20260807.md) | | 3 | `ROAD-V1-C3` | MTP k=1 + GDN speculative path, then DFlash, DSpark and heterogeneous-vocabulary TLI | [engine matrix](engine-matrix.md), [coverage view §8](feature-matrix.md#8-speculative-decoding) | MTP and DFlash specs exist. **M-mtp-0 CLOSED 2026-07-24** - the standalone MTP draft head is oracle-parity-proven on BOTH gate checkpoints (op-level parity vs a dumped k=1 vLLM oracle, not a token-generation SACRED gate). **I2 SCHEDULER-HALF LANDED 2026-07-24** ([mtp-spec-decode §2.7](specs/mtp-spec-decode.md)): host-side spec-decode scheduler/engine plumbing + the FROZEN spec-metadata ABI that I3 (rejection sampler) and I5 (verify/propose runner) build against - `SpeculativeConfig`, `DraftTokenIds`, `Request::spec_token_ids`/`NumTokensWithSpec`, the first population of `scheduled_spec_decode_tokens`, `Scheduler::update_draft_token_ids`, the `take_draft_token_ids` seam, `EngineCore::post_step`, `InputBatch::num_accepted_tokens`/`update_req_spec_token_ids`; DEFAULT-OFF and INERT (no `SpeculativeConfig` => `num_lookahead_tokens == 0` => byte-identical engine). `SPEC-MTP` **STAYS `GATING`** because M-mtp-1..4 (greedy rejection, GDN spec slots, k>1, CUDA graphs) are still open, so spec decode remains user-invisible. DSpark is user-promoted scope with DeepSeek-V4/Qwen3 draft models, reduced-vocabulary handling and full-CUDA-graph behavior inventoried under `SPEC-DSPARK`; tokenizer-agnostic target<->draft mapping is separately inventoried as `SPEC-TLI`. Their dedicated spikes are not written **I3 GREEDY REJECTION SAMPLER LANDED 2026-07-24** (`SPEC-REJECTION` `READY` -> `ACTIVE`): per-request logits expansion to `1 + k_i` rows plus the greedy accept rule (accept a draft iff it equals the target argmax; on the first mismatch emit the target argmax and stop; bonus token when all k accept), CUDA==CPU bit-exact at vocab 248320. **I4 GDN-HALF LANDED 2026-07-24** (`SPEC-GDN-SEGMENTS` -> `ACTIVE`): the GDN spec metadata split + decode->prefill reclassification, the `T>1`/`IS_SPEC` recurrence with per-timestep snapshots, the conv sliding window honouring `num_accepted`, k+1 slot allocation - bit-exact rollback. **I5a GDN LAYER ROUTING + runner spec-metadata upload LANDED 2026-07-24** (`CLAIM-SPEC-MTP-I5A`): `GdnBlockPaged` now routes a pure-spec batch through `vt::GdnSpecDecode`/`vt::CausalConv1dSpecUpdate` and the runner uploads I4's six spec device tensors - first of the scoped M-mtp-1 sub-increments (I5a GDN wiring -> I5b prepare_prefill -> I5c MTP paged propose -> I5d config+runner-loop+27B token gate, spec §5), DEFAULT-OFF INERT, bit-exact vs the I4 ops, no e2e loop yet. **I5b `prepare_prefill_inputs` LANDED 2026-07-24** (`CLAIM-SPEC-MTP-I5B`): the drafter prefill input-prep - shift each request's `input_ids` left one within its query span, splice the just-sampled next token, `query_len -= num_rejected`, emit last-token index / query_start_loc / seq_lens into the `SpecPrefillInputs` struct; a HOST routine (no new CUDA kernel; mirrors our DEVICE-NEUTRAL `prepare_inputs`/`combine_sampled_and_draft_tokens` family), unit-gated 7 cases / 27 assertions RED-first, DEFAULT-OFF INERT, additive by construction. **I5c MTP PAGED PROPOSE + DRAFT KV LAYER LANDED 2026-07-24** (`CLAIM-SPEC-MTP-I5C`): `Qwen3_5MTPModel::ForwardPaged` runs the head + one full_attention decoder layer over the head's OWN paged draft KV layer (ReshapeAndCache + PagedAttention over the target's block table / slot mapping); `MakeQwen3_5KVCacheSpec(num_spec>0)` adds that draft KV layer (`fa_draft` FullAttentionSpec group, index num_hidden_layers); `ForwardDeviceTap` exposes the `[T,H]` post-final-norm hidden tap (INERT); and `MtpProposePrefill` is the callable k=1 propose (I5b shift-splice -> one paged forward -> argmax at last_token_indices, early-exit). CORE PROOF: the paged forward reproduces I1's standalone head logits/argmax on BOTH gate checkpoints; a two-step drive proves the draft-KV write/read (RED control diverges). DEFAULT-OFF INERT (no spec config -> draft KV layer not allocated, tap nullptr, target forward byte-identical); NOT wired into the runner step loop. **I5d-pre REGISTRY/FORWARD-SEAM ENABLING REFACTOR LANDED 2026-07-25** (`CLAIM-SPEC-MTP-I5D-PRE`): a scoping pass found the model seam is fully TYPE-ERASED, so the runner cannot reach the concrete target weights / hidden-state tap / loaded MTP weights the I5d loop needs. Four ADDITIVE, inert-when-spec-off access paths + one latent-bug fix - the `hidden_tap` out-field on the type-erased `ModelForwardInput` (routes to the existing `ForwardDeviceTap`), a `LoadedModel::BuildMtpDraft` virtual (typed path to the draft, null for non-MTP), MTP weight loading + shard retention in `FromModelDir` behind `EngineParams::speculative_config`, and the `GPUModelRunner` ctor widened with optional draft/draft-KV/`SpeculativeConfig`; PLUS the latent `initialize_kv_cache` fix (select the FIRST non-eagle full-attn group as the target so a third `fa_draft` group can't displace it; byte-identical at num_spec==0). DEFAULT-OFF INERT, unit-gated RED-first, spec-off SACRED gates byte-identical. **I5d CONFIG RUNTIME + VERIFY/PROPOSE RUNNER LOOP LANDED as a spec-off-byte-identical PARTIAL 2026-07-25** (`CLAIM-SPEC-MTP-I5D`): `--speculative-config` JSON parse -> `EngineParams` -> `LoadedEngine` resolution (widened KV `MakeQwen3_5KVCacheSpec(num_spec>0)`, `BuildMtpDraft`, forced sync scheduling, `MakeScheduler(spec)`, `EngineCore(check_for_draft=true)`) + the full runner loop (draft splice, hidden-tap capture, GDN builder spec-overload feed, k+1 GDN state-slot remap + widened conv cache + draft-KV alloc, `MtpProposePrefill`, `take_draft_token_ids`, acceptance telemetry). CUDA `-Werror` 0 warnings, cutlass-ON; spec-OFF SACRED byte-identical (27B 235/235, 35B 315/315, Coder 138/138 + spec unit tests ALL PASS). **The three-way 27B token gate is NOT yet passing**: the spec-ON engine RUNS the loop end to end and MEASURES the blocker (`test_qwen27_spec_decode`) - it throws on the FIRST prefill step at `gdn_state_gather: working/cache row shapes must match` (`src/vt/ops.cpp:1773`) because I4's spec conv rollback needs the conv row widened to `(K-1)+num_spec` while the non-spec GDN conv ops assume `(K-1)`. `SPEC-MTP` STAYS `GATING`. **I5e LANDED 2026-07-25 (`CLAIM-SPEC-MTP-I5E`): the non-spec GDN conv ops made widened-cache-aware (mirror vLLM `state_len=KERNEL_WIDTH-1` + physical `stride_conv_state_tok`, leading `(K-1)` sub-window, byte-identical at `num_spec==0`) AND the async input-combine forced off under spec (it overwrote the verify batch's draft position with the committed token -> 0 acceptance, RCA'd on the real 27B). **THE THREE-WAY 27B SINGLE-REQUEST GREEDY GATE PASSES**: our-spec-ON == vLLM `--speculative-config mtp` greedy == our-spec-OFF token-for-token, **acceptance 16/16 drafts accepted** (~16 target steps saved); spec-OFF SACRED byte-identical (27B 235/235, 35B 315/315, Coder 138/138), compute-sanitizer 0 on the spec step. `SPEC-MTP` LEAVES `GATING` (single-request greedy correctness PROVEN); NOT `DONE` - the MIXED `GdnBlockPaged` concurrency split/merge + the throughput A/B vs vLLM same-config are I6. **I6 LANDED 2026-07-25 (`CLAIM-SPEC-MTP-I6`, `benchmark_binding=true`): the §5 c1 THROUGHPUT GATE — ours spec-ON AT/ABOVE vLLM spec-ON on EVERY measured axis at c1** (TPOT 66.2/62.95 vs 69.1/65.3 ms prose/code, ours ~1.04x faster; output tput +4.6%/+3.9%; ITL/TTFT lower; acceptance ours 0.85/0.92 vs vLLM 0.838, within noise; spec helps both ~1.5-1.6x TPOT; ours ~4% faster spec-OFF too), via an additive example-only `--speculative-config` bench flag (NO engine code touched). STAYS `ACTIVE`: the c>1 mixed-batch path is still refused + owes a c>1 A/B, and no server-facing spec flag yet. | `ACTIVE` | M-mtp-0, I2 scheduler-half, I3 rejection sampler, I4 GDN spec slots, I5a GDN layer routing, I5b prepare_prefill, I5c MTP paged propose + draft KV, I5d-pre the registry/forward enabling seam, and I5d config runtime + verify/propose runner loop (spec-off byte-identical) are landed; next (before `SPEC-MTP` leaves `GATING`) is closing the measured I5d gate blocker - make the non-spec GDN conv ops widened-cache-aware (mirror vLLM `causal_conv1d` `state_len=width-1+(seqlen-1)`) + the MIXED `GdnBlockPaged` split/merge - then the passing M-mtp-1 27B k=1 greedy three-way token gate + acceptance, then M-mtp-2 35B, then DFlash, the DSpark spike/gates and TLI. **DFlash D0-redo + D1 LANDED 2026-07-26 (`CLAIM-DFLASH-D0D1`, [dflash-spec-decode §0](specs/dflash-spec-decode.md)): `SPEC-DFLASH` UNBLOCKED + `ACTIVE` on the advanced pin `555967922`/vLLM 0.26.0.dev0.** The prior 0.25.0 ORACLE-BLOCKED verdict is SUPERSEDED — under `VLLM_USE_V2_MODEL_RUNNER=1` (vllm#40898 resolved) the mixed-SWA/full z-lab 27B draft CONSTRUCTS and the drafter is ALIVE (acceptance 2.21/8.80/4.75/4.57 > 1, `num_spec=16`, flashinfer-native fp8-KV; goldens committed). Gate FORM measured STRICT MODE-MATCHED (vLLM-ON run-deterministic K>=3 but != vLLM-OFF at k=16 near-ties — NOT the MTP three-way identity). **D1 `DF-AUX-TAPS` DONE:** the single hidden tap is generalized to the multi-tap `[T,H×taps]` (`ForwardDeviceMultiTap` capturing `(hidden+res)` at `target_layer_ids`), config-gated byte-identical off; unit gate 598 assertions (RED-first), CUDA 697/697 + sanitizer 0, 27B MTP e2e 9/9 + 27B SACRED 235/235 byte-identical (inertness). **D2-D5 LANDED 2026-07-26 (`CLAIM-DFLASH-D2`/`D3`/`D4D5`/`D5`):** the drafter model + the project's first non-causal in-block attention (D2, GPU parity vs the real vLLM draft), context-KV precompute + `prepare_dflash_inputs` (D3, GPU numeric-parity 61/61), the non-autoregressive whole-block propose brick + `dflash` config-select (D4), and the RUNNER-LOOP INTEGRATION + 27B e2e (D5): the full verify/propose loop is wired (separate z-lab draft load + target-shared bf16 embed/lm_head, aux-tap capture, per-request combined-feature context accumulation honoring num_rejected, `propose_drafts_dflash`) and RUNS end to end - `test_qwen27_dflash_spec_decode` 2/4 STRICT token-exact vs the vLLM-DFlash-ON golden + acceptance ~ vLLM on ALL 4 (19/39/29/25 vs 17/39/30/25). The 2 divergences are SINGLE bf16 near-tie flips (ratified near-tie ROOT = the D3 inline context-KV recompute envelope), NOT a wiring bug; inertness SACRED 235/235 + MTP 9/9 byte-identical; CUDA `-Werror` clean, no new kernel. NOT a clean strict-4/4 pass - STRICT 4/4 token-identity + the speed A/B = D6 (persistent paged draft-KV bit-matching vLLM's fused projections + the uniform-1+k FULL CG). Capture tool + goldens: `scripts/spec/d{0,2,3}_dflash_*.py`, `tests/parity/goldens/dflash_27b{,_draft,_kvprep}/`. **D6-D9 SPEED CAMPAIGN 2026-07-27:** D6 c1 A/B + STRICT-4/4 bf16-irreducibility RCA; D7 device-resident within-step forward (bit-identical); D8 acceptance RCA + FINAL golden A/B (ours 0.69× vLLM). **D9 (`CLAIM-DFLASH-D9`) PERSISTENT PAGED DRAFT-KV LANDED (bit-identical, +22.7%): `AppendContextKVHost` + `ForwardBlockLogitsWithPrecomputedKV` replace the O(context²) per-step recompute with an append-only per-request store; ours-ON 20.99→25.75 tok/s = 0.917× vLLM-ON (28.09, was 0.69×); e2e 27/27 SAME tokens, SACRED 235/235 + MTP 9/9 byte-identical, CUDA `-Werror` clean, no new kernel. D8's "bf16 acceptance ceiling" REFUTED — same-trajectory per-step acceptance == vLLM (ratio 1.00) and ours realized acceptance (3.68/step) > vLLM (3.31); the SOLE residual (~8%) is the FULL uniform-(1+k) CUDA graph (eager-vs-graphed), a closeable increment. SPEC-DFLASH stays `ACTIVE` (speed not yet ≥ vLLM).** **SPEC-MTP → `DONE` 2026-07-26 (`CLAIM-SPEC-MTP-DONE`, records-only, closing commit I7 `72f9fb1`):** the user ratified the c>1 near-tie+SPEED criterion, closing both I6-owed items (mixed-batch concurrency + server/CLI/C-ABI `--speculative-config`); MTP k=1 is COMPLETE + gated. **M-mtp-2 CLOSED 2026-07-26 (`CLAIM-SPEC-MTP-M-MTP-2`): the 35B `Qwen3_5MoeMTP` full e2e three-way token gate PASSES** — our spec-ON == our spec-OFF == vLLM 0.25.0 `--speculative-config mtp` greedy == vLLM spec-OFF, 16/16 vs the `greedy_ids` anchor (STRICT, c1), acceptance 16/16 both sides; c1 spec-ON 1.19x TPOT / +16.3% output-tput vs spec-OFF (acceptance 0.908) — MoE speedup transfers; spec-OFF byte-identical (test+docs-only). MTP is now `DONE` on BOTH gate models (`MODEL-SPEC-qwen3-5-mtp-qwen3-5-moe-mtp` `GATING`→`DONE`). **DFlash D11+D12 2026-07-27 — the FULL uniform-(1+k) CUDA graph is being built in three parts:** D11 (`CLAIM-DFLASH-D11`) landed Part A (the device-store primitive, CPU-gated); **D12 (`CLAIM-DFLASH-D12`) landed A-wire (the D11 device store is now the PRODUCTION path; GPU-gated e2e 27/27 all-exact acceptance 19/39/29/25 + SACRED 235/235 + MTP 9/9 byte-identical) + Part B (`vt::DFlashPagedBlockAttention`, the capture-safe paged kernel; `test_ops_dflash_paged_block_attn` 795648/795648 CPU==CUDA + cross-check vs materialized `DFlashBlockAttention` + compute-sanitizer 0; NO function-local host `cu_seqlens` upload = capture-UAF fixed).** Speed UNCHANGED 0.917× (A-wire eager + Part B not yet wired). The SOLE remaining piece is Part C (static-shape capture + device mask-scatter + `BeginCapture`/replay + the ≥vLLM c1 A/B); if ours-ON-graphed ≥ vLLM-ON → SPEC-DFLASH DONE → C3 complete. C3 stays `ACTIVE` (DFlash Part C + DSpark/TLI remain) | | 4 | `ROAD-V1-C4` | Quantization: llama.cpp breadth/speed, NVFP4/FP8/MX, MLX native | [quantization matrix](quantization-matrix.md) | coverage spike merged; `QUANT-GGUF-CPU-THREADPOOL` W1-W3 implemented and correctness-gated, still `GATING` (its reproduction now exists — same-binary 1-vs-20-thread A/B is prefill 12.47x / decode 8.05x / RSS 1.000x, so **decode misses the >=10x bar**). **GGUF COMPUTE-IN-QUANT IS NOW LIVE AND DEFAULT-ON (2026-07-22, `CLAIM-QUANT-GGUF-CIQ-G4-1`):** [compute-in-quant GEMM](specs/gguf-compute-in-quant-gemm.md) **G1-G4** — block dtypes + traits, the Q8_0/Q8_K activation quantizers, the six generic `vec_dot`, `kMatmulBTQuant`, and now the ROUTING (`vt::MatmulBT` dispatches a block-dtype weight to it) — plus [keep-quant loader](specs/gguf-keep-quant-loader.md) **L1-L4**, whose master switch defaults ON wherever that op is registered for the running device (CPU today; a CUDA build still expands). Six encodings (Q4_0, Q8_0, Q3_K, Q4_K, Q5_K, Q6_K) now carry `C` = `Y`. **Correctness held exactly**: the 35B GGUF gate is 16/16 token-exact vs the same-file llama.cpp oracle with the quant path on, and the bench model's output tokens are byte-identical across the pre-G4, post-G4 and `VT_CPU_REF=1` arms — no golden regenerated. **Binding CPU A/B** (idle dgx aarch64, one flock, same binary, 3 reps): decode **3.45x**, prefill **4.16x**, peak RSS **1.16x less**; vs llama.cpp we went from 11.7x / 34.1x / 2.66x behind to **3.38x / 8.20x / 2.29x**. The projected 9-17x did **not** hold, for a measured reason: 60 % of that file's weight bytes are `f16`, which no block encoding covers. Keep-quant loader **L4** is therefore MEASURED-and-NOT-MET on RSS; other leaf specs open. **THAT #1 LEVER IS NOW LANDED (2026-07-22, `CLAIM-KERNEL-CPU-ELEM-GEMM-1`, new row [`KERNEL-GEMM-CPU-ELEM`](kernel-matrix.md)):** [the elementwise CPU GEMM](specs/cpu-elementwise-gemm.md) **E1-E4** — per-dtype specialization out of the K loop, 16 independent accumulators instead of one, AArch64 NEON + x86-64 SSE2/F16C tiers behind a runtime probe, and M-blocking — all **BYTE-IDENTICAL** to the historical kernel (`memcmp` gate, exhaustive 65,536-pattern widening check, same token md5), so nothing was regenerated. Binding same-binary A/B: prefill **3.41x**, decode **3.11x**; op-level bf16 18-24 -> 69-351 GFLOP/s. **vs llama.cpp: decode AT PARITY (1.03x), prefill 2.34x behind, RSS 2.29x worse. THEN loader L5 LANDED (2026-07-23, `CLAIM-QUANT-GGUF-KEEPQ-L5-1`):** [keep-quant loader](specs/gguf-keep-quant-loader.md) **L5** — mmap in-place residency (borrow kept blocks out of the read-only mapping, refcounted), tied-head sharing (one bf16 vocab matrix for embed+lm_head), and a read-once page release (port of llama.cpp `unmap_fragment`) — took **peak RSS 6.401 -> 3.884 GiB (2.29x -> 1.39x llama.cpp)** with decode UNCHANGED and output tokens byte-identical (md5 `d235db12f2cd304007530286a1755c95`). The remaining ~1.09 GiB over llama.cpp is the f16 expansion (no block encoding covers f16). | `PARTIAL` | **THE OWED FRESH PROFILE IS DONE (2026-07-23) and it re-ranks the plan.** A `vt::GetOp` hook (100% of wall time, reverted before binding) on the CURRENT binary: prefill is **no longer GEMM-bound** — kMatmulBTQuant 37%, **kGdnPrefill 25%**, kMatmul 12%, kMatmulBT 10%, **kPagedAttention 10%**; the two non-GEMM kernels (GDN linear-attention recurrence + paged attention) run **SINGLE-THREADED** on the CPU and are now the top prefill levers. Decode is memory-bound matmul at parity, no kernel work owed. **RE-RANK: G5/G6/G7 all only speed the already-fast quant GEMM and rank BELOW the two serial non-GEMM kernels; the new #1 CPU lever is threading kGdnPrefill + kPagedAttention.** **THAT #1 LEVER IS NOW LANDED (2026-07-23, `CLAIM-CPU-THREAD-GDN-PAGED-1`, [two-kernel threading](specs/cpu-thread-gdn-paged-2026-07-23.md)):** kGdnPrefill chunks over the (sequence, value-head) axis and kPagedAttention over query-token rows, both via the existing `ParallelForRows`, both **byte-identical** (qwen35 output-token md5 `d235db12f2cd304007530286a1755c95` unchanged at threads 1/4/20 + `VT_CPU_REF=1`, determinism battery extended, CPU ctest 158/158). **Binding dgx aarch64 (idle): prefill 1.382× same-binary (73.0→100.9 t/s), 2.43×→1.76× behind llama.cpp pp128; decode at parity; op-scaling 1→20 GdnPrefill 7.08× / PagedAttention 8.96×; fresh profile shows the two kernels 35%→8.6% of prefill and re-ranks the NEW bottleneck to the GEMMs (kMatmulBTQuant 50% + kMatmul 16% + kMatmulBT 14% = 80%) ⇒ next CPU lever is the SIMD/repack GEMM tiers (G5/G6/G7).** **THE FIRST SUCH TIER IS NOW LANDED (2026-07-23, `CLAIM-QUANT-GGUF-CIQ-G6-1`, [compute-in-quant GEMM](specs/gguf-compute-in-quant-gemm.md) G6):** the Arm **i8mm mmla `nrc==2`** `vec_dot` 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 (decode M=1 → portable, unchanged); runtime `HWCAP2_I8MM` probe + `VT_CPU_QUANT_MMLA` defeat + per-file `+i8mm`. **BYTE-IDENTICAL** where the math allows (q8_0/q4_0 bit-exact to the scalar tier, q4_K/q6_K within NMSE ≤ 5e-4), bit-identical across threads 1/2/4/20, e2e token md5 `d235db12f2cd304007530286a1755c95` byte-identical (mmla on/off/`VT_CPU_REF=1`), 35B GGUF gate 16/16 vs llama.cpp on both files. **Op-level portable→i8mm: q4_K 7–8.4×, q6_K 3.8–4.5×, q8_0 ~1.2×**; e2e prefill same-binary **1.084×** on the q8_0-dominant bench file (1.56×→**1.44× behind** llama.cpp pp128, Amdahl-bounded — the big k-quant win lands on the APEX 35B files). Fresh bottleneck: the elementwise f16/f32 GEMM (~30%, unchanged) is now co-dominant on this mixed file. CUDA `-Werror` 0-warn, regression set UNCHANGED. `docs/BENCHMARKS.md` ACCEPTED. RSS deficit closed to 1.39x by L5; the last RSS lever is an f16 keep-as-is compute path, not this loader. **THEN the GDN split-projection orientation LANDED (2026-07-23, `CLAIM-CPU-GDN-ORIENT-1`, [GDN projection orientation](specs/cpu-gdn-proj-orientation-2026-07-23.md)):** a fresh op-dispatch profile of the current binary (warm prefill, `vt::GetOp` hook + per-GEMM shape histogram, reverted before binding) found the four GDN input projections (`in_proj_qkv/z/b/a`, 72 GEMMs, **17.9%** of prefill: `kMatmulBTQuant 50.7% / kMatmul 17.9% / kMatmulBT 14.9%`) were the LAST weight family `LoadGdnGguf` still transposed into [K,N] (nk=false → the N-striding `kMatmul`, no M-blocking) after G4's `expand_nk` gave every other expanded weight the file's own [N,K] order. New `GgufLoadPolicy::gdn_expand_nk` + `MakeGdnProj` keep them [N,K] nk=true → the M-blocked `kMatmulBT`; **BYTE-IDENTICAL** (same sequential f32 K-reduction, only the weight offset differs — token md5 `d235db12f2cd304007530286a1755c95` unchanged across on/`VT_GGUF_GDN_NK=0`/`VT_CPU_REF=1` and threads 1/4/20), `test_qwen36_gguf_engine` 2/2·28/28·16/16 on APEX. **Binding same-binary prefill 1.090× / decode 1.09× (44.1→40.4 ms TPOT = 1.01× llama tg32, at parity), 1.44×→1.32× behind llama.cpp pp128, RSS unchanged.** Fresh post-change profile: `kMatmul` **17.9%→0% (72→0 calls, ELIMINATED)**, absorbed into `kMatmulBT` (14.9%→27.7%); **next CPU prefill lever = the quant GEMM (kMatmulBTQuant, now 55%): G7 repack-at-load.** **G7 LANDED 2026-07-23 (`CLAIM-QUANT-GGUF-CIQ-G7-1`, [compute-in-quant GEMM](specs/gguf-compute-in-quant-gemm.md) G7):** q8_0 repacked once at load into the i8mm `block_q8_0x4` interleave (ported llama.cpp `repack.cpp` `q8_0_4x8`), `kMatmulBTQuant` dispatches a pre-shuffled gemm/gemv with no per-block register shuffles. **BIT-IDENTICAL** (byte-permute weight + non-fused `vmlaq_f32`, 305-assertion memcmp across decode/prefill/bf16-out/threads, token md5 `d235db12f2cd304007530286a1755c95` unchanged on/`VT_CPU_QUANT_REPACK=0`/`VT_CPU_REF=1`; a `ResidentWeight`/`MakeTensor` flag-drop that produced all-zero tokens was caught by the E2E gate and fixed). Op-level q8_0 **3.7–5.9×** (518→2401 / 583→3456 / 514→1902 GFLOP/s); **E2E prefill 1.92× same-binary (1096→572 ms), 223.8 t/s vs llama.cpp pp128 177.3 = 1.26× — AT/BEYOND PARITY** (was ~1.5× behind), decode at parity, RSS 3.884 GiB unchanged. Fresh profile: q8_0 GEMM 55%→~21%; **the CPU prefill-lever search is CLOSED — the sole remaining gap to llama.cpp is peak RSS (1.39×), not prefill.** CUDA-inert (gated off any non-CPU-quant device), CUDA `-Werror` 0-warn, regression set UNCHANGED. **THE RSS GAP IS NOW CORRECTLY ATTRIBUTED (2026-07-23, `CLAIM-QUANT-GGUF-KEEPF16-L6-1`, [keep-quant loader](specs/gguf-keep-quant-loader.md) L6): it is NOT the f16 expansion.** L6 implemented keep-f16 residency (keep the file's 56 F16 weights + tied head resident as F16 and compute on them, mirroring llama.cpp `ggml_vec_dot_f16`) and MEASURED it **RSS-NEUTRAL** (3.884 → 3.832 GiB, −52 MB) and prefill-regressive (TTFT 577 → ~1000 ms, from 1.25× ahead of llama.cpp to 0.72× behind) — because L5's page-release had ALREADY dropped the f16 file pages, so keep-f16 only swaps an anonymous bf16 buffer for equal-size file-backed f16 pages. smaps attribution proves our weight residency is at llama.cpp parity (file-backed 2.63 ≈ 2.68 GiB); **the residual ~1.08 GiB is the engine's ANONYMOUS activation/KV workspace, not weights — the real, separate CPU RSS lever.** keep-f16 ships DEFAULT OFF (`VT_GGUF_KEEP_F16=1` opt-in), tokens byte-identical (md5 `d235db1…`), `test_gguf_keep_quant` 35/35 (x86+aarch64), regressions UNCHANGED (27B 235/235, 35B 315/315, Coder 138, dense 184, OPT 63, DeepSeek 223, Llama 92, GGUF engine 28/28). **NEXT CPU RSS lever: profile + shrink the engine's activation/KV working set, NOT the weight loader** | | 5 | `ROAD-V1-C5` | Sliding window, local attention, YaRN/long context | [engine matrix](engine-matrix.md), [coverage view §§2,11](feature-matrix.md#2-kv-cache--memory), [joint spike](specs/sliding-local-yarn-long-context.md) | **CUDA GPU CLOSURE 2026-07-27 (`CLAIM-ROADMAP-C5`, dgx GB10 sm_121a, clean build of `489f7771`, oracle vLLM 0.26.0.dev0):** the shared scaled-RoPE + local-mask CUDA path COMPILES `-Werror`-clean and RUNS on GB10; the C5 feature-positive correctness gates that were the stated `GATING` blocker now PASS — SWA (Gemma-2/Gemma-3 48/48), LongRoPE (Phi-4-mini 16/16, RED-first), llama3-rope (Llama-3.2-1B 16/16), dynamic-NTK (InternLM2 16/16); both RoPE 0.26-oracle recaptures BIT-IDENTICAL to goldens (zero drift). Leaves `ATTN-SLIDING-WINDOW`/`ATTN-ROPE-{LLAMA3,LONGROPE,DYNAMIC-NTK}`/`ATTN-YARN` → `ACTIVE` | `PARTIAL` | (RI) **Honest residual (vehicle-blocked, not skipped):** YaRN model e2e (no cached Nomic/gpt-oss consumer) + chunked-local model e2e (no Llama4 row) are REACHABLE-BLOCKED — operator/formula stay GPU/G3-gated; long-context positive-mask (prompt > W) SWA model e2e + the KV memory-optimization G8; and the roadmap-wide every-axis SPEED tail (all C5 leaves correctness-complete, speed-pending, mirroring their model consumers). Not row-DONE until speed + the blocked vehicles close | diff --git a/.agents/specs/sm120-qwen35-pareto-2026-08-09.md b/.agents/specs/sm120-qwen35-pareto-2026-08-09.md index fff40af6f..2807967f8 100644 --- a/.agents/specs/sm120-qwen35-pareto-2026-08-09.md +++ b/.agents/specs/sm120-qwen35-pareto-2026-08-09.md @@ -3,9 +3,11 @@ **Issue:** [#206 — RTX 5070 Ti: close Qwen3.5-4B TTFT, TPOT, and VRAM gaps vs vLLM](https://github.com/mudler/vllm.cpp/issues/206) -**Rows:** `KERNEL-SSM-MAMBA`, feeding `ROAD-V1-C2-LOCAL-BF16` +**Rows:** `KERNEL-SSM-MAMBA`, `SERVE-CLI-BENCH`, feeding +`ROAD-V1-C2-LOCAL-BF16` -**Lifecycle:** SPEC COMMITTED; code transplant and fresh measurement pending +**Lifecycle:** combined prefill opt-in accepted locally; frontend timing parity +correction specified before the next cross-engine claim **Target:** local NVIDIA GeForce RTX 5070 Ti (`sm_120`), Qwen3.5-4B plain BF16. This vehicle does not substitute for the hardware-unavailable 27B/35B @@ -21,8 +23,8 @@ prompts and production configuration; vLLM is not run with `--enforce-eager`. | Axis | vllm.cpp | pinned vLLM | Status | |---|---:|---:|---| -| total throughput | **6784.743 tok/s** | 6643.593 tok/s | **PASS**, 1.021246x | -| mean TTFT | 1018.040 ms | **937.584 ms** | **OPEN**, 1.085812x | +| total throughput | 6784.743 tok/s | 6643.593 tok/s | **VOID**: frontend mismatch; corrected run pending | +| mean TTFT | 1018.040 ms | 937.584 ms | **VOID**: frontend mismatch; corrected run pending | | mean TPOT / ITL | 34.740 ms | **33.906 ms** | **OPEN**, 1.024597x | | peak VRAM | 13053.3 MiB | **12820 MiB** | **OPEN**, 1.018199x | @@ -34,6 +36,12 @@ whose improvement preserves every passing axis. Binding baseline and reproduction detail: [Qwen3.5-4B sm_120 evidence](../../docs/bench-evidence/qwen35-4b-sm120-main-20260807.md). +The values remain historical observations, but the throughput and TTFT ratios +are not gateable: local includes string tokenization after its timer starts, +whereas pinned vLLM pre-tokenizes before its timer. The same-binary local +kernel A/B remains valid; TPOT remains comparable because it starts after the +first token. + ## Source ground and existing evidence Current local exact-chunk dispatch and causal-conv kernel are grounded at @@ -87,6 +95,42 @@ what was tried, including negative results. In particular: 6. Attempt BF16 vector writeback only after the TTFT/Pareto work above, and only under its own compiler, correctness and rejection gates. +## Pretokenized benchmark-front-end parity contract + +The TTFT attribution selected a frontend correction before any further kernel +work. Local starts `t0` before admission and currently calls the string +`AsyncLLM::add_request` overload (`examples/bench/bench_core.h:508-529`), while +the oracle encodes all prompt strings before `run_closed_loop` starts and +submits `TokensPrompt` IDs +(`tools/bench/vllm_closed_loop_metrics.py:59-82,137-167`). The implementation +under issue #206 must satisfy this contract: + +1. After the tokenizer/model is available, encode **all** selected prompts + before `t0`. Preserve the prompt order and exact token IDs, then admit via + the existing pretokenized `AsyncLLM::add_request` overload + (`include/vllm/v1/engine/async_llm.h:80-89`). No tokenization may occur in + the default timed admission interval. +2. Add a pure inline production dispatch seam, e.g. + `DispatchBenchPromptAdmission(env, pretokenized_cb, timed_string_cb)`. An + unset `VT_BENCH_PRETOKENIZE` and `=1` invoke **only** the token-ID callback; + exact rollback `VT_BENCH_PRETOKENIZE=0` invokes **only** the existing timed + string callback. Do not add a hot-path counter or debug-state branch. +3. Focused CPU tests own parser/default/rollback behavior and callback identity: + default/unset and `=1` select only the pretokenized callback, `=0` selects + only the string callback. Mutating the parser, inverting callback selection, + or invoking both callbacks must fail. An end-to-end synthetic default/rollback + A/B must prove prompt-token IDs and output-token IDs identical in submission + order. +4. Measure the production vehicle in one immutable binary using counterbalanced + `STRINGa -> TOKENSa -> TOKENSb -> STRINGb`, with exact prompt/output token + hashes. The default may become ON only after both token-ID legs win TTFT and + total throughput, and neither leg regresses TPOT/ITL, E2E, correctness, peak + VRAM, host PSS, or load behavior outside the calibrated noise band. +5. After acceptance, run a fresh local-token-ID versus pinned-vLLM comparison + on identical pretokenized prompts and the existing production configurations. + Only that rerun may restore cross-engine throughput or TTFT ratios. Retain + `VT_BENCH_PRETOKENIZE=0` as the exact same-binary rollback. + ## Correctness, tests, review and acceptance Before product code, a fresh implementer ports or restores the smallest tests diff --git a/docs/BENCHMARKS.md b/docs/BENCHMARKS.md index 7ca38eb92..aed942b9c 100644 --- a/docs/BENCHMARKS.md +++ b/docs/BENCHMARKS.md @@ -41,7 +41,7 @@ The binding comparison. vLLM runs its **production graphed config**, never | Qwen3.6-27B | NVFP4 (`nvidia` @`0893e160`, ModelOpt `modelopt_mixed`) | 0.25.0 | 0/4 | **BEHIND, uniformly 0.85x** on decode throughput (was 0.72x before the FP8 tower fix); greedy continuation IDENTICAL to vLLM. A different model from the `unsloth` row (NVFP4 MLP + FP8 W8A8 GDN/attn tower) | | Qwen3.6-35B-A3B | NVFP4 `modelopt_mixed` | 0.25.0 | 2/18 | 3-rep grid 2026-08-05 @`1ea26427`: 0.93-1.03x (c4 wins), c16 0.93x. Both c16 levers A/B'd NEG: drain event -1.9%, mirror 0.999x. ★ probe found a prod async batch-1 greedy DEGENERATION bug the mirror fixes | | DeepSeek-V2-Lite | bf16 MLA | 0.25.0 | 4/25 | Attributed miss, row stays `ACTIVE` | -| Qwen3.5-4B | bf16 direct-load | 0.26.0.dev0 | throughput + host PSS | Exact chunks: throughput **1.021x PASS**; TTFT **1.086x**, TPOT **1.025x**, VRAM **1.018x OPEN**. #206 opt-ins restored; combined A/B pending ([data](bench-evidence/qwen35-4b-sm120-main-20260807.md)) | +| Qwen3.5-4B | bf16 direct-load | 0.26.0.dev0 | host PSS; frontend correction pending | Combined local A/B accepted: +0.718% tput, -1.340% TTFT, -0.556% TPOT; IDs exact. Cross-engine tput/TTFT **VOID/PENDING**; TPOT 1.025x and VRAM 1.018x OPEN ([data](bench-evidence/qwen35-4b-sm120-main-20260807.md)) | ### GDN prefill causal-convolution by GPU diff --git a/docs/STATUS.md b/docs/STATUS.md index 479774aac..36c068e7d 100644 --- a/docs/STATUS.md +++ b/docs/STATUS.md @@ -78,7 +78,7 @@ token-for-token correctness against the pinned oracle. | Qwen3.6-27B (NVFP4) text generation | Correctness-complete; speed is CHECKPOINT-dependent | Token-exact GB10 on both. `unsloth` @`890bdef7` beats vLLM 0.25.0 every c (1.007-1.045x), 115/124; `nvidia` @`0893e160` (ModelOpt FP8 tower) is **0.85x BEHIND**, decode ~100% GPU-busy | | Qwen3.6-35B-A3B (NVFP4, GDN MoE) | Correctness-complete; binding grid @`a0fa12c7` FLAT 0.935x-0.979x over c1-c32 (CoV <0.81%); the prior 0.87x c2 / 0.92x c8 "weak cells" were harness mismatch, not code; memory PSS 3.81x, GPU 1.40x | Token-exact SYNC+ASYNC; `VT_ASYNC_DEVICE_MIRROR` ON fixes async batch-1 token-0 degeneration; `VT_ASYNC_EXECUTOR` Option A NEUTRAL → OFF | | Qwen3 / Qwen2 dense (BF16) | Correctness-complete, speed-pending. Async-serving P0 FIXED (`ROW-SERVE-ASYNC-DENSE-MIRROR`): classic-dense `Qwen3ForCausalLM` now honors the async device token-ids mirror; CPU-only -Werror test-guard fixes x2 | Near-tie-robust token-exact vs vLLM (Qwen3-0.6B, Qwen3-4B); c1 effective parity, c8 decode residual. **Async device-mirror (`ROW-SERVE-ASYNC-DENSE-MIRROR`, `f9c969ae`): the #31 fix ported to the classic dense family, dgx-VERIFIED.** The shared dense `EmbedInto` (qwen3.cpp) raced the async combine's device input-ids write against a stale host upload → token-0 degeneration on the depth-2 AsyncLLM serving path (quant-independent). `EmbedInto` now consumes the device override published by `ForwardQwen3ForCausalLM`'s `DeviceTokenIdsScope` (27B-dense template); gate `test_qwen3_dense_async_serving` RED on `VT_ASYNC_DEVICE_MIRROR=0`, GREEN default, byte-identical mirror-off. dgx GB10: async gate RED→GREEN 0.6B+4B, SACRED 0.6B+4B 184/184 unchanged (byte-neutral sync path), memcheck 0 errors; Yi30/Qwen3-8B-MXFP4 default-config e2e coherent + 3/4 token-exact (p2 = oracle-ratified near-tie, gap 0.0000), closing the QUANT-CT-MXFP4 async-default residual. RESIDUAL: sibling InternLM2/Mistral/Llama scope one-liner; W4 bench RAN; FA2 GQA-swap default-ON, c2-c8 <1.0x. `FLASH-PTXAS` #82: codegen at PARITY (no ptxas lever); gap=engine context. **D1 (2026-07-31, `CLAIM-D1-BF16-MERGED-QKV`): the bf16 merged-QKV path (`Qwen3QkvMergeEnabled`/`VT_QWEN3_QKV_MERGE`) is now default-ON** — one `vt::MatmulBT` over the merged `[qdim+2kdim,H]` owner + a contiguous `vt::QkvSplit` (OLMo-2 exemplar), replacing three per-shard GEMMs. Bit-exact GEMM math (A/B unit `test_ops_qkv_merge` byte-identical, RED-first); the wider-N cuBLASLt K-reduction flips the 0.6B genuine bf16 near-tie so the SACRED 0.6B golden was regenerated (all tokens within the near-tie band, max 0.125 nats), while Qwen3-4B is byte-neutral (0 diffs, stays STRICT). Re-gated 0.6B 16/16 + 4B 16/16; consistency/launch-count fold (measured NEUTRAL on 4B decode), no new throughput owed | -| Qwen3.5-4B plain BF16 direct loading on discrete CUDA | Correctness-complete; throughput passes, latency/VRAM open | Exact chunks are byte-identical. vLLM: throughput 1.021x PASS; TTFT 1.086x, TPOT 1.025x, VRAM +233 MiB OPEN. #206 opt-ins restored; combined A/B pending ([data](bench-evidence/qwen35-4b-sm120-main-20260807.md)) | +| Qwen3.5-4B plain BF16 direct loading on discrete CUDA | Correctness-complete; frontend correction pending; TPOT/VRAM open | Combined local A/B: +0.718% tput, -1.340% TTFT, -0.556% TPOT; IDs exact. Cross-engine tput/TTFT VOID pending pretokenized timing; TPOT 1.025x, VRAM +233 MiB OPEN ([data](bench-evidence/qwen35-4b-sm120-main-20260807.md)) | | Qwen3-Coder-30B-A3B MoE (BF16) | Correctness-complete, speed-pending | Near-tie-robust token-exact 6/6; 11 of 16 binding grid cells at or above vLLM. **D1 (2026-07-31): inherits the default-ON bf16 merged-QKV via the shared dense `AttnBlock` — byte-neutral (0 token diffs, golden UNCHANGED); re-gated 6/6** | | Llama-3.x dense (BF16) | Correctness-complete, speed-pending | Near-tie-robust token-exact 16/16 (Llama-3.2-1B); llama3 RoPE scaling | | Mistral dense (BF16) | Correctness-complete, speed-pending | Paged-engine token-exact 16/16 (Mistral-7B-v0.3) | From 67767c268f2c04482b14b7dd8957bff7097f4b7a Mon Sep 17 00:00:00 2001 From: Richard Palethorpe Date: Sun, 9 Aug 2026 19:58:19 +0100 Subject: [PATCH 05/17] perf(bench): pretokenize prompts before TTFT timing (#206) Match the pinned vLLM benchmark frontend by encoding all prompts before the timed interval and admitting token IDs by default. Keep VT_BENCH_PRETOKENIZE=0 as the exact timed-string rollback and report the resolved path. Focused CPU gate: test_bench 6/6 tests, 57/57 assertions. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: Codex:gpt-5.6-sol [Codex] --- .agents/NOW.md | 4 +- .agents/engine-matrix.md | 2 +- .agents/roadmap_v1.md | 2 +- .../specs/sm120-qwen35-pareto-2026-08-09.md | 13 +- docs/BENCHMARKS.md | 2 +- docs/ENVIRONMENT.md | 1 + docs/STATUS.md | 2 +- examples/bench/bench_core.h | 77 ++++++++++-- tests/examples/test_bench.cpp | 113 ++++++++++++++++++ 9 files changed, 202 insertions(+), 14 deletions(-) diff --git a/.agents/NOW.md b/.agents/NOW.md index 78463d0ba..8e0ebb162 100644 --- a/.agents/NOW.md +++ b/.agents/NOW.md @@ -21,7 +21,7 @@ Work: exact-chunks on main `1ce0d662b`; sm_120 measured at `3d2581551`. | MiniMax-H3 | **PRUNED ckpts RUN (#241): Q8_0 renders, seam 0.9941** | same-binary A/B | | Kimi-Linear-48B | 122/128 held; grouped router parallelised, e2e NOT ESTABLISHED | ckpt is tiktoken-only: no warm server | | 35B binding grid | @`a0fa12c7`: **flat 0.935-0.979x, NO c2/c8 weak cell** (CoV <0.81%); mem PSS 3.81x | Attribute the flat ~5% mid-band; TTFT c2 0.872x | -| Qwen3.5-4B sm_120 | #206 K4+postconv Pareto win; IDs exact. Cross-engine tput/TTFT VOID; TPOT/VRAM OPEN | Pretokenize; rerun | +| Qwen3.5-4B sm_120 | #206 pretoken CPU green; IDs exact vs rollback. Cross-engine tput/TTFT VOID; TPOT/VRAM OPEN | Review; GPU A/B | | RPi5 A76 CPU | **R5 asm GREEN; llama NOT MET**: 0.461x pf, 0.653x dec | W6: BF16 GEMM | | MXFP4 parity | c1 1.020, c2-c8 0.962-0.969. **#82 CLOSED: ptxas-lineage REFUTED** | TERMINAL: at parity | | SERVE-ASYNC-DENSE-MIRROR | **LANDED+VERIFIED** (`f9c969ae`): async mirror, dense Qwen3; SACRED 184/184 | Sibling scope one-liner | @@ -53,7 +53,7 @@ latency/memory on every axis, both gate models, reproduced 2–3x idle. See of roof. Dense-marlin +0.5%; Triton-AOT GDN a WASH. 2. **Spike the Parakeet encoder row** (vLLM: `nano_nemotron_vl.py`; the transducer half is NOT in vLLM: separate call). -3. **Qwen3.5-4B #206:** pretokenize benchmark; rerun cross-engine. +3. **Qwen3.5-4B #206:** pretoken green; review/operator; GPU/vLLM rerun. 2. **Merge the invocation-parity prevention** (CI guard + AGENTS.md checklist); CUDA build-verify the byte-exact `kGemvHeuristicAlgos` refactor on dgx. 4. **Restore `local-ai-worker`** on dgx at campaign end (`--restart=always`). diff --git a/.agents/engine-matrix.md b/.agents/engine-matrix.md index 159698900..4a2cdc939 100644 --- a/.agents/engine-matrix.md +++ b/.agents/engine-matrix.md @@ -204,7 +204,7 @@ claims it. | `SERVE-HTTP-TRANSPORT` | Serving-socket transport parity: mirror vLLM's uvicorn/asyncio default `TCP_NODELAY` on every accepted SSE socket so per-token stream frames are not held by Nagle against the peer's delayed ACK. Implemented + CPU-tested; the non-binding localhost A/B sizing is COMPLETE and NEUTRAL within noise on c1/c2 ITL/TPOT/throughput (loopback ACKs are instant, so Nagle never coalesces ~100 ms-cadence token frames) — no gate-axis credit expected; the mirror stays for real-network parity. Future keep-alive / read-write-timeout / listening-socket option parity noted, not done | T0 | vLLM serves via uvicorn over asyncio `vllm/entrypoints/launcher.py:71,76`, `vllm/entrypoints/openai/api_server.py:591,630`; asyncio disables Nagle per accepted TCP stream socket `asyncio/base_events.py:192-197` (`_set_nodelay`) called from `asyncio/selector_events.py:950`; cpp-httplib default-off `third_party/httplib/httplib.h:142`, applied on accept only when set `third_party/httplib/httplib.h:12083` | `src/vllm/entrypoints/openai/api_server.cpp:69` (`set_tcp_nodelay(true)` in the ApiServer setup) | behavioral accepted-socket `getsockopt(TCP_NODELAY)` case `tests/vllm/entrypoints/openai/test_api_server.cpp:1076` (helper `:380`); RED accepted `TCP_NODELAY` 0 → GREEN 1, full `test_openai_api_server` **22/22 cases / 242 assertions**; non-binding sizing root `~/work/vllm.cpp-tcpnodelay-sizing/ff915e8…` (raw-set SHA `f5b52900…2128`) neutral within noise; closure [ledger](parity-ledger.md#L451) | [serve-tcp-nodelay.md](specs/serve-tcp-nodelay.md) | `DONE` | `ff915e8` | | `SERVE-C-ABI` | Stable LocalAI-style C FFI (**19** exported `VLLM_API` symbols at `VLLM_ABI_VERSION 10`; blocking and nonblocking request handles. Count corrected 2026-07-24 from a stale `17`, which predated ABI v4/v5 adding `tool_parser`/`reasoning_parser` and the chat entry points; `include/vllm.h` is the source of truth and README:231 already said 19). **ABI v9 2026-07-28 (`CLAIM-CAPI-ENGINE-CONFIG-V9`): the ABI carried strictly LESS engine config than `EngineParams` does** - `max_num_batched_tokens`, the scheduler `scheduling_policy` (`fcfs` / `priority` / `lpm`), and `kv_transfer_config` (the external KV connector / LMCache JSON) were reachable from the bundled server's flags and from NO embedder. All three added, inert at their defaults (zero-filled v8 growth == byte-identical pre-v9 engine); the connector NAME is validated against `KVConnectorFactory` at load, mirroring the server's startup check. `tokenizer_config_path` stopped being a declared-since-v1 no-op and now selects the chat template's source file. Malformed `speculative_config`/`kv_transfer_config` documents now report `VLLM_ERR_INVALID_ARGUMENT` (the contract vllm.h documented since v6) instead of `VLLM_ERR_MODEL_LOAD`, via a catch scoped to the parse block so a real `FromModelDir` failure still reports MODEL_LOAD. Driver: the LocalAI vllm-cpp backend could not expose LMCache or the prefill budget in a model config) | T0 | Original project ABI; pinned vLLM has no C ABI | `include/vllm.h:143,181,207`; `src/capi/vllm_c.cpp:229,264,327,391` | `tests/capi/test_capi.cpp:320,428,505,574,606,640`; `tests/capi/test_dlopen.cpp:77,86`; `tests/capi/c_header_compile.c:1` | [c-api-library.md](specs/c-api-library.md) | `ANCHOR-BACKFILL` | `CLAIM-SERVE-C-ABI-SPIKE` | | `SERVE-CPP-API` | Rich `LLM` and `AsyncLLM` C++ API | T1 | `vllm/entrypoints/llm.py:66,422`; `vllm/v1/engine/async_llm.py:70` | - | - | `planned: specs/cpp-api.md` | `INVENTORIED` | - | -| `SERVE-CLI-BENCH` | Serve and latency/throughput/serve benchmark modes | T0 | `vllm/entrypoints/cli/serve.py:44`; `vllm/entrypoints/cli/benchmark/main.py:29`; production queue `vllm/v1/engine/core.py:200-231,622-669`; pinned comparison pretokenizes before timing `tools/bench/vllm_closed_loop_metrics.py:137-167` | separate binaries + explicit scheduler-capacity flags `examples/server/main.cpp:63,96,116,170`; production `AsyncLLM` benchmark frontend + auditable scheduler depth `examples/bench/bench_core.h:426,495,595`; current string admission starts after `t0` at `examples/bench/bench_core.h:508-529`, so #206 cross-engine throughput/TTFT is VOID pending the committed pretokenized-admission correction; `examples/bench/main.cpp:51` | existing server-help/production-frontend/metric assertions `examples/CMakeLists.txt:34`, `tests/examples/test_bench.cpp:18,29-32,61,81,97`; owed #206 default/rollback callback-identity and exact prompt/output-token A/B tests | [CLI/serve/benchmark spike](specs/cli-serve-bench.md); [#206 campaign contract](specs/sm120-qwen35-pareto-2026-08-09.md) | `PARTIAL` | - | +| `SERVE-CLI-BENCH` | Serve and latency/throughput/serve benchmark modes | T0 | `vllm/entrypoints/cli/serve.py:44`; `vllm/entrypoints/cli/benchmark/main.py:29`; production queue `vllm/v1/engine/core.py:200-231,622-669`; pinned comparison pretokenizes before timing `tools/bench/vllm_closed_loop_metrics.py:137-167` | separate binaries + explicit scheduler-capacity flags `examples/server/main.cpp:63,96,116,170`; production `AsyncLLM` benchmark frontend + auditable scheduler depth; #206 default pre-encodes every prompt before `t0` and admits token IDs, exact `VT_BENCH_PRETOKENIZE=0` retains timed-string admission, and `BenchResult` reports the resolved path (`examples/bench/bench_core.h:190,537-592,667,753`) | server-help/production-frontend/metric assertions plus #206 parser/callback-identity and synthetic default/rollback exact prompt/output-ID A/B: `tests/examples/test_bench.cpp:74,199` 6/6, 57 assertions CPU-GREEN; fresh mutation review and production GPU A/B remain pending | [CLI/serve/benchmark spike](specs/cli-serve-bench.md); [#206 campaign contract](specs/sm120-qwen35-pareto-2026-08-09.md) | `PARTIAL` | - | | `SERVE-GATE-ONLINE` | Same-corpus online correctness, TTFT/TPOT/ITL, throughput and peak-memory gate vs vLLM v0.25.0 | T0 | `vllm/benchmarks/serve.py:1,581-615`; [v0.25 audit](sync/2026-07-12-702f481.md); `tests/benchmarks/test_serve_cli.py:1` | Schema-v5 harness plus [trace controller](../include/vt/cuda/cuda_profiler_control.h#L13), [production component driver](../scripts/dgx-gdn-packed-component.sh), and fail-closed [component finalizer](../tools/bench/gdn_packed_component.py) | **BINDING `9ecd9d0`: 114/124** (async default ON; mem 4/4, c1 20/20, c2 20/20, c16 19/20, c4 & c32 18/20, c8 15/20; `benchmark_binding` refers here, superseding `3f256ab` 55/124 and `246a23c` 49/124, both retained immutable). Two-grid totality with `f0fb727` (111/124) is 115/124 effective parity vs vLLM 0.25.0 (27B). Async CLOSED the c16/c32 ITL tails (ours now BEATS vLLM: c16 p99 1.055, c32 p90 1.034/p99 1.078) and leaves a stable c8 `p99_itl` ~0.86 residual, ROOT-CAUSED (2026-07-18, `CLAIM-C8-P99-TAIL-1`, [spec](specs/c8-p99-itl-tail-2026-07-18.md)) as IRREDUCIBLE-AS-MIRRORED: our deterministic synchronous forward keeps co-admitted c8 requests in byte-identical lockstep where vLLM's async-future jitter de-phases them; the c16/c32 INVERSION proves this is the trailing edge of the per-step determinism that wins c16/c32 + throughput, not a capability gap (scheduler + async placeholder byte-identical, `tests/vllm/v1/test_scheduler_wave.cpp:265`, [tail spec](specs/tail-stall-analysis-2026-07-16.md)). Full grid + per-binding forensics: roadmap_v1.md + parity ledger; no packed speed credit | [online serving gate](specs/cuda-online-serving-gate.md); [merged GDN projections](specs/gdn-merged-input-projections.md); [packed decode](specs/gdn-packed-decode.md) | `ANCHOR-BACKFILL` | CLAIM-SERVE-GATE-1 | | `SERVE-E2E-NIGHTLY` | Server conformance and real-model nightly suites for all release gates | T0 | `tests/entrypoints/openai/`; `tests/v1/e2e/`; `.buildkite/test-pipeline.yaml` | current unit/conformance tests only; no scheduled DGX suite | `tests/vllm/entrypoints/openai/test_conformance.cpp:1`; `tests/parity/test_qwen36_paged_engine.cpp:78`; `tests/parity/test_qwen27_paged_engine.cpp:110` | `planned: specs/server-e2e-nightly.md` | `INVENTORIED` | - | | `ENG-RELEASE-BINARIES` | Downloadable host-ABI-specific `vllm-server` bundles: adaptive CPU and fat CUDA primary artifacts, optional per-SM diagnostics, and literal-static feasibility boundary | T0 | vLLM release lanes `.buildkite/release-pipeline.yaml:1-18,34-170` @ `555967922`; release-image dependency boundary `docker/Dockerfile.cpu:262-290` | Required W1-W11/W13 implementation is complete in draft PR #196: ten-SM gencode/AOT, adaptive CPU tiers, extracted-archive validation and supply chain, least-privilege immutable handoff, eight primary bundles, byte-derived indexes, attestation, and exact-file publication; W12 remains optional/non-primary | Local manifest/archive/workflow/mutation gates, full adaptive x86 tier execution, clean CPU archive, and Vulkan 35/35 backend + 11/11 cross-device archive run are green; hosted ten-SM completion, full eight-tuple dry run, matching-hardware gates, and tagged publication remain pending, so no published-binary claim exists | [release-binary-matrix.md](specs/release-binary-matrix.md) | `ACTIVE` | `CLAIM-ENG-RELEASE-BINARIES-W1-W13` | diff --git a/.agents/roadmap_v1.md b/.agents/roadmap_v1.md index 54feca634..4aa10dba2 100644 --- a/.agents/roadmap_v1.md +++ b/.agents/roadmap_v1.md @@ -101,7 +101,7 @@ models we already ship + benchmark. Full seam map + M0–M5 W-plan: | MEM | `ROAD-V1-MEM` | **Memory budgeting: auto-size to the declared workload by default, optional total-footprint cap, pre-flight error instead of an OOM (user-directed 2026-08-06, [#83](https://github.com/mudler/vllm.cpp/issues/83))** — the user-facing wart that every operator hits before they hit any perf question: vLLM makes you compute your own VRAM budget, express it as a PERCENT, and nail it or OOM | [coverage view §2](feature-matrix.md#2-kv-cache--memory), [porting inventory](porting-inventory.md) | — **M1+M2 LANDED 2026-08-08** ([`specs/kv-sizing.md`](specs/kv-sizing.md)): the absolute `--kv-cache-memory` knob sizes the pool via a group-aware `KVBytesPerBlock` divisor, `--num-blocks` is the override, both mirrored on the C ABI at v16; `ResolveNumBlocks` precedence `num_blocks > bytes > 256`, CPU-gated (`KVBytesPerBlock` 5/5 + `test_capi` v16). M3 (the `gpu_memory_utilization` profile run) stays dgx-gated. (M0 design grounded in vLLM `config/cache.py` + `gpu_worker.py:497-599`; GB10 unified-pool caveat carried) | `M1+M2 DONE` | **Source-verified 2026-08-06 (records-only, NO code).** WE ARE CURRENTLY BEHIND vLLM ON THIS AXIS, not ahead: there is NO memory profiling at all and the KV pool is a RAW BLOCK COUNT the user types by hand — `EngineParams::num_blocks = 256` (`include/vllm/entrypoints/model_loader.h:58`, beside `block_size = 32` `:57` / `max_model_len` `:59` / `max_num_seqs = 8` `:60`), exposed verbatim as `--num-blocks N` (`examples/server/main.cpp:100,203-204,370`), carried on the C ABI as `vllm_model_params.num_blocks` at the same 256 default (`src/capi/vllm_c.cpp:429,485`), landing as `BlockPool(num_gpu_blocks, ...)` which asserts `> 0` and otherwise TRUSTS it (`include/vllm/v1/core/block_pool.h:96,223`; `src/vllm/v1/core/block_pool.cpp:51`). So a user must convert "40 GB free, 32k context, concurrency 8" into a block count themselves — strictly worse ergonomics than a percentage. UPSTREAM HAS THREE KNOBS, all `config/cache.py`, all T0, all rowed at [porting-inventory.md](porting-inventory.md) §T0: `gpu_memory_utilization` (fraction of TOTAL, default 0.9), `kv_cache_memory_bytes` (absolute KV pool) and `num_gpu_blocks_override` (exact block pin), sized off a profile run as `total x utilization - non-torch - peak activation`. **Mirroring that is necessary but does NOT solve the complaint**, for three source-grounded reasons: (1) the fraction is of TOTAL not FREE, so on any shared card the right fraction is a function of what someone else already holds — exactly the arithmetic the user is being asked to do; (2) WEIGHTS LOAD BEFORE THE KNOB ENGAGES (utilization sizes the KV pool AFTER the model is resident), so an oversized model OOMs during load and never reaches the check — which is the failure operators actually hit; (3) 0.9 is taken whether or not it is needed (a 4B model on an 80 GB card reserves 72 GB it will never touch and blocks everything else on the device). THE DESIGN (user-ratified 2026-08-06) is three modes: **Mode 1 default = SIZE TO THE DECLARED WORKLOAD** — per-allocation-class accounting BEFORE allocating anything (weights from checkpoint metadata / safetensors header / GGUF manifest, known before reading a byte of tensor data; CUDA context measured at creation; peak activation from a profile run at `max_num_batched_tokens`; KV for `max_model_len x max_num_seqs` at the resolved `block_size`/KV dtype; CUDA-graph capture-set footprint) → allocate exactly that and LEAVE THE REMAINDER OF THE DEVICE FREE. This is the SURPASS over vLLM, which takes its 90% regardless of whether the workload needs 8 GiB or 80. **Mode 2 = a cap on the TOTAL ENGINE FOOTPRINT** (weights + activations + KV + graph pools + context), NOT on the KV pool alone — load-bearing, because a KV-only cap cannot prevent the weight-load OOM; three spellings of the same cap: `--memory-limit 40GiB` (absolute, the primary form), `--gpu-memory-utilization 0.85` (vLLM's exact flag name and fraction semantics so existing vLLM launch lines port unchanged, per [[mirror-vllm-always-no-asking]]) and `--num-gpu-blocks-override N` (upstream's reproducibility escape hatch — this is where today's `--num-blocks` GOES, demoted from primary knob to explicit override), with precedence spelled out and TESTED, not left to argument order. **Mode 3 = REFUSE BEFORE ALLOCATING** with the full per-class breakdown and remedies COMPUTED FROM THE ACTUAL BUDGET (`--max-model-len N` / `--max-num-seqs M` / `--kv-cache-dtype fp8` / smaller quant, each with the GiB it recovers) — "you are 43.9 GiB over and here are the three levers that close it" is the difference between a usable error and a stack trace. **UNIFIED-MEMORY HAZARD (not hypothetical):** on GB10 the ~119 GiB pool is UNIFIED, a fraction-of-total setting reserves HOST RAM too, and `gpu_memory_utilization=0.85` has HARD-REBOOTED our DGX three separate times ([[gb10-unified-memory-oom-reboots-box]]) — hence absolute bytes is the PRIMARY form with the percentage kept only for vLLM compatibility, and hence the accounting must know whether the pool is unified, which makes free/total + an is-unified predicate a PLATFORM-SEAM question belonging behind `ROAD-V1-C1`'s abstraction (note `Platform::needs_weight_staging()` was deliberately NOT `is_unified_memory()` because the latter FLIPS GB10 — the distinction matters here) rather than a CUDA-specific branch, since discrete and unified devices need different safety margins. CORRECTNESS: pool size changes preemption/scheduling TIMING but not emitted tokens, so the SACRED token-exact gates are unaffected — and M2's gate makes that explicit by re-running them with no block flag at all. **Next gate = M0 the `specs/kv-sizing.md` spike (accounting model + precedence rules + upstream `file:line`); then M1 a `MemoryBudget` computing required bytes per class WITHOUT allocating plus the Platform free/total + is-unified seam, unit-gated predicted-vs-actual weight bytes; M2 auto-sizing as the default with `--num-blocks` demoted to `--num-gpu-blocks-override`, gated by every existing model gate running with NO block flag and staying token-exact; M3 the three caps + precedence through the server flags and the C ABI, gated by our KV pool matching vLLM's own at a matched `--gpu-memory-utilization`; M4 the pre-flight refusal, gated by a deliberately over-subscribed config exiting cleanly on GB10 (non-zero exit, no OOM, NO BOX REBOOT) covering the unified-pool path specifically; M5 (optional) a runtime guard failing the REQUEST rather than the engine.** Docs (README, [STATUS](../docs/STATUS.md)) update in the SAME change as whichever milestone shifts externally-visible behaviour, per [[keep-readme-current]] | | 1 | `ROAD-V1-C1` | **Extensibility-first (USER PRIORITY 2026-07-18):** drop-in kernel ABI + the MISSING Platform seam + model self-registration — make new GPUs/archs/models ADDITIVE (plan: [extensibility-platform-seam-2026-07-18.md](specs/extensibility-platform-seam-2026-07-18.md)) | [`BACKEND-ABI-VT`](backend-matrix.md), [kernel matrix](kernel-matrix.md) | exhaustive kernel/dependency inventory and [raw-pointer adapter ABI](specs/dropin-kernel-abi.md) accepted; additive W0 implemented and CPU 94/94. `CLAIM-BACKEND-ABI-W0-GPU-1` repaired the GCC13/doctest blocker without runtime changes; exact sm_121a all-target build, focused CUDA/ABI sanitizer, and both gate-model tests pass at `1141b79`. Cross-arch/trace/A-B and scalar-forwarder/backend-shim debts remain explicit | `PARTIAL` | **★ NEW ORDER-1 HEAD (user-directed 2026-07-19): the PORTABLE AUTOMATIC OP-FUSION FRAMEWORK (`KERNEL-FUSION-FRAMEWORK`, spike [portable-fusion-framework.md](specs/portable-fusion-framework.md), `SPIKE`).** The extensibility cornerstone: fusions DECLARED ONCE (backend-agnostic `constexpr FusedRecipe` catalog above `vt::`, transcribing vLLM's finite pattern-pass set `passes/fusion/*` @ `pass_manager.py:138-200`, mirroring the `CustomOp` `forward_native`/`forward_cuda` seam `custom_op.py:103`) and REALIZED PER-BACKEND through the existing `vt::` op table (Tier-0 composite = the CPU oracle inherited free by any backend; Tier-1 interpreter = one kernel port per backend lights up every recipe). Makes a new vLLM fusion PR a ONE-declaration port, a new GPU an additive catalog realization, a new model an additive pattern declaration — the PR-#4 remedy, composed with the Platform/attn-registry/model-registry seams below. The TDR Phase-0 skeleton is already LANDED (`fused_recipe.h`/`recipes.h` one recipe + `OpId::kFusedChain` Tier-0/1 on CPU+CUDA + byte-exact `test_ops_fused_chain.cpp`). **W0 ADOPTED 2026-07-19 (`CLAIM-FUSION-FRAMEWORK-W0`):** the seam is now used in production at ONE real site — the 35B `RunLayerPaged` post-attention layernorm routes its plain add+residual+gemma-RMSNorm through `vt::FusedChain(kFusedAddRmsNorm)` (`VT_FUSED_CHAIN_ADOPT` default-ON / `=0` rollback), behaviour-preserving + byte-identical to the prior hand-call (DGX: clean CUDA `-Werror` 0 warn, byte-exact composite==interp==golden incl. H=2048, 35B 315/315 + 27B 235/235 token-exact BOTH arms, memcheck 0 errors) — proving the declare-once/realize-per-backend seam end-to-end; the current 3-opcode/4-role POD sufficed byte-identically for the plain add+rmsnorm class, so W0 needed NO generalization. **W1 GENERALIZED the POD 2026-07-20 (`CLAIM-FUSION-FRAMEWORK-W1`, `1115648`):** full activation/norm/quant/rope opcode set + indexed operand table; all 5 quant-fused chains declared byte-exact; Tier-0 composite = ONE device-agnostic walker (kills CPU/CUDA oracle drift); infrastructure only, no call site changed (DGX: `-Werror` 0-warn, byte-exact CPU 196 + CUDA 361, memcheck 0, 27B 235/235 + 35B 315/315 both arms). **W2 MIGRATED the hand-fusions 2026-07-20 (`CLAIM-FUSION-FRAMEWORK-W2`):** the framework now OWNS the fusion dispatch — each recipe binds (new backend-agnostic `FusedRecipe.fast_op`) to its EXISTING single-launch fused kernel, so `FusedChain(recipe)` dispatches to the SAME fast kernel the model called directly pre-migration (byte-identical + perf-neutral by construction; composite is the graceful fallback + oracle). SIX call sites migrated behind `VT_FUSED_CHAIN_ADOPT` (`kSiluMulFp4Quant`/`kSigmoidGateFp4Quant`/`kRmsNormGatedQuantFp8`×2/`kRmsNormQuantFp8`/`kAttnQkNormRopeGate`×2). DGX: `-Werror` 0-warn, byte-exact CPU 228 + CUDA 420, memcheck 0, 27B 235/235 + 35B 315/315 BOTH arms. **W3 MECHANICAL-SYNC PROOF LANDED 2026-07-20 (`CLAIM-FUSION-FRAMEWORK-W3`):** ported a NEW, previously-unported vLLM fusion pass — `SiluMulFp8StaticQuantPattern` (`act_quant_fusion.py:81` → `_C.silu_and_mul_quant`, the static-per-tensor-FP8 sibling of `kSiluMulFp4Quant`) — as ONE `constexpr FusedRecipe kSiluMulQuantFp8` + its byte-exact test, touching EXACTLY 2 shared files (`recipes.h` + `test_ops_fused_chain.cpp`), NO kernel/dispatch/composite-walker/model-site edit and NO new primitive (composite = existing `vt::MoeSiluMul` + `vt::QuantFp8Static`; `fast_op=kNoFastOp`). The PR-#4 additivity test made concrete: a whole new fusion pattern = one declaration. DGX: `-Werror` 0-warn, byte-exact CUDA 432, memcheck 0, no token regression (recipe declared-only) 27B 235/235 + 35B 315/315. **W4 BACKEND-ADDITIVITY PROOF LANDED 2026-07-20 (`CLAIM-FUSION-FRAMEWORK-W4`) — the W-series proof milestone is DONE:** made the additivity claim EXECUTABLE — new test `test_fused_chain_additivity.cpp` treats the EXISTING CPU backend AS the 'second backend' relative to CUDA (no mock `DeviceType` — that would edit the core enum + every switch, ironically non-additive) and, in ONE generic loop over the WHOLE catalog (all 7 recipes), asserts each runs byte-exact on the CPU backend via the Tier-0 composite — 4 CPU-full end-to-end + 3 fp8-prefix (byte-exact prefix + the FULL composite asserted to THROW on CPU, documenting the CUDA-only static-fp8 backend-negotiated tail, §3b/§6). Additivity evidence: catalog `recipes.h` grew 1→6→7 while the composite walker stayed ONE per-OPCODE function (12 `FOp::` cases) + the CPU/CUDA `kFusedChain` registration ONE line each + `cpu_ops.cpp` never `#include`s `recipes.h` — W3's whole new recipe `kSiluMulQuantFp8` is in ZERO backend TUs, inherited free. CPU `-Werror` 0-warn, `test_fused_chain_additivity` 17/17 + `test_ops_fused_chain` 228/228; engine byte-identical (no `src/`/`include/` change) ⇒ 27B 235/235 + 35B 315/315 structurally unchanged; memcheck N/A. Honest deferred (named, non-blocking the ORDER-1 milestone): Tier-1 perf interpreter for the quant chains (composite-only today), a REAL Metal/Vulkan catalog realization (M4 HW-blocked), and per-recipe fast single-launch kernels. Honest payoff: perf ceiling ~3.5%/step compute-bound on 35B (NOT a perf lever — tasks #61/#62; W0 is perf-neutral by construction); primary value = extensibility + mechanical upstream-sync + CPU/CUDA oracle-drift elimination. Incremental W0 adopt-one **(DONE)** → W1 generalize POD **(DONE)** → W2 migrate hand-fusions **(DONE)** → W3 mechanical-sync proof **(DONE)** → W4 backend-additivity proof **(DONE)** → Wn honest re-measure (optional, off the extensibility critical path). **W-SERIES ORDER-1 PROOF MILESTONE DONE 2026-07-20 (`CLAIM-FUSION-FRAMEWORK-W4`).** **PRIOR extensibility items (all LANDED, the seams this composes with):** **#1 extensibility item — extract the Platform seam — LANDED 2026-07-18** (`BACKEND-PLATFORM` `ACTIVE`, `CLAIM-BACKEND-PLATFORM-1`): `include/vllm/platforms/interface.h` + `src/vllm/platforms/{platform,cpu,cuda}.cpp` mirror `vllm/platforms/interface.py:134-229` 1:1; `CurrentPlatform()` self-registered per `DeviceType`; the 7 memory-model/residency `device.type == kCUDA` sites (of PR #4's ~37) now route through it → new-GPU memory model is ONE additive `platforms/.cpp`. Behavior-preserving (clean CPU build + `test_platform` + full CPU CTest green; DGX 235/235 + 315/315 pending). **Item 2 residency-as-Platform-capability LANDED 2026-07-19** (`CLAIM-BACKEND-PLATFORM-2`): the host-free / load-stream / DevicePool-cap decisions in `qwen3_5.cpp` now READ `GetPlatform(.device.type).residency_policy()` (per-device) via the pure `ShouldReleaseHostWeights`/`ShouldInterleaveLoadStream` helpers + `device_pool_cap_bytes`, not an inline `device.type`/env gate; `CudaPlatform.release_host_weights_after_upload` flipped false→true (now CONSUMED ⇒ reproduces today's GB10 host-free-after-Marlin + ~4 GiB load peak EXACTLY); `MarlinMoeEnabled()` stays the orthogonal kernel-path gate. **A new (discrete) GPU sets `residency_policy()` values ⇒ ZERO model edit.** Behavior-preserving (clean CPU build + `test_platform` consumption cases 7/43 + full CPU CTest + tools 164/164 green; **DGX-CONFIRMED @ `62fc0e0`: clean CUDA `-Werror`, 27B 235/235 + 35B 315/315 token-exact, 35B VmHWM ≈ 4.0 GiB load-stream win preserved, memcheck 0 errors**). Then item 3 drop-in ABI family migration. **Item 4 attn-backend registry LANDED 2026-07-19** (`CLAIM-ATTN-REGISTRY-1`, `BACKEND-ATTN-REGISTRY`): NEW `include/vllm/v1/attention/registry.{h,cpp}` `(DeviceType,name)` registry + `SelectAttentionBackendName` selector (mirror `registry.py` self-registration + `cuda.py:361-470` `get_attn_backend_cls`/`_get_backend_priorities`); `Platform::get_attn_backend_priority()` filled (was the item-1 STUB) → capability-ordered name lists on `CudaPlatform` (major-10 vs else) + `CpuPlatform`; FLASH_ATTN/GDN self-register. **Adding a backend's attention = 1 self-registering TU + 1 priority slot, ZERO selector/model/runner edit.** Behavior-preserving — the walk returns FLASH_ATTN on CUDA+CPU (the same FA2 attention runs); clean CPU build + `test_attn_backend_registry` (8/25) + full CPU CTest, tools 164/164, checkers green; **DGX-CONFIRMED @ `2c732e7`: 27B 235/235 + 35B 315/315 token-exact (FA2 sm_121a), memcheck 0/315**. **Item 5 model self-registration LANDED 2026-07-19** (`CLAIM-MODEL-SELFREG-1`, `MODEL-FACTORY-registry`): the fixed `kRegistrations` array → `REGISTER_VLLM_MODEL(...)` static-`Registrar` idiom (`model_registry.h:167-189`) + Qwen dense/MoE arch entry points split into per-variant TUs (`qwen3_5_dense.cpp`/`qwen3_5_moe.cpp`) over shared `qwen3_5_common.{h,cpp}`, so **adding a model = 1 new TU + 1 REGISTER line, ZERO shared-array edit**; behavior-preserving (clean CPU build + `test_model_registry` extension + full CPU CTest, tools 164/164, checkers green; DGX 27B/35B token-exact pending). Deep `qwen3_5.cpp` machinery factoring deferred. Metal/MLX bring-up proves the seams (needs M4). **★ THE ARCH HALF OF THIS ITEM IS NOW PROVEN BY MEASUREMENT, NOT ARGUED (2026-07-22, `CLAIM-CUDA-SM120-BRINGUP`, [spec §W8](specs/cuda-arch-additivity.md)):** consumer-Blackwell `sm_120a` was brought up as a BUILD-supported target and required **ZERO kernel, model, runner, sampler or feature-table edits** — the additive seams (per-arch FEATURE TABLE, capability-keyed tactic registry keyed on `major == 12`, Platform auto-probe, `pageable && integrated` residency classification) already covered it, so the entire diff is build configuration, a configure-tier test and records. That is the PR-#4 additivity test passed on a real second architecture. It is deliberately NOT a runtime-support claim: no sm_120 board exists here. **★ THE MODEL/QUANT HALF NOW ADVANCES TOO (2026-07-23, `CLAIM-BACKEND-SEAM-S4-1`): the `model_executor/layers/` `LinearMethod`/`QuantizationConfig` seam the [accelerator-seam audit](specs/accelerator-seam-audit.md) §9 named ABSENT now EXISTS in part.** `S4` landed byte-identical: the dense model's projections route through a `method.Apply()` chosen ONCE from the checkpoint (retiring the per-call `IsNvfp4()` tensor-name probe), and 18 shared-layer `device==kCUDA` availability gates became `vt::OpRegistered` op-table queries — the policy(scheme)/implementation(kernel) split the audit's binding rule requires. **DSR 86 → 67**; all six SACRED gates byte-identical on dgx (27B/35B/Coder/dense/OPT/DeepSeek-V2); the fragile 27B-W4A4/fp8-recipe gates are correctly deferred to `S6` behind `S5`'s reference tier. **★ `S6` ASSESSED 2026-07-23 (`CLAIM-BACKEND-SEAM-S6-1`) → NO-OP / BLOCKED, DSR stays 67 (§11):** the deferred fp4/fp8 gates convert ZERO sites byte-identically — every one bottoms out at a **dual-registered** (CPU+CUDA) bespoke op (none CUDA-only, unlike S4's convertible gates), so `OpRegistered(op,dev)` is TRUE on `kCPU` ⇒ the class-A swap is bit-changing on the CPU reference/emulation path (two numerics per device); S5's reference tier does not change this (those CPU kernels are present natively, never a miss). No `src/`/`include/`/test byte changed, no baseline moved. The genuine byte-identical unlock is re-scoped to **`S3`** (Platform capability fields mirroring `supports_fp8`/`cutlass_fp4_supported` — the audit's own class-D fix) and **`S7`** (layer extraction); the plan's `~37` S6 target assumed the class-A `OpRegistered` swap was byte-identical, which holds only for CUDA-only ops (all taken by S4). **★ `S3` LANDED 2026-07-23 (`CLAIM-BACKEND-SEAM-S3-1`) — the byte-identical unlock S6 re-scoped:** mirrors vLLM's `Platform` capability surface (`supports_fp8`/`cutlass_fp4_supported`/`opaque_attention_op`/`is_integrated_gpu`/`support_static_graph_mode`/`is_device_capability_family`, base false in `interface.h`; CudaPlatform answers GB10 values in `cuda.cpp`/`platform.cpp`) and converts **12** deferred `qwen3_5.cpp` gates onto it (7 fp4-act `cutlass_fp4_supported`, 3 fp8-fused `supports_fp8`, 2 decode-graph `support_static_graph_mode`) — byte-identical because a capability answers the base false off CUDA, exactly what `device==kCUDA` did (where S6's `OpRegistered` was TRUE on `kCPU`), and it DECOUPLES (a future accelerator answers for itself). **DSR 67 → 55** (`kcuda` 25→13), baseline lowered same commit, ratchet + 24-case suite green; all seven SACRED gates byte-identical on dgx (27B 235/235 · 35B 315/315 · Coder · dense-32B · OPT · DeepSeek-V2 · Llama), `test_platform` CUDA-leg proves each predicate == former `device==kCUDA` on GB10, memcheck 0 errors, clean CUDA+CPU `-Werror`. Residency/stream/FA2-dtype/merged-layout sites LEFT for `S7`. **★ `S7` LANDED 2026-07-23 (`CLAIM-BACKEND-SEAM-S7-1`) — the seam campaign's TERMINAL runtime-decoupling state, the closest this extensibility work comes to a finish line:** ALL 23 remaining runtime `kCUDA`/`is_cuda()` sites in the shared model layer hoisted onto capabilities — new `Platform::needs_weight_staging()` (the CUDA device-resident staging policy, NOT `is_unified_memory()` which would FLIP GB10; covers residency/merged-GDN/packed-decode/direct-load), `Platform::supports_fa2_attention()` (FA2 dtype), `Backend::SupportsAuxStream()` (MoE aux-stream), reusing S3's `supports_fp8`/`cutlass_fp4_supported`/`support_static_graph_mode`/`is_integrated_gpu` (runner combine/scatter) and `vt::OpRegistered(kMoeGroupedGemmBf16)` (a CUDA-only op). Each returns the former `device==kCUDA` value on GB10 → byte-identical. **DSR 55 → 32 — the IRREDUCIBLE build-gate floor:** the shared model layer holds ZERO runtime device tests; the 32 residual are all `#ifdef VT_*` compile-time gates for kernels that only build on one GPU family (a kernel that only compiles on one arch is legitimately irreducible), so the audit's `<10` is NOT reachable and this is the honest answer to "how additive can the shared layer get" — every runtime device coupling is gone. baseline lowered same commit, ratchet + 24-case suite green; all seven SACRED gates byte-identical on dgx (27B 235/235 · 35B 315/315 · Coder 6/6 · dense-32B 16/16 · OPT 6/6 · DeepSeek-V2 8/8 · Llama 16/16), new `test_platform`/`test_backend`/`test_cuda_backend` cases green, memcheck 0, clean CUDA+CPU `-Werror`. The `layers/`-library physical relocation (shrinking `qwen3_5.cpp` toward `qwen3_next.py`'s 802-line shape) is a follow-on refactor; the device coupling it was to remove is already gone | | 2 | `ROAD-V1-C2` | Model families: Llama/Qwen3/Mistral, MoE, Qwen3-Next | [model matrix](model-matrix.md) | current pin has 353 static IDs; v0.25.0 adds three sync-target rows (MOSS-Transcribe-Diarize, Laguna DFlash, Bailing hybrid MTP), yielding 356 after pin advance. **FIRST ADDITIVE-MODEL BRING-UP W0-W4 LANDED 2026-07-20 — CORRECTNESS COMPLETE (0.6B + 4B gates PASS 16/16; SPEED pending)** ([first-additive-model-qwen3-dense.md](specs/first-additive-model-qwen3-dense.md), `MODEL-TEXT-qwen3-qwen3-for-causal-lm` `ACTIVE`(correctness COMPLETE, speed pending), runner generalization `ENG-RUNNER-MODELSHAPE`, `CLAIM-MODEL-QWEN3-DENSE`) **MLA CAMPAIGN SPIKED 2026-07-21** ([mla-deepseek-campaign](specs/mla-deepseek-campaign.md), `CLAIM-MLA-DEEPSEEK`): five rows `INVENTORIED` -> `SPIKE` (DeepSeek V2 / V3+V3.2 / v1-MHA, Kimi-Linear, MiniMax-M2). **KIMI-LINEAR-48B W0 DEDICATED SPIKE 2026-08-05** ([kimi-linear.md](specs/kimi-linear.md), `CLAIM-KIMI-LINEAR-W0`): full dedicated W0 spike for `MODEL-TEXT-kimi-linear-*` (stays `SPIKE` — actively claimed) — the ONE Kimi text model that FITS one GB10 (91.5 GiB, 0.77x pool) with a real e2e SACRED gate; HEAVY reuse (MLA + sigmoid/`noaux_tc` MoE + GDN family + KDA host refs landed), NET-NEW = the KDA device kernel + NoPE-MLA branch + hybrid schedule/loader; W1 implementation can start. Answers the Tier-3 "MLA = new attention, new campaign" item in [breadth-sweep-plan](specs/breadth-sweep-plan.md) §B.3. Key determinations: GB10/sm_121 selects **`TRITON_MLA`** for dense MLA decode and **`FLASH_ATTN`** for MLA prefill (`vllm/platforms/cuda.py:129-133`, `mla/prefill/selector.py:74-77`), so the sm90/sm100-only MLA kernels are out of reach and out of scope; the cross-cutting cost is the **compressed-latent KV cache** (one 576-wide latent per token, `num_kv_heads=1`, no separate V), which our allocator and `vt::ReshapeAndCache`/`vt::PagedAttention` cannot express; and **only DeepSeek-V2-Lite (~29.3 GiB bf16) fits GB10** — V3/V3.2, Kimi-K2.5, MiniMax-M2/M3 are HW-BLOCKED e2e, Kimi-Linear-48B is HW-MARGINAL. W0-W10 plan recorded; nothing implemented. **GLM + DSA + LATEST-DEEPSEEK SPIKED 2026-07-21** ([glm-dsa-latest-deepseek](specs/glm-dsa-latest-deepseek.md), `CLAIM-GLM-DSA-LATEST-DEEPSEEK`): seven rows `INVENTORIED` -> `SPIKE` (ChatGLM, Glm, Glm4, Glm4Moe, Glm4MoeLite, GlmMoeDsa, DeepSeek-V4). Answers the user's "also glm, and deepseek latest versions" priority. Headline: **`Glm4MoeLiteForCausalLM` / `zai-org/GLM-4.7-Flash` (31.2B, 58.2 GiB bf16) FITS GB10 and is a SECOND MLA gate vehicle that closes BOTH coverage gaps the MLA campaign named as unit-gated-only** (it has `q_lora_rank=768` and `noaux_tc`/`e_score_correction_bias`, which DeepSeek-V2-Lite lacks). **DSA is DOUBLY BLOCKED on GB10:** the sparse XOR filter eliminates `TRITON_MLA` for sparse models leaving `FLASHINFER_MLA_SPARSE_SM120` as the sole candidate, and that path is non-functional on flashinfer 0.6.12 (XQA backend is dense-only, discards `sparse_mla_top_k`); GLM-5 is 1404 GiB and V3.2 642 GiB regardless. DeepSeek-V4 is a NEW architecture (Sinkhorn-normalized Manifold Hyper-Connections, CSA/HCA compressor with recurrent state, hash-routed MoE) and HW-BLOCKED at 148.7 GiB — but its TOKENIZER risk is LOW (standard HF fast BPE; only the chat template needs porting, with upstream golden fixtures). Glm4/Glm need two primitives we have NONE of: partial rotary factor and sandwich norms. Nothing implemented. **NVFP4A16 (W4A16)** on the already-done dense `Qwen3ForCausalLM` (`RedHatAI/Qwen3-32B-NVFP4A16`, 64L) — the QUANT-SCHEME additivity experiment, serving user priorities #2 (models) and #4 (quants) at once. KERNEL LAYER FULLY ADDITIVE (ZERO new kernel code: vLLM FORCES Marlin for `use_a16`, OBSERVED `Using MarlinNvFp4LinearKernel`, and that is the GEMM we already vendored for the 35B). **CORRECTNESS CLOSED 2026-07-21 (W4b):** the strict gate's 4/6 was diagnosed by the ratified TEACHER-FORCING isolation — all 29 divergent positions gap <= 0.0625 nats with 28/29 EXACTLY 0.0, one root flip an EXACT bf16 tie at which vLLM's teacher-forced argmax is OURS and vLLM contradicts its own greedy. **NOT a W4A16 defect: it is the pre-existing dense-forward bf16 near-tie drift, recorded against `MODEL-TEXT-qwen3-qwen3-for-causal-lm`.** Gate closes **6/6** under the ratified near-tie-robust bar with the nats evidence committed. SPEED still pending ⇒ `ACTIVE`, not `DONE`. **GEMMA FAMILY SPIKED 2026-07-24** ([sweep-gemma](specs/sweep-gemma.md), `CLAIM-SWEEP-GEMMA`): four rows `INVENTORIED` → `SPIKE` (Gemma 1/2/3/4). Answers the user's "and then we do gemma" ("gemma 4") next-target. **The newest registered Gemma is Gemma 4** (real, public, but ALL checkpoints multimodal-wrapped `Gemma4*ForConditionalGeneration`, ≥12B, 0.25.0 oracle-support unverified, needs a PLE/YOCO/MoE/k_eq_v stack) — it leads the characterization but is gate-BLOCKED as a first vehicle. **The recent-first gate vehicle that FITS + is oracle-certain is Gemma 3** (`Gemma3ForCausalLM` on `google/gemma-3-1b-it`). Headline: Gemma reduces MOSTLY to landed infra — gemma-RMSNorm, sandwich norms (glm4 `b568d20`), SentencePiece (names "Gemma"), sliding-window (FA-2 + SlidingWindow/ChunkedLocalAttention specs), the `kAttnQkNormRopeGate` QK-norm+rope recipe, tied embeddings are ALL REUSE; the one genuinely-new compute kernel is GeGLU (`gelu_pytorch_tanh`+mul; we have only SiLU), plus the final logit soft-cap + qpas/embed-scale scalars + dual-rope routing. Per-version delta: Gemma-2 has an attn logit soft-cap, Gemma-3 removed it for QK-norm. **GEMMA-3 W0-W2 LANDED 2026-07-24 — CORRECTNESS COMPLETE, the FIRST Gemma family** (`MODEL-TEXT-gemma3-gemma3-for-causal-lm` `ACTIVE`, speed pending): `Gemma3ForCausalLM` on `google/gemma-3-1b-it`. W1 = two additive default-inert vt ops `kGeluAndMul` (GeGLU `gelu_pytorch_tanh`) + `kMulScalar` (bf16 embed-scale), CUDA+CPU, unit 12/12. W2 = `gemma3.{h,cpp}`/`gemma3_weights.cpp`/`gemma3_registry.cpp` reusing the GLM-4 sandwich-norm layout + `dense_attn_block.h` + FA-only KV: GemmaRMSNorm `(1+w)`, per-head Gemma q/k norm, dual per-layer RoPE theta, qpas scale, per-layer sliding window, GeGLU MLP, `sqrt(hidden)` embed-scale, tied lm_head. **SACRED gate STRICT token-exact 48/48** greedy vs vLLM 0.25.0 (K=5 ALL-DETERMINISTIC → STRICT; BOS-verified; tokenizer-free like Mistral's `LOAD-SENTENCEPIECE` path). Loader 340 tensors, registry 23/23, clean `-Werror` 0 warn. **GEMMA-2 + GEMMA-1 W3-W6 LANDED 2026-07-24 — CORRECTNESS COMPLETE** (`MODEL-TEXT-gemma2-gemma2-for-causal-lm` + `MODEL-TEXT-gemma-gemma-for-causal-lm` `ACTIVE`, speed pending): W3 = the logit soft-cap primitives (`vt::SoftCap` final cap + `PagedAttentionArgs.logits_soft_cap` attention cap threaded into the native/CPU/FA-2 attention, default-inert). W4 `Gemma2ForCausalLM` (gemma-2-2b-it) = the inverse of Gemma-3 (BOTH soft-caps, NO QK-norm, single rope) — **near-tie-band SACRED 48/48** (44/48 strict + 4/48 at 0.0-nat vLLM-own ties, 0 forward-divergent; soft-cap PROVEN applied by a cap-on≠cap-off A/B). W5 `GemmaForCausalLM` (gemma-2b) = the original Gemma (two fused norms, head_dim scale) — **STRICT 48/48**. W6 = Gemma-4 honesty pass (HW/DEP-BLOCKED, not registered). Regressions byte-identical (Gemma-3 48/48, Qwen3-dense 184/184, OPT 63/63, Llama 92/92, Mistral 92/92) + DeepSeek-V2 asserts-on 223/223; compute-sanitizer 0; clean `-Werror` 0 warn. Gemma-4 stays `BLOCKED`. | `PARTIAL` | **ACTIVE: the first additive-model bring-up = Qwen3 dense (`Qwen3ForCausalLM`) on `Qwen3-0.6B` BF16 — W0+W1 landed 2026-07-20.** W0 (config+registry stub: new TU `qwen3_dense.cpp`+`qwen3.h`, one `REGISTER_VLLM_MODEL`, full-attention-only KV spec, forward stub) + W1 (the RUNNER GENERALIZATION `ENG-RUNNER-MODELSHAPE`) are DONE and gated: dgx CUDA `-Werror` 0-warn, **27B 235/235 + 35B 315/315 token-exact UNCHANGED** (behaviour-preserving), new CPU runner tests RED(SIGSEGV)→GREEN, registry resolves `Qwen3ForCausalLM`, ASan/UBSan clean. The runner is now MODEL-SHAPE-AGNOSTIC (a full-attention-only KV config allocates+steps without the hybrid GDN path) → every future dense/non-hybrid arch adds new-files-only. Qwen3-0.6B is the only standard-dense arch with a checkpoint + runnable vLLM 0.25.0 oracle on dgx TODAY (no Llama/Mistral checkpoint present → Llama-first needs a download, sequenced as W-next for genuine cross-family additivity). **W2 loader + W3 forward LANDED 2026-07-20:** dense forward `qwen3.cpp` (`Qwen3DenseModel::Forward/ForwardDevice`) composed from vt:: ops + the fusion catalog (2 new byte-exact recipes: `kFusedAddRmsNormStd`, `kAttnQkNormRope`); bf16 attention numerics mirror vLLM. The first pure-dense bf16 model forced out + FIXED 2 genuine latent bugs: tokenizer `kQwen2Classic` (classic Qwen2/Qwen3 pre-tokenizer was hard-rejected) and `cuda_paged_attn.cu` WMMA prefill mistokenizing at head_dim≠256 (now gated to the validated d=256). **W4 CORRECTNESS COMPLETE 2026-07-20 — near-tie-robust gate PASSES on 0.6B AND a bigger 4B.** The 2026-07-20 razor's "vLLM greedy non-deterministic" premise was a BATCHING artifact: per-prompt (batch=1, the gate regime) vLLM 0.25.0 greedy is DETERMINISTIC (0.6B 0-multi/K=10, 4B 0-multi/K=5). Forward correctness is PROVEN by teacher-forcing vLLM on OUR exact prefix (`scripts/qwen3-neartie-gap.py`): at all-but-2 positions vLLM's own argmax given our prefix IS our token (gap 0.0000, bit-identical logprobs — our forward matches vLLM's prefill logits); residual flips are bf16 near-ties (0.6B ≤0.125 nats, 4B ≤0.25) where vLLM's own prefill argmax disagrees with its decode. Gate = our token within 0.5 nats of vLLM's teacher-forced argmax (strict where equal): **Qwen3-0.6B 16/16** (strict 12 + near-tie 4) and the **bigger-model complete-correctness proof Qwen3-4B (36L, GQA 32/8, hidden 2560, same forward code) 16/16** (strict 10 + near-tie 6). Regression 27B 235/235 + 35B 315/315 UNCHANGED, CUDA `-Werror` 0-warn, memcheck 0. Correctness-complete. **SPEED — d128 FA2 PREFILL + DECODE DEFAULT-ON 2026-07-20 (`Qwen3-4B` vs vLLM 0.25.0 production/graphed, in1024/out128) — big gap-close, still below vLLM, `MODEL-TEXT-qwen3-qwen3-for-causal-lm` stays `ACTIVE` NOT `DONE`:** implemented the dominant prefill lever (a d128 FlashAttention-2 varlen prefill — generalized the vendored FA2 launcher to head_dim 128, `VT_FA2_PREFILL_QWEN3` default-ON) and flipped the FA2 varlen d128 decode default ON (near-tie gate re-passes 16/16 on 0.6B + 4B). Total tput now 0.90× (c1)/0.62× (c8) (up from 0.80×/0.48×), c1 decode at parity (TPOT 1.04×, ITL P99 0.98× win); prefill A/B = +7%/+41% total, −55%/−48% TTFT. STILL failing TTFT median 5.85×/10.2× + total <1×: the full prefill STEP (not the attention kernel, now vLLM's FA2 family) is ~6× vLLM = non-attention glue (GEMM/MLP fusion) + host-side launch overhead (un-graphed prefill); plus c8 decode batch efficiency (TPOT 1.38×). Dominant residual lever = portable prefill-step fusion + graphed prefill (secondary = c8 split-KV decode occupancy). **RoPE cos/sin cache flipped DEFAULT-ON 2026-07-20** (`VT_QWEN3_ROPE_CACHE`): the opt-in blocker (an alleged FA2-split-KV-combine run-to-run nondeterminism) was GROUNDED + DISPROVEN — the paged engine is byte-deterministic run-to-run and the combine never launches on the gate (`num_splits==1`); goldens regenerated on the canonical `$HOME/cutlass-4.5.0` build (the flashinfer cutlass copy tips the 27B tok6 razor to 234/235; cutlass-4.5.0 = 235/235), gate 16/16 both, 27B 235/235 + 35B 315/315 unchanged. RoPE-ON closes total tput 0.90×→0.97× (c1) / 0.62×→0.82× (c8) and c1 TTFT ratio 5.85×→2.27×, still `ACTIVE`. **SPEED RE-BOUND 2026-07-21 (same-session, matching-recipe) — TTFT residual RESOLVED, cutlass claim CORRECTED:** the "TTFT 2.27×/5.85×" + "c8 ITL 4.3×" were BAD-DENOMINATOR/num-prompts artifacts — a fresh same-session vLLM capture gives c1 TTFT ~152 ms & c8 ITL P99 ~130 ms, and OURS WINS TTFT at both concurrencies (c1 0.90×, c8 0.38×). **c1 = effective every-axis parity** (tput 0.98× / TPOT 1.01× / TTFT+ITL wins); **c8 residual = decode** (tput 0.93× / TPOT 1.10× / ITL P99 1.12×), which nsys shows is 93% GPU-busy/compute-bound (small-M=8 `cutlass_80_wmma` projections). The **qkv-merge** (new GQA `QkvSplit` op mirroring vLLM `QKVParallelLinear`) was implemented + MEASURED NEUTRAL (doesn't cut decode FLOPs) ⇒ default-OFF. **CUTLASS CLAIM CORRECTED: 27B `test_qwen27_paged_engine` = 235/235 on the FLASHINFER cutlass build** (the "flashinfer ⇒ 234/235" was a build artifact). Stays `ACTIVE`; named residual = c8 decode-GEMM efficiency (a decode-fusion sub-campaign). **SWEEP MODEL #1 — Qwen3-Coder-30B-A3B (`Qwen3MoeForCausalLM`) W0+W1 LANDED 2026-07-21** ([sweep-qwen3-coder-30b.md](specs/sweep-qwen3-coder-30b.md), `MODEL-TEXT-qwen3-moe-qwen3-moe-for-causal-lm` `ACTIVE`, `CLAIM-MODEL-QWEN3-CODER`): the first full-attention BF16 MoE, composed from the done dense attention + the done 35B MoE experts (ZERO runner change). W0 = registry stub (`qwen3_moe_registry.cpp` + `qwen3_moe.h`, one `REGISTER_VLLM_MODEL`, full-attn-only KV, `is_dense_model=false`, W2/W3 throwing stubs). W1 = three behaviour-preserving refactors making the two done pieces reusable: (#1) dense `AttnBlock` + glue extracted to `dense_attn_block.h` (Qwen3-dense byte-identical), (#2) bf16 `MoeBlock` exposed cross-TU via `RunMoeBlock` (`qwen3_5_moe_block.h`; 35B untouched), (#3) no-shared-expert guard (inert for the 35B). Gated: dgx CUDA `-Werror` 0-warn; Qwen3-dense 0.6B+4B 16/16 + 27B 235/235 + 35B 315/315 UNCHANGED; registry resolves `Qwen3MoeForCausalLM`; memcheck 0. W2 bf16 loader → W3 forward → W4 near-tie token-exact → W5 fast bf16 grouped-MoE GEMM remain. Then Llama dense (download), Mistral, MoE families **SWEEP MODEL — GLM-4-9B-0414 (`Glm4ForCausalLM`) G2 LANDED 2026-07-24 — CORRECTNESS COMPLETE** ([glm-dsa-latest-deepseek](specs/glm-dsa-latest-deepseek.md), `MODEL-TEXT-glm4-glm4-for-causal-lm` now `READY` per the [live-state audit](specs/live-state-audit-2026-08-06.md), `CLAIM-GLM-DSA-LATEST-DEEPSEEK` amended) **GLM-4.7-Flash (`Glm4MoeLiteForCausalLM`, 31.2B MLA+MoE) G1 LANDED 2026-07-24 — SACRED gate 8/8, CORRECTNESS COMPLETE** (`MODEL-TEXT-glm4-moe-lite-glm4-moe-lite-for-causal-lm` `ACTIVE`, speed pending): reuses the DeepSeek-V2 MLA stack; first e2e coverage of the q_lora branch + noaux_tc sigmoid router, closing the MLA campaign's two C2 coverage gaps: the FIRST GLM-family model. SACRED gate 16/16 vs vLLM 0.25.0 (STRICT 13/16 + near-tie 3/16, max gap 0 nats; vLLM K=5 self-deterministic ⇒ STRICT bar), speed PENDING. The two "new primitives" the spike named reduced to EXISTING infra: partial + interleaved `RopeFromCache` (`is_neox_style=false`, the DeepSeek-V2 decoupled-rope path) over `rotary_dim=64`; standalone `vt::RmsNorm` sandwich norms. Biased qkv, no QK-norm, GQA 32/2, untied lm_head. New files + one REGISTER, reusing the shared dense glue. **SWEEP MODEL — Llama-3.2 (`LlamaForCausalLM`) W0-W4 LANDED 2026-07-23 — CORRECTNESS COMPLETE** ([sweep-llama-3.2](specs/sweep-llama-3.2.md), `MODEL-TEXT-llama-llama-for-causal-lm` `ACTIVE`, `CLAIM-MODEL-LLAMA-3.2`): the roadmap's explicit "Llama-first" increment and the first mainstream non-Qwen/non-OPT dense arch. `LlamaForCausalLM` (Llama-3.2-1B) = the Qwen3-dense forward with exactly two ADDITIVE deltas — NO qk-norm (shared `AttnBlock` skips it when q_norm/k_norm empty) + llama3 rope-scaling (4 default-0 `RopeArgs` fields + a `Llama3ScaleFreq` kernel helper, no-op elsewhere; formula verified 2e-7 rel vs vLLM) — reusing the shared dense forward VERBATIM (`LlamaModel == Qwen3DenseModel`). 3 new Llama files, ZERO edit to runner/scheduler/platforms/attn-registry/`hf_config`/any qwen3-opt model. vLLM 0.25.0 greedy MEASURED DETERMINISTIC (K=6, 0 multi-valued cells) ⇒ STRICT token-exact bar, PASS **16/16 (12 strict + 4 near-tie band, max gap 0.0000 nats, 0 divergent)** — at all 13 divergent positions vLLM's own teacher-forced argmax given our prefix IS our token. A correctness-fatal tokenizer bug (Llama's `Sequence` post_processor wrapping `TemplateProcessing` ⇒ BOS 128000 never prepended, silently 1/16) was isolated by a CUDA prefill-argmax diagnostic (forward proven 4/4 correct given vLLM's exact tokens) and fixed byte-preservingly (Qwen/OPT/DeepSeek unaffected — all ByteLevel/top-level-TemplateProcessing). Regressions 27B 235/235 · 35B 315/315 · Coder 6/6 · Qwen3-dense 16/16 · OPT 6/6 · DeepSeek-V2 8/8 UNCHANGED; `-Werror` 0-warn; memcheck 0; DSR 67. SPEED pending (head_dim 64 → generic paged path; Llama-3.2-3B head_dim-128 is the FA2-toggle W-next). **MLA CAMPAIGN W0+W1 LANDED 2026-07-21** (`CLAIM-MLA-DEEPSEEK`; rows STAY `SPIKE` — W0/W1 make no model supported). **W0 grounded every fact the spike flagged as an unverified source read; ALL CONFIRMED, none contradicted:** DeepSeek-V2-Lite fetched to dgx (30 GB, 4 shards) and loading in the vLLM 0.25.0 oracle; the oracle's own DEBUG startup on sm_121 prints `Using TRITON_MLA attention backend out of potential backends: ['TRITON_MLA']` and `Using FLASH_ATTN MLA prefill backend` — so the dense-MLA decode + MLA-prefill targets are OBSERVED, not inferred, and the sm90/sm100-only MLA kernel class stays out of scope; the real `config.json` confirms every §5.1 number (`kv_lora_rank=512`, `qk_nope=128`, `qk_rope=64` -> the **576-wide latent**, `v_head_dim=128`, `q_lora_rank=null`, `n_group=topk_group=1`, softmax/greedy, 64+2 experts, 27 layers) plus `is_neox_style=False` and the mscale2 scale correction; and BOTH recorded coverage gaps (no `fused_qkv_a_proj` branch, no `e_score_correction_bias`) are confirmed real, so those pieces stay unit-gated only. **W1 = the behaviour-preserving spec-driven KV allocation, ZERO MLA math:** the attention cache is now sized `num_blocks * spec->page_size_bytes()` and viewed from the spec's own `block_size`/`num_kv_heads`/`head_size`/`dtype` instead of the hardcoded `2 * block * Hkv * Dh` with shape reconstructed from the HF config (`runner.cpp`), plus `MLAAttentionSpec` with upstream's factor-1 single-tensor page formula (`kv_cache_interface.py:397-398`) registered against the ORDINARY `FullAttentionManager` (`single_type_kv_cache_manager.py:1539`) — the spike's key finding, which is why block manager/prefix caching/eviction need no change. Gated: dgx clean CUDA `-Werror` 0 warnings/0 errors; **27B 235/235 + 35B 315/315 + Qwen3-Coder 6/6 + Qwen3-dense 16/16 ALL UNCHANGED** (behaviour-preserving proven, not assumed); `test_runner` 15/15, `test_kv_cache_interface` 21/21 (4 new MLA-spec cases), `test_llm_engine` 5/5; the new path is proven EXERCISED (not merely compiled) by `fa_page_size_bytes()` + a `page_size_padded` case no HF-config formula can produce. **W2 + W3 LANDED 2026-07-21** (base `a05437f`; rows STAY `SPIKE` — still no MLA attention math, no MLA model, no forward). **W2 = the MLA branch of `_get_backend_priorities` the pre-W2 comment deferred, ported as DATA:** the whole of `cuda.py:84-176` (BOTH branches — MLA sm_100 including the `:96-115` adaptive sparse tail, MLA sm_12x, MLA `else`, and the two pre-existing non-MLA arms) is now a TABLE in the new header `include/vllm/platforms/cuda_attn_priority.h`, one row per upstream arch arm keyed on `(use_mla, major)`, so a future arch is a ROW rather than a code path; putting it in a header (not the CUDA-only TU) also let the CPU test tier assert the REAL table and DELETED the hand-copied `FakeCudaPlatform` duplicate. On sm_121 a `use_mla=true` request now RESOLVES to `TRITON_MLA`, matching the W0 oracle observation. **The sparse/DSA seam is left OPEN and unit-proven:** GB10's row keeps both upstream entries and the sparse one loses to a real FILTER — `AttentionBackend::is_mla()`/`is_sparse()` checked against the request (`backend.py:307-360 validate_configuration`) — so a future DSA backend is selected purely by declaring `is_sparse() == true`, with ZERO edit to the table or the selector. `TritonMLABackend` lands the NAME plus upstream's 3-D `get_kv_cache_shape` (no K/V axis; `num_kv_heads != 1` REFUSED), `get_impl_cls()` deliberately still `nullptr`. MLA prefill priority ported too (GB10 -> `[FLASH_ATTN]` alone). **W3 = the two new `vt::` ops, both CPU-reference-gated.** `vt::ConcatAndCacheMla` mirrors `csrc/libtorch_stable/cache_kernels.cu:401-442` — and per the whole-chain rule this was VERIFIED, not assumed, to be vLLM's OWN csrc kernel (`_custom_ops.py:2532` -> `torch.ops._C_cache_ops`), with no flashinfer/cutlass variant in the dense-bf16 path; it concatenates the latent + rope part into ONE 576-wide entry, the write `ReshapeAndCache`'s K/V-pair signature cannot express, stride-driven so a per-layer cache slice and the two column halves of `kv_a_proj_with_mqa` both work copy-free. The **grouped-topk (`noaux_tc`) router** extension — flagged in `coordination.md` as SHARED with `CLAIM-GLM-DSA-LATEST-DEEPSEEK` and "must not be implemented twice" — is landed HERE and that claim now consumes it: additive `MoeRouterTopKArgs` fields + an optional `e_score_correction_bias`, with `num_expert_group == 0` still dispatching the ORIGINAL kernel so the 27B/35B/Coder/dense routers are byte-identical BY CONSTRUCTION. **Stated plainly: the `noaux_tc` correctness evidence is UNIT-ONLY.** V2-Lite has `n_group=topk_group=1` and no bias, so the e2e vehicle exercises none of it; the gate is `tests/vt/test_ops_moe_router_grouped.cpp` at DeepSeek-V3's REAL dimensions (256 experts, n_group=8, topk_group=4, sigmoid, scaling 2.5, WITH the bias) against an INDEPENDENT sort-based transcription of the upstream formula. **W4 LANDED 2026-07-22** (base `ed2c342`; rows STAY `SPIKE` — W4 adds a kernel and fills a `nullptr`, it makes no model supported). **`vt::MlaDecodeAttention` — the MQA decode over the compressed latent (QK 576 / V 512, `num_kv_heads=1`), a structure port of the two-stage split-KV pair W0 OBSERVED EXECUTING:** `MlaDecodeStage1` <- `_fwd_grouped_kernel_stage1` (`triton_decode_attention.py:278-458`, the `IS_MLA` branch whose `v = tl.trans(k)` at `:424-431` is the whole MLA trick — V is the leading 512 columns of the SAME latent row already loaded as K, so one shared-memory tile serves as both), `MlaDecodeStage2` <- `_fwd_kernel_stage2` (`:575-639`), `ComputeNumKvSplits` <- `_compute_num_kv_splits` (`triton_mla.py:40-47`), split workspace <- `_reserve_attn_logits_workspace` (`:57-78`) realized as the house grow-only per-stream scratch. **Honest reuse verdict:** our FA-2 split+combine machinery fit at the ALGORITHM level (the split schedule, the LSE merge algebra, the fixed-ascending no-atomicAdd determinism rule) and NOT at the code level — the vendored FA-2 launcher takes separate 4-D k/v caches and is instantiated for symmetric head_dim {128,256}, which cannot express a 3-D single-buffer cache with QK 576 / V 512; that is recorded in the TU header rather than forced. **Evidence is unit-level and deliberately strong** (there is no e2e model until W7): [`tests/vt/test_ops_mla_attn.cpp`](../tests/vt/test_ops_mla_attn.cpp), a port of `tests/kernels/attention/test_mla_decode_cpu.py` whose `ref_mla` becomes an INDEPENDENT TWO-PASS oracle (a different algorithm from the streaming online-softmax both impls use) plus its NaN-padding out-of-bounds detector, run at the REAL V2-Lite geometry (576/512/64, block 16, mscale^2 scale) over ragged / multi-block / single-block / every split boundary (`num_kv_splits` ∈ {1..512} incl. splits > seq_len) / 128-head V3 / non-BLOCK_H head counts / a 288-256 block-32 geometry / bf16 + f32, with run-to-run BIT-exactness. dgx sm_121: 11/11 cases, 2,303,193 assertions; `compute-sanitizer` memcheck **0 errors**, racecheck **0 hazards**, synccheck **0 errors**; clean CUDA build **0 warnings / 0 errors**; **27B 235/235 + 35B 315/315 + Coder 6/6 + Qwen3-dense 16/16 + OPT 6/6 ALL UNCHANGED**. `TritonMLABackend::get_impl_cls()` is no longer `nullptr` — it returns a real `TritonMLAImpl` whose `forward_mqa` is the 1:1 counterpart of `triton_mla.py:189-260`; PREFILL is W5 and `forward()` refuses a prefill-shaped batch BY NAME rather than producing wrong numbers. NO speed number — decode perf is W9. **W5 LANDED 2026-07-22** (base `5395203`; rows STAY `SPIKE`). **MLA PREFILL + the workspace-bounded CHUNKED-CONTEXT loop.** Three new ops — `vt::MlaPrefillAttention` (<- `mla/prefill/flash_attn.py:153-248`, the ONLY MLA prefill backend reachable on sm_121 and the one W0 OBSERVED the oracle logging), `vt::GatherMlaCache` (<- `cache_kernels.cu:992-1064`) and `vt::MergeAttnStates` (<- `merge_attn_states.cu:18-192`, both `-inf` edge cases verbatim) — plus the loop itself in the new `mla_chunked_context.h` (<- `mla_attention.py:1422-1451,1667-1745,2094-2199,2344-2425`), which is what keeps a long-context prefill inside a bounded workspace instead of materializing a 3 GB up-projected context. **The vendored FA-2 launcher WAS generalized, and W4's prediction that it would be tractable held for a reason worth recording: upstream does not ask FA-2 for asymmetric head dims either.** `requires_v_padding` is TRUE on GB10, so upstream ZERO-PADS V from 128 to 192 and slices the output back — the kernel stays a plain SYMMETRIC head_dim-192 instantiation. The whole change is two new explicit instantiations of an UNCHANGED generic template, one new launcher entry for the contiguous-varlen mode, and the pad/slice pair; the paged launcher every non-MLA prefill calls is textually untouched (211 insertions / **0 deletions**), and 27B 235/235 + 35B 315/315 + Coder 138/138 + Qwen3-dense 664/664 + OPT 36/36 are all UNCHANGED. Evidence is UNIT-ONLY (there is still no model): 4/4 cases / **2,377,052 assertions** and 5/5 / **306,037 assertions** on dgx sm_121 at the real QK 192 / V 128 geometry, against an INDEPENDENT double-precision two-pass oracle and — for the loop — a SINGLE-SHOT whole-sequence oracle that never chunks, over exact / +1 / -1 chunk boundaries, zero-context and zero-key-in-chunk requests, ADVERSARIAL reverse-interleaved block tables, NaN-poisoned outputs and run-to-run bit-exactness; memcheck **0**, racecheck **0 hazards**, synccheck **0**. A genuine upstream FA-2 quirk was found on the way and worked around rather than papered over: its EMPTY-K early exit ignores the unpadded-LSE flag, which a zero-key chunk request would turn into an out-of-bounds LSE write. **W6 LANDED 2026-07-22** (base `2846467`; rows STAY `SPIKE` — W6 adds an attention LAYER, not a model). **The MLA attention BLOCK + LOAD-TIME WEIGHT ABSORPTION — the piece that finally COMPOSES W3's cache write, W4's MQA decode and W5's MHA prefill into one layer:** the projections with BOTH `q_lora_rank` branches (`fused_qkv_a_proj` -> `q_a_layernorm` -> `q_b_proj`, or the direct `q_proj`), the two RMSNorms (the rope part deliberately NOT normed), the DECOUPLED RoPE (`is_neox_style=False`, only the trailing 64-dim slice rotates) with its YaRN cos/sin cache and the SEPARATE mscale^2 softmax-scale correction, the `kv_b_proj -> W_UK/W_UV` split, the prefill-MHA / decode-MQA dispatch with decode tokens packed FIRST, and the `kv_b_proj` up-projection callback W5 left open. **The spike's most useful prediction held: absorption needed NO new attention kernel** — it is a LOAD-TIME weight transform plus TWO batched GEMMs, so the entire new-kernel surface is two general primitives, `vt::BatchedMatmul` (<- `torch.bmm` at `mla_attention.py:789,1034`; on CUDA torch resolves that to cuBLAS `gemmStridedBatchedEx`, and ours is the cuBLASLt strided-batched form of the same GEMM) and `vt::ConcatMlaNopeRope` (<- `concat_mla_q`, generalized so one op also serves `_concat_k_nope_k_pe`). **The absorbed-vs-unabsorbed equivalence — the heart of W6 — is PROVEN NUMERICALLY, three independent ways, rather than argued:** an INDEPENDENT double-precision block oracle computes the attention BOTH ways and agrees to < 1e-11 (the identity itself); our absorbed decode reproduces the UNABSORBED oracle to < 2e-4 in f32; and — the strongest — the SAME batch is driven once through our ABSORBED MQA decode kernel (QK 576 / V 512, one KV head, K/V never materialized) and once through our UNABSORBED materialized-MHA prefill path (QK 192 / V 128 plus the chunked-context loop), agreeing to < 3e-4 (CPU f32) / < 4e-2 (CUDA bf16) with nothing but the weights shared between them. Evidence on dgx sm_121: `test_mla_attention_block.cpp` 10/10 cases / 2,372,644 assertions and `test_ops_mla_absorb.cpp` 9/9 / 1,644,807 (CUDA cases proven to EXECUTE; NaN-poisoned outputs; run-to-run BIT-exact), porting `tests/kernels/test_concat_mla_q.py` in both arms. memcheck 0, racecheck 0 hazards, synccheck 0 (the last needing `--num-cuda-barriers 65536` — the default table overflows on a binary driving this many kernel families and the tool then reports a bogus launch failure, a trap worth knowing). Clean CUDA build 0 warn / 0 err; **27B 235/235 + 35B 315/315 + Coder 138/138 + Qwen3-dense 664/664 + OPT 36/36 ALL UNCHANGED**. **Coverage stated plainly: the `q_lora` query branch has NO e2e coverage and cannot get any on GB10** — DeepSeek-V2-Lite has `q_lora_rank=null`, so it is unit-gated at DeepSeek-V3's real dimensions only; GLM-4.7-Flash (`q_lora_rank=768`, 58.2 GiB, fits) is what would close it. **W7 LANDED 2026-07-22** (base `ce43c51`; the row STILL stays `SPIKE`). **THE DEEPSEEK-V2 MODEL — registry + config parse + loader + forward: the first MLA model in the tree, and the first one that runs a real MLA checkpoint end to end.** Four new files plus ONE shared-code edit (a two-line additive condition in `runner.cpp` recognising a `kMlaAttention` KV group as the model's attention group — upstream maps MLA onto the ordinary `FullAttentionManager`, so block tables/prefix caching/eviction are untouched). **LOADER GATE PASSED on the real 4-shard DeepSeek-V2-Lite: 5291/5291 checkpoint tensors accounted for, ZERO unmapped and ZERO leftover** (4/4 cases / 37,331 assertions), every shape asserted including the LOAD-TIME `kv_b_proj -> W_UK_T [16,128,512]` / `W_UV [16,512,128]` absorption split — the same transform, at the same lifecycle point, as upstream's `process_weights_after_loading`. **V2-Lite takes the DIRECT `q_proj` query branch** (`q_lora_rank: null`), asserted with the fused branch EMPTY on every layer. **FORWARD GATE PASSED and obviously right, not merely finite: the real checkpoint prefill of `The capital of France is` -> argmax ` Paris`** (top-5 ` Paris`/` the`/` a`/` one`/` also`, run-to-run bit-exact) — the direct analogue of the Qwen3-Coder W3 sanity case. **BATCH-ORDERING GATE:** the ordering invariant W6 measured 0.86 relative error from is now VALIDATED, not assumed — `BuildMlaBatchSplit` throws (naming the request and citing the upstream line) if a decode follows a prefill or a with-context prefill follows a context-free one. **SHARED EXPERTS — new for this family and UNGATED unlike Qwen3.6's sigmoid-gated one — gated two ways:** a MoE layer with every routed expert zeroed is BIT-IDENTICAL to a dense layer holding the same MLP, and turning the shared expert off CHANGES the logits. **The CUDA path is EXERCISED, not merely compiled:** a case at the real MLA head geometry drives the CUDA MLA kernels and the CUDA-only grouped bf16 MoE GEMM, bit-exact on device and within 0.0061 worst relative logit error of the CPU reference path. 11/11 forward cases; memcheck/racecheck/synccheck all **0**; clean CUDA build **0 warn / 0 err**; **regression set UNCHANGED**. **Only `DeepseekV2ForCausalLM` is REGISTERED** — `DeepseekForCausalLM` (plain MHA), V3 (fp8/671B) and V3.2 (DSA indexer) are REFUSED BY NAME in the config parse rather than falsely claimed. A pre-existing tree-wide hazard was found on the way and recorded: the shared `DevicePool` is a process-wide singleton keyed only on a byte size class, so a single process driving BOTH a CPU and a CUDA forward hands the second backend the first's recycled pointers. **NEXT: W8 — the SACRED token-exact gate on DeepSeek-V2-Lite** (wire the paged engine to produce the MLA batch order the model already validates, capture oracle goldens, run the STRICT form W0 determined). A loading, forwarding model is NOT a supported model, so no model row moves until that gate passes. **W8 LANDED 2026-07-22 — THE SACRED CORRECTNESS GATE PASSES 8/8, and `MODEL-TEXT-deepseek-v2-deepseek-v2-for-causal-lm` moves `SPIKE` -> `ACTIVE` (correctness COMPLETE, speed PENDING). NOT `DONE` — that additionally requires vLLM-speed parity on every axis, which is W9 and has NO number yet; the other four campaign rows stay `SPIKE`.** An 8-prompt battery is driven through the FULL paged `LLMEngine` and compared to the pinned vLLM 0.25.0 oracle: **8/8 PASS — STRICT token-exact 5/8, near-tie band 3/8, 92/128 tokens strictly exact, max teacher-forced gap 0.25 nats, 0 forward-divergent** (223 assertions). **The bar was ARRIVED AT by measurement, not chosen:** vLLM is DETERMINISTIC on this model at batch=1 (W0's K=5 8/8, re-confirmed by W8's own capture at T=16 with 0 multi-valued cells), so the STRICT form ran FIRST and came out 5/8; the ratified TEACHER-FORCING diagnostic then showed **36 divergent positions with 35 at gap EXACTLY 0.0000 nats** — vLLM's own argmax GIVEN OUR PREFIX is our token, so they are the downstream tail of one earlier flip — **exactly ONE root flip with any gap at all (prompt[3] tok 9, 0.2500 nats, inside the ratified 0.5-nat band and equal to the landed Qwen3-dense 4B gate's worst)**, and **ZERO tokens outside vLLM's top-20**, with the per-position nats COMMITTED as goldens and anything beyond the band still FAILING. **W8's first job — the scheduler/runner wiring — turned out to need NO new code, for a non-accidental reason:** `runner.cpp:671` already reorders with `decode_threshold = 1`, exactly MLA's `reorder_batch_threshold` (`mla_attention.py:1420`), and its `decode -> short_extend -> long_extend -> pure_prefill` ordering satisfies BOTH MLA invariants (decodes form a batch prefix; with-context prefills lead the prefill tail). W8 PROVES that end to end rather than duplicating it, with new DIAGNOSTIC `MlaBatchSplitStats` counters and a non-vacuity bar: the battery is admitted CONCURRENTLY with staggered arrival, producing **7 genuinely MIXED decode+prefill steps at up to 8 concurrent requests** with `BuildMlaBatchSplit` (which throws naming the request) never firing, plus a prefix-cache-driven **with-context prefill**, and a phase-0 check that the engine really allocated the compact MLA cache (`fa_page_size_bytes = 36864`, no factor 2). **THE REAL BLOCKER WAS THE TOKENIZER, NOT THE MODEL:** the first run REFUSED to load (`unsupported normalizer "Sequence"`), and behind it sat a whole NEW pre-tokenizer family — DeepSeek's is a HF `Sequence` PIPELINE of SEVEN stages (five `Split(Isolated)` over ENUMERATED codepoint ranges, then `Digits(individual_digits=true)`, then `ByteLevel(use_regex=false)`), whose stage ORDER is load-bearing because stage 2's punctuation class spans 0x3A-0x7E and CONTAINS A-Z/a-z. Landed as `SplitPattern::kDeepSeek` with the five patterns compared VERBATIM at load, and MEASURED token-for-token against the REAL HF `tokenizers` library over a stage-stress corpus (**6/6 cases / 2461 assertions**). **The TOKENIZATION goldens earned their keep by REFUTING a fix that was already written:** `tokenizer_config.json` declares `add_bos_token: true`, which reads as exactly the OPT missing-BOS bug — but vLLM's resolved tokenizer (`TokenizersBackend`) adds NO BOS, our loader already matched bit-for-bit, and the "fix" would have BROKEN a passing gate; it was reverted and the measured behaviour PINNED by a guard case ([[ground-premises-before-dispatching]]). Regression set UNCHANGED (27B 235/235, 35B 315/315, Coder 6/6, Qwen3-dense 16/16, OPT 6/6, plus every tokenizer test — W8 touches SHARED tokenizer code, so that was proved, not assumed); clean CUDA rebuild 0 warn/0 err; local CPU suite 151/151; memcheck/racecheck/synccheck 0. Batch invariance is REPORTED (6/8) and deliberately NOT a bar, because the ORACLE itself changed on 3/8 of this battery under batched generation (W0). One W9 input recorded: the oracle must run `moe_backend='triton'` — vLLM's auto-selected FlashInfer CUTLASS unquantized MoE REBOOTED dgx three times on GB10's unified memory. **W9 SPEED CLOSE LANDED 2026-07-22 — the track has its FIRST binding speed number, and it is an ATTRIBUTED MISS: `MODEL-TEXT-deepseek-v2-deepseek-v2-for-causal-lm` STAYS `ACTIVE` (correctness COMPLETE, speed SHORT), NOT `DONE`** ([grid](../docs/BENCHMARKS.md), [spike §W9](specs/mla-deepseek-campaign.md)). Denominator SETTLED with evidence — CUTLASS MoE has now rebooted dgx **five times** (two more at W9, the second on a pristine box with a 0 GiB page cache and every mitigation applied, both deaths at the identical post-`torch.compile` phase), so `--moe-backend triton` IS vLLM's best STABLE GRAPHED configuration here and is the bar; the substitution does not flatter us, we lose to it. `nsys` (both sides, `--cuda-graph-trace=node`) overrode the plan: the lever was not the planned MLA fusion recipes but `MlaDecodeStage1` sitting at **44.7% of all GPU time and ~180x off its own memory-bound floor** on a **2-CTA grid at batch 1**; applying upstream's own occupancy target made it **18.3x faster** (837 -> 45.8 us) for **+69.5%/+53.3%/+32.0%/+19.5%** end-to-end at c1/c2/c4/c8, while the planned decode-graph sibling is worth only ~+2% (this decode is GPU-bound). Grid vs vLLM: output throughput **0.87/0.95/0.86/0.88** (was 0.50 at c1), TTFT **1.06/1.14/0.96/0.88** (we WIN at c4/c8), TPOT **1.11/0.97/1.16/1.17**. SACRED gate **8/8 UNCHANGED** with both levers default-ON; a real latent CUDA-graph use-after-free in the MLA metadata upload was found and fixed (its whole class now guarded); regression set UNCHANGED; clean rebuild 0 warn/0 err; sanitizers 0. **NEXT LEVER, NAMED: route the batch-1 dense projections off cuBLAS `gemvx` (31.8% of our GPU time) onto a tensor-core GEMM — vLLM splits the same work `gemvx` 12.7% + `nvjet_sm121_tst_mma_*` 6.6%.** **W10 BLOCKED-ROW HONESTY PASS LANDED 2026-07-22 — the campaign's W-plan is COMPLETE; records only (no code, no build, no GPU work, nothing downloaded, no number claimed).** Rows set to their final honest state: `MODEL-TEXT-deepseek-v2-deepseek-v3-for-causal-lm` (V3 + V3.2, and Kimi-K2/K2.5's text backbone by config composition) and `MODEL-TEXT-minimax-m2-mini-max-m2-for-causal-lm` move `SPIKE` -> `BLOCKED`, joined cross-claim by `MODEL-TEXT-deepseek-v2-glm-moe-dsa-for-causal-lm` (GLM-5) under `CLAIM-GLM-DSA-LATEST-DEEPSEEK`; each is HW-BLOCKED on 119 GiB (~642 GiB fp8 / ~428 GiB / 1404 GiB) and the two DSA models are additionally DEP-BLOCKED — for a SPARSE model the XOR filter eliminates `TRITON_MLA`, leaving `FLASHINFER_MLA_SPARSE_SM120` alone, whose sm12x dispatch goes to flashinfer's DENSE-ONLY XQA backend that discards `sparse_mla_top_k` (upstream's own test monkeypatches the probe and asserts nothing numerical). `MODEL-TEXT-deepseek-v2-deepseek-for-causal-lm` stays `SPIKE` with the record repaired to say it is plain MHA and needs NO MLA; Kimi-Linear stays `SPIKE` (MLA half unlocked, KDA a separate kernel campaign, HW-MARGINAL). Each blocked row states what CAN still be gated (config resolution, weight-map on a slice, unit parity at real dimensions) versus what CANNOT (anything e2e). **Two PERMANENT coverage gaps now stated in the rows:** the `noaux_tc` grouped router and the `q_lora` query branch have NO e2e coverage and are unit-gated only, because V2-Lite is `n_group=topk_group=1`/softmax with no `e_score_correction_bias` and `q_lora_rank=null`. **NAMED NEXT VEHICLE: GLM-4.7-Flash** (`MODEL-TEXT-glm4-moe-lite-glm4-moe-lite-for-causal-lm`, 31.2B / 58.2 GiB, FITS GB10) — the only reachable checkpoint that closes BOTH gaps. **BLOCK NOT CLOSEABLE, nothing archived:** the DeepSeek-V2 row is `ACTIVE`, not `DONE`, so the plan/spec stay LIVE; the single open item is the named `gemvx` -> tensor-core dispatch lever. **MISTRAL FIFTH FAMILY W0-W3 LANDED 2026-07-23** ([sweep-mistral](specs/sweep-mistral.md), `MODEL-TEXT-mistral-mistral-for-causal-lm` `ACTIVE`, `CLAIM-MODEL-MISTRAL`): the closest-to-Llama dense arch (vLLM `mistral.py` = "Mistral adaptation of the LLaMA architecture") — plain rope θ1e6 (no rope_scaling) + qk-norm-optional + untied lm_head + null sliding_window, all PRE-EXISTING ⇒ NO new primitive, 3 new files + additive CMake/registry-test rows only, ZERO shared-code edit. **MODEL forward gate 30/30 greedy tokens vs vLLM 0.25.0** (tokenizer-free: fed vLLM's exact prompt tokens through our CUDA prefill; 29 STRICT token-exact + 1 near-tie, 0 forward-divergent; vLLM greedy det 4/5 K=3). W2 loader real-weights 1541 assertions. **REAL FINDING:** Mistral's SentencePiece/Metaspace tokenizer is unsupported by our ByteLevel-BPE tokenizer → the FULL paged-engine SACRED gate is BLOCKED, the pre-inventoried `LOAD-SENTENCEPIECE` row (SentencePiece tokenizer family). `-Werror` 0-warn, DSR 32, regressions UNCHANGED (Llama paged 16/16, Qwen3-dense forward 1031, registry 299; MoE/GDN gates unaffected by construction). SPEED + full paged gate both PENDING (row `ACTIVE`, not `DONE`). **OLMo-2 SPIKED 2026-07-24** ([sweep-olmo2](specs/sweep-olmo2.md), `CLAIM-SWEEP-OLMO2`): one row `INVENTORIED` → `SPIKE` (`MODEL-TEXT-olmo2-olmo2-for-causal-lm`, covering `Olmo2ForCausalLM` + its `Olmo3ForCausalLM` alias). Answers the breadth-sweep §B.3 Tier-2 rank-8 "GLM4 / Olmo2-3" item (GLM-4 + Gemma landed; OLMo-2 next). **HEADLINE: OLMo-2 is the cleanest dense bring-up yet — ZERO new compute kernels.** The two distinctive facts both reduce to WIRING over landed ops: (1) the **pure post-norm (`norm_after`) placement** is a strict SUBSET of the GLM-4/Gemma sandwich (keeps ONLY the standalone-output-norm op `glm4.cpp:174-178` — the exact primitive flagged — DROPS the pre-norms, plain residual add); (2) the **QK-norm is FULL-WIDTH not per-head** → reuses `vt::RmsNorm` at a `[T,q_size]`/`[T,kv_size]` shape but CANNOT use the fused per-head `kAttnQkNormRopeGate`. Everything else REUSES (plain RMSNorm, SiLU SwiGLU, NeoX rope, GQA paged glue, Gemma-3 sliding-window for Olmo-3, tied embeddings, packed loader, ByteLevel BPE). Gate vehicle `allenai/OLMo-2-0425-1B` (1.485B, ~2.77 GiB, fits GB10 tight ~30 GiB free); Olmo-3 rides the same row (0.25.0 oracle-support UNVERIFIED). OLMo-1 (non-parametric LayerNorm), OLMoE/FlexOlmo (MoE), OlmoHybrid (SSM) stay `INVENTORIED`. Nothing implemented. | -| 2a | `ROAD-V1-C2-LOCAL-BF16` | Close local discrete-Blackwell Qwen3.5 plain-BF16 production parity | [`MODEL-MM-qwen3-5-qwen3-5-for-conditional-generation`](model-matrix.md), [`LOAD-SAFETENSORS-DIRECT-DENSE`](engine-matrix.md), [`SERVE-CLI-BENCH`](engine-matrix.md), [`KERNEL-SSM-MAMBA`](kernel-matrix.md) | Exact `(sequence, 8-token chunk)` GDN conv dispatch is byte-identical and default ON. Rebased-main same-binary reprofile: **718.704→233.955 ms (3.072x)** and whole run **+2.272%**; sealed vLLM remains 145.421 ms (**1.609x residual**). The reviewed combined default-off K4 causal-conv + 16-token post-conv local A/B is exact and improves total throughput **0.7175%**, TTFT **1.340%**, TPOT **0.556%** and E2E **0.716%**. Prior cross-engine throughput and TTFT ratios are **VOID/PENDING** because only local timed string tokenization; TPOT and VRAM remain OPEN | `GATING` | Implement the committed pretokenized benchmark frontend with `VT_BENCH_PRETOKENIZE=0` rollback, review, then rerun local/vLLM. Latency, VRAM and 27B/35B correctness remain gates; no 4B-to-gate-model extrapolation. Evidence: [campaign spec](specs/sm120-qwen35-pareto-2026-08-09.md), [production baseline](../docs/bench-evidence/qwen35-4b-sm120-main-20260807.md) | +| 2a | `ROAD-V1-C2-LOCAL-BF16` | Close local discrete-Blackwell Qwen3.5 plain-BF16 production parity | [`MODEL-MM-qwen3-5-qwen3-5-for-conditional-generation`](model-matrix.md), [`LOAD-SAFETENSORS-DIRECT-DENSE`](engine-matrix.md), [`SERVE-CLI-BENCH`](engine-matrix.md), [`KERNEL-SSM-MAMBA`](kernel-matrix.md) | Exact `(sequence, 8-token chunk)` GDN conv dispatch is byte-identical and default ON. Rebased-main same-binary reprofile: **718.704→233.955 ms (3.072x)** and whole run **+2.272%**; sealed vLLM remains 145.421 ms (**1.609x residual**). The reviewed combined default-off K4 causal-conv + 16-token post-conv local A/B is exact and improves total throughput **0.7175%**, TTFT **1.340%**, TPOT **0.556%** and E2E **0.716%**. The #206 frontend now defaults to pre-timed tokenization with exact timed-string rollback and is CPU-GREEN with identical prompt/output IDs. Cross-engine throughput and TTFT remain **VOID/PENDING** until the reviewed GPU rerun; TPOT and VRAM remain OPEN | `GATING` | Fresh mutation review, operator CPU gate, then counterbalanced `STRING→TOKENS→TOKENS→STRING` GPU A/B and pinned-vLLM rerun. Latency, VRAM and 27B/35B correctness remain gates; no 4B-to-gate-model extrapolation. Evidence: [campaign spec](specs/sm120-qwen35-pareto-2026-08-09.md), [production baseline](../docs/bench-evidence/qwen35-4b-sm120-main-20260807.md) | | 3 | `ROAD-V1-C3` | MTP k=1 + GDN speculative path, then DFlash, DSpark and heterogeneous-vocabulary TLI | [engine matrix](engine-matrix.md), [coverage view §8](feature-matrix.md#8-speculative-decoding) | MTP and DFlash specs exist. **M-mtp-0 CLOSED 2026-07-24** - the standalone MTP draft head is oracle-parity-proven on BOTH gate checkpoints (op-level parity vs a dumped k=1 vLLM oracle, not a token-generation SACRED gate). **I2 SCHEDULER-HALF LANDED 2026-07-24** ([mtp-spec-decode §2.7](specs/mtp-spec-decode.md)): host-side spec-decode scheduler/engine plumbing + the FROZEN spec-metadata ABI that I3 (rejection sampler) and I5 (verify/propose runner) build against - `SpeculativeConfig`, `DraftTokenIds`, `Request::spec_token_ids`/`NumTokensWithSpec`, the first population of `scheduled_spec_decode_tokens`, `Scheduler::update_draft_token_ids`, the `take_draft_token_ids` seam, `EngineCore::post_step`, `InputBatch::num_accepted_tokens`/`update_req_spec_token_ids`; DEFAULT-OFF and INERT (no `SpeculativeConfig` => `num_lookahead_tokens == 0` => byte-identical engine). `SPEC-MTP` **STAYS `GATING`** because M-mtp-1..4 (greedy rejection, GDN spec slots, k>1, CUDA graphs) are still open, so spec decode remains user-invisible. DSpark is user-promoted scope with DeepSeek-V4/Qwen3 draft models, reduced-vocabulary handling and full-CUDA-graph behavior inventoried under `SPEC-DSPARK`; tokenizer-agnostic target<->draft mapping is separately inventoried as `SPEC-TLI`. Their dedicated spikes are not written **I3 GREEDY REJECTION SAMPLER LANDED 2026-07-24** (`SPEC-REJECTION` `READY` -> `ACTIVE`): per-request logits expansion to `1 + k_i` rows plus the greedy accept rule (accept a draft iff it equals the target argmax; on the first mismatch emit the target argmax and stop; bonus token when all k accept), CUDA==CPU bit-exact at vocab 248320. **I4 GDN-HALF LANDED 2026-07-24** (`SPEC-GDN-SEGMENTS` -> `ACTIVE`): the GDN spec metadata split + decode->prefill reclassification, the `T>1`/`IS_SPEC` recurrence with per-timestep snapshots, the conv sliding window honouring `num_accepted`, k+1 slot allocation - bit-exact rollback. **I5a GDN LAYER ROUTING + runner spec-metadata upload LANDED 2026-07-24** (`CLAIM-SPEC-MTP-I5A`): `GdnBlockPaged` now routes a pure-spec batch through `vt::GdnSpecDecode`/`vt::CausalConv1dSpecUpdate` and the runner uploads I4's six spec device tensors - first of the scoped M-mtp-1 sub-increments (I5a GDN wiring -> I5b prepare_prefill -> I5c MTP paged propose -> I5d config+runner-loop+27B token gate, spec §5), DEFAULT-OFF INERT, bit-exact vs the I4 ops, no e2e loop yet. **I5b `prepare_prefill_inputs` LANDED 2026-07-24** (`CLAIM-SPEC-MTP-I5B`): the drafter prefill input-prep - shift each request's `input_ids` left one within its query span, splice the just-sampled next token, `query_len -= num_rejected`, emit last-token index / query_start_loc / seq_lens into the `SpecPrefillInputs` struct; a HOST routine (no new CUDA kernel; mirrors our DEVICE-NEUTRAL `prepare_inputs`/`combine_sampled_and_draft_tokens` family), unit-gated 7 cases / 27 assertions RED-first, DEFAULT-OFF INERT, additive by construction. **I5c MTP PAGED PROPOSE + DRAFT KV LAYER LANDED 2026-07-24** (`CLAIM-SPEC-MTP-I5C`): `Qwen3_5MTPModel::ForwardPaged` runs the head + one full_attention decoder layer over the head's OWN paged draft KV layer (ReshapeAndCache + PagedAttention over the target's block table / slot mapping); `MakeQwen3_5KVCacheSpec(num_spec>0)` adds that draft KV layer (`fa_draft` FullAttentionSpec group, index num_hidden_layers); `ForwardDeviceTap` exposes the `[T,H]` post-final-norm hidden tap (INERT); and `MtpProposePrefill` is the callable k=1 propose (I5b shift-splice -> one paged forward -> argmax at last_token_indices, early-exit). CORE PROOF: the paged forward reproduces I1's standalone head logits/argmax on BOTH gate checkpoints; a two-step drive proves the draft-KV write/read (RED control diverges). DEFAULT-OFF INERT (no spec config -> draft KV layer not allocated, tap nullptr, target forward byte-identical); NOT wired into the runner step loop. **I5d-pre REGISTRY/FORWARD-SEAM ENABLING REFACTOR LANDED 2026-07-25** (`CLAIM-SPEC-MTP-I5D-PRE`): a scoping pass found the model seam is fully TYPE-ERASED, so the runner cannot reach the concrete target weights / hidden-state tap / loaded MTP weights the I5d loop needs. Four ADDITIVE, inert-when-spec-off access paths + one latent-bug fix - the `hidden_tap` out-field on the type-erased `ModelForwardInput` (routes to the existing `ForwardDeviceTap`), a `LoadedModel::BuildMtpDraft` virtual (typed path to the draft, null for non-MTP), MTP weight loading + shard retention in `FromModelDir` behind `EngineParams::speculative_config`, and the `GPUModelRunner` ctor widened with optional draft/draft-KV/`SpeculativeConfig`; PLUS the latent `initialize_kv_cache` fix (select the FIRST non-eagle full-attn group as the target so a third `fa_draft` group can't displace it; byte-identical at num_spec==0). DEFAULT-OFF INERT, unit-gated RED-first, spec-off SACRED gates byte-identical. **I5d CONFIG RUNTIME + VERIFY/PROPOSE RUNNER LOOP LANDED as a spec-off-byte-identical PARTIAL 2026-07-25** (`CLAIM-SPEC-MTP-I5D`): `--speculative-config` JSON parse -> `EngineParams` -> `LoadedEngine` resolution (widened KV `MakeQwen3_5KVCacheSpec(num_spec>0)`, `BuildMtpDraft`, forced sync scheduling, `MakeScheduler(spec)`, `EngineCore(check_for_draft=true)`) + the full runner loop (draft splice, hidden-tap capture, GDN builder spec-overload feed, k+1 GDN state-slot remap + widened conv cache + draft-KV alloc, `MtpProposePrefill`, `take_draft_token_ids`, acceptance telemetry). CUDA `-Werror` 0 warnings, cutlass-ON; spec-OFF SACRED byte-identical (27B 235/235, 35B 315/315, Coder 138/138 + spec unit tests ALL PASS). **The three-way 27B token gate is NOT yet passing**: the spec-ON engine RUNS the loop end to end and MEASURES the blocker (`test_qwen27_spec_decode`) - it throws on the FIRST prefill step at `gdn_state_gather: working/cache row shapes must match` (`src/vt/ops.cpp:1773`) because I4's spec conv rollback needs the conv row widened to `(K-1)+num_spec` while the non-spec GDN conv ops assume `(K-1)`. `SPEC-MTP` STAYS `GATING`. **I5e LANDED 2026-07-25 (`CLAIM-SPEC-MTP-I5E`): the non-spec GDN conv ops made widened-cache-aware (mirror vLLM `state_len=KERNEL_WIDTH-1` + physical `stride_conv_state_tok`, leading `(K-1)` sub-window, byte-identical at `num_spec==0`) AND the async input-combine forced off under spec (it overwrote the verify batch's draft position with the committed token -> 0 acceptance, RCA'd on the real 27B). **THE THREE-WAY 27B SINGLE-REQUEST GREEDY GATE PASSES**: our-spec-ON == vLLM `--speculative-config mtp` greedy == our-spec-OFF token-for-token, **acceptance 16/16 drafts accepted** (~16 target steps saved); spec-OFF SACRED byte-identical (27B 235/235, 35B 315/315, Coder 138/138), compute-sanitizer 0 on the spec step. `SPEC-MTP` LEAVES `GATING` (single-request greedy correctness PROVEN); NOT `DONE` - the MIXED `GdnBlockPaged` concurrency split/merge + the throughput A/B vs vLLM same-config are I6. **I6 LANDED 2026-07-25 (`CLAIM-SPEC-MTP-I6`, `benchmark_binding=true`): the §5 c1 THROUGHPUT GATE — ours spec-ON AT/ABOVE vLLM spec-ON on EVERY measured axis at c1** (TPOT 66.2/62.95 vs 69.1/65.3 ms prose/code, ours ~1.04x faster; output tput +4.6%/+3.9%; ITL/TTFT lower; acceptance ours 0.85/0.92 vs vLLM 0.838, within noise; spec helps both ~1.5-1.6x TPOT; ours ~4% faster spec-OFF too), via an additive example-only `--speculative-config` bench flag (NO engine code touched). STAYS `ACTIVE`: the c>1 mixed-batch path is still refused + owes a c>1 A/B, and no server-facing spec flag yet. | `ACTIVE` | M-mtp-0, I2 scheduler-half, I3 rejection sampler, I4 GDN spec slots, I5a GDN layer routing, I5b prepare_prefill, I5c MTP paged propose + draft KV, I5d-pre the registry/forward enabling seam, and I5d config runtime + verify/propose runner loop (spec-off byte-identical) are landed; next (before `SPEC-MTP` leaves `GATING`) is closing the measured I5d gate blocker - make the non-spec GDN conv ops widened-cache-aware (mirror vLLM `causal_conv1d` `state_len=width-1+(seqlen-1)`) + the MIXED `GdnBlockPaged` split/merge - then the passing M-mtp-1 27B k=1 greedy three-way token gate + acceptance, then M-mtp-2 35B, then DFlash, the DSpark spike/gates and TLI. **DFlash D0-redo + D1 LANDED 2026-07-26 (`CLAIM-DFLASH-D0D1`, [dflash-spec-decode §0](specs/dflash-spec-decode.md)): `SPEC-DFLASH` UNBLOCKED + `ACTIVE` on the advanced pin `555967922`/vLLM 0.26.0.dev0.** The prior 0.25.0 ORACLE-BLOCKED verdict is SUPERSEDED — under `VLLM_USE_V2_MODEL_RUNNER=1` (vllm#40898 resolved) the mixed-SWA/full z-lab 27B draft CONSTRUCTS and the drafter is ALIVE (acceptance 2.21/8.80/4.75/4.57 > 1, `num_spec=16`, flashinfer-native fp8-KV; goldens committed). Gate FORM measured STRICT MODE-MATCHED (vLLM-ON run-deterministic K>=3 but != vLLM-OFF at k=16 near-ties — NOT the MTP three-way identity). **D1 `DF-AUX-TAPS` DONE:** the single hidden tap is generalized to the multi-tap `[T,H×taps]` (`ForwardDeviceMultiTap` capturing `(hidden+res)` at `target_layer_ids`), config-gated byte-identical off; unit gate 598 assertions (RED-first), CUDA 697/697 + sanitizer 0, 27B MTP e2e 9/9 + 27B SACRED 235/235 byte-identical (inertness). **D2-D5 LANDED 2026-07-26 (`CLAIM-DFLASH-D2`/`D3`/`D4D5`/`D5`):** the drafter model + the project's first non-causal in-block attention (D2, GPU parity vs the real vLLM draft), context-KV precompute + `prepare_dflash_inputs` (D3, GPU numeric-parity 61/61), the non-autoregressive whole-block propose brick + `dflash` config-select (D4), and the RUNNER-LOOP INTEGRATION + 27B e2e (D5): the full verify/propose loop is wired (separate z-lab draft load + target-shared bf16 embed/lm_head, aux-tap capture, per-request combined-feature context accumulation honoring num_rejected, `propose_drafts_dflash`) and RUNS end to end - `test_qwen27_dflash_spec_decode` 2/4 STRICT token-exact vs the vLLM-DFlash-ON golden + acceptance ~ vLLM on ALL 4 (19/39/29/25 vs 17/39/30/25). The 2 divergences are SINGLE bf16 near-tie flips (ratified near-tie ROOT = the D3 inline context-KV recompute envelope), NOT a wiring bug; inertness SACRED 235/235 + MTP 9/9 byte-identical; CUDA `-Werror` clean, no new kernel. NOT a clean strict-4/4 pass - STRICT 4/4 token-identity + the speed A/B = D6 (persistent paged draft-KV bit-matching vLLM's fused projections + the uniform-1+k FULL CG). Capture tool + goldens: `scripts/spec/d{0,2,3}_dflash_*.py`, `tests/parity/goldens/dflash_27b{,_draft,_kvprep}/`. **D6-D9 SPEED CAMPAIGN 2026-07-27:** D6 c1 A/B + STRICT-4/4 bf16-irreducibility RCA; D7 device-resident within-step forward (bit-identical); D8 acceptance RCA + FINAL golden A/B (ours 0.69× vLLM). **D9 (`CLAIM-DFLASH-D9`) PERSISTENT PAGED DRAFT-KV LANDED (bit-identical, +22.7%): `AppendContextKVHost` + `ForwardBlockLogitsWithPrecomputedKV` replace the O(context²) per-step recompute with an append-only per-request store; ours-ON 20.99→25.75 tok/s = 0.917× vLLM-ON (28.09, was 0.69×); e2e 27/27 SAME tokens, SACRED 235/235 + MTP 9/9 byte-identical, CUDA `-Werror` clean, no new kernel. D8's "bf16 acceptance ceiling" REFUTED — same-trajectory per-step acceptance == vLLM (ratio 1.00) and ours realized acceptance (3.68/step) > vLLM (3.31); the SOLE residual (~8%) is the FULL uniform-(1+k) CUDA graph (eager-vs-graphed), a closeable increment. SPEC-DFLASH stays `ACTIVE` (speed not yet ≥ vLLM).** **SPEC-MTP → `DONE` 2026-07-26 (`CLAIM-SPEC-MTP-DONE`, records-only, closing commit I7 `72f9fb1`):** the user ratified the c>1 near-tie+SPEED criterion, closing both I6-owed items (mixed-batch concurrency + server/CLI/C-ABI `--speculative-config`); MTP k=1 is COMPLETE + gated. **M-mtp-2 CLOSED 2026-07-26 (`CLAIM-SPEC-MTP-M-MTP-2`): the 35B `Qwen3_5MoeMTP` full e2e three-way token gate PASSES** — our spec-ON == our spec-OFF == vLLM 0.25.0 `--speculative-config mtp` greedy == vLLM spec-OFF, 16/16 vs the `greedy_ids` anchor (STRICT, c1), acceptance 16/16 both sides; c1 spec-ON 1.19x TPOT / +16.3% output-tput vs spec-OFF (acceptance 0.908) — MoE speedup transfers; spec-OFF byte-identical (test+docs-only). MTP is now `DONE` on BOTH gate models (`MODEL-SPEC-qwen3-5-mtp-qwen3-5-moe-mtp` `GATING`→`DONE`). **DFlash D11+D12 2026-07-27 — the FULL uniform-(1+k) CUDA graph is being built in three parts:** D11 (`CLAIM-DFLASH-D11`) landed Part A (the device-store primitive, CPU-gated); **D12 (`CLAIM-DFLASH-D12`) landed A-wire (the D11 device store is now the PRODUCTION path; GPU-gated e2e 27/27 all-exact acceptance 19/39/29/25 + SACRED 235/235 + MTP 9/9 byte-identical) + Part B (`vt::DFlashPagedBlockAttention`, the capture-safe paged kernel; `test_ops_dflash_paged_block_attn` 795648/795648 CPU==CUDA + cross-check vs materialized `DFlashBlockAttention` + compute-sanitizer 0; NO function-local host `cu_seqlens` upload = capture-UAF fixed).** Speed UNCHANGED 0.917× (A-wire eager + Part B not yet wired). The SOLE remaining piece is Part C (static-shape capture + device mask-scatter + `BeginCapture`/replay + the ≥vLLM c1 A/B); if ours-ON-graphed ≥ vLLM-ON → SPEC-DFLASH DONE → C3 complete. C3 stays `ACTIVE` (DFlash Part C + DSpark/TLI remain) | | 4 | `ROAD-V1-C4` | Quantization: llama.cpp breadth/speed, NVFP4/FP8/MX, MLX native | [quantization matrix](quantization-matrix.md) | coverage spike merged; `QUANT-GGUF-CPU-THREADPOOL` W1-W3 implemented and correctness-gated, still `GATING` (its reproduction now exists — same-binary 1-vs-20-thread A/B is prefill 12.47x / decode 8.05x / RSS 1.000x, so **decode misses the >=10x bar**). **GGUF COMPUTE-IN-QUANT IS NOW LIVE AND DEFAULT-ON (2026-07-22, `CLAIM-QUANT-GGUF-CIQ-G4-1`):** [compute-in-quant GEMM](specs/gguf-compute-in-quant-gemm.md) **G1-G4** — block dtypes + traits, the Q8_0/Q8_K activation quantizers, the six generic `vec_dot`, `kMatmulBTQuant`, and now the ROUTING (`vt::MatmulBT` dispatches a block-dtype weight to it) — plus [keep-quant loader](specs/gguf-keep-quant-loader.md) **L1-L4**, whose master switch defaults ON wherever that op is registered for the running device (CPU today; a CUDA build still expands). Six encodings (Q4_0, Q8_0, Q3_K, Q4_K, Q5_K, Q6_K) now carry `C` = `Y`. **Correctness held exactly**: the 35B GGUF gate is 16/16 token-exact vs the same-file llama.cpp oracle with the quant path on, and the bench model's output tokens are byte-identical across the pre-G4, post-G4 and `VT_CPU_REF=1` arms — no golden regenerated. **Binding CPU A/B** (idle dgx aarch64, one flock, same binary, 3 reps): decode **3.45x**, prefill **4.16x**, peak RSS **1.16x less**; vs llama.cpp we went from 11.7x / 34.1x / 2.66x behind to **3.38x / 8.20x / 2.29x**. The projected 9-17x did **not** hold, for a measured reason: 60 % of that file's weight bytes are `f16`, which no block encoding covers. Keep-quant loader **L4** is therefore MEASURED-and-NOT-MET on RSS; other leaf specs open. **THAT #1 LEVER IS NOW LANDED (2026-07-22, `CLAIM-KERNEL-CPU-ELEM-GEMM-1`, new row [`KERNEL-GEMM-CPU-ELEM`](kernel-matrix.md)):** [the elementwise CPU GEMM](specs/cpu-elementwise-gemm.md) **E1-E4** — per-dtype specialization out of the K loop, 16 independent accumulators instead of one, AArch64 NEON + x86-64 SSE2/F16C tiers behind a runtime probe, and M-blocking — all **BYTE-IDENTICAL** to the historical kernel (`memcmp` gate, exhaustive 65,536-pattern widening check, same token md5), so nothing was regenerated. Binding same-binary A/B: prefill **3.41x**, decode **3.11x**; op-level bf16 18-24 -> 69-351 GFLOP/s. **vs llama.cpp: decode AT PARITY (1.03x), prefill 2.34x behind, RSS 2.29x worse. THEN loader L5 LANDED (2026-07-23, `CLAIM-QUANT-GGUF-KEEPQ-L5-1`):** [keep-quant loader](specs/gguf-keep-quant-loader.md) **L5** — mmap in-place residency (borrow kept blocks out of the read-only mapping, refcounted), tied-head sharing (one bf16 vocab matrix for embed+lm_head), and a read-once page release (port of llama.cpp `unmap_fragment`) — took **peak RSS 6.401 -> 3.884 GiB (2.29x -> 1.39x llama.cpp)** with decode UNCHANGED and output tokens byte-identical (md5 `d235db12f2cd304007530286a1755c95`). The remaining ~1.09 GiB over llama.cpp is the f16 expansion (no block encoding covers f16). | `PARTIAL` | **THE OWED FRESH PROFILE IS DONE (2026-07-23) and it re-ranks the plan.** A `vt::GetOp` hook (100% of wall time, reverted before binding) on the CURRENT binary: prefill is **no longer GEMM-bound** — kMatmulBTQuant 37%, **kGdnPrefill 25%**, kMatmul 12%, kMatmulBT 10%, **kPagedAttention 10%**; the two non-GEMM kernels (GDN linear-attention recurrence + paged attention) run **SINGLE-THREADED** on the CPU and are now the top prefill levers. Decode is memory-bound matmul at parity, no kernel work owed. **RE-RANK: G5/G6/G7 all only speed the already-fast quant GEMM and rank BELOW the two serial non-GEMM kernels; the new #1 CPU lever is threading kGdnPrefill + kPagedAttention.** **THAT #1 LEVER IS NOW LANDED (2026-07-23, `CLAIM-CPU-THREAD-GDN-PAGED-1`, [two-kernel threading](specs/cpu-thread-gdn-paged-2026-07-23.md)):** kGdnPrefill chunks over the (sequence, value-head) axis and kPagedAttention over query-token rows, both via the existing `ParallelForRows`, both **byte-identical** (qwen35 output-token md5 `d235db12f2cd304007530286a1755c95` unchanged at threads 1/4/20 + `VT_CPU_REF=1`, determinism battery extended, CPU ctest 158/158). **Binding dgx aarch64 (idle): prefill 1.382× same-binary (73.0→100.9 t/s), 2.43×→1.76× behind llama.cpp pp128; decode at parity; op-scaling 1→20 GdnPrefill 7.08× / PagedAttention 8.96×; fresh profile shows the two kernels 35%→8.6% of prefill and re-ranks the NEW bottleneck to the GEMMs (kMatmulBTQuant 50% + kMatmul 16% + kMatmulBT 14% = 80%) ⇒ next CPU lever is the SIMD/repack GEMM tiers (G5/G6/G7).** **THE FIRST SUCH TIER IS NOW LANDED (2026-07-23, `CLAIM-QUANT-GGUF-CIQ-G6-1`, [compute-in-quant GEMM](specs/gguf-compute-in-quant-gemm.md) G6):** the Arm **i8mm mmla `nrc==2`** `vec_dot` 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 (decode M=1 → portable, unchanged); runtime `HWCAP2_I8MM` probe + `VT_CPU_QUANT_MMLA` defeat + per-file `+i8mm`. **BYTE-IDENTICAL** where the math allows (q8_0/q4_0 bit-exact to the scalar tier, q4_K/q6_K within NMSE ≤ 5e-4), bit-identical across threads 1/2/4/20, e2e token md5 `d235db12f2cd304007530286a1755c95` byte-identical (mmla on/off/`VT_CPU_REF=1`), 35B GGUF gate 16/16 vs llama.cpp on both files. **Op-level portable→i8mm: q4_K 7–8.4×, q6_K 3.8–4.5×, q8_0 ~1.2×**; e2e prefill same-binary **1.084×** on the q8_0-dominant bench file (1.56×→**1.44× behind** llama.cpp pp128, Amdahl-bounded — the big k-quant win lands on the APEX 35B files). Fresh bottleneck: the elementwise f16/f32 GEMM (~30%, unchanged) is now co-dominant on this mixed file. CUDA `-Werror` 0-warn, regression set UNCHANGED. `docs/BENCHMARKS.md` ACCEPTED. RSS deficit closed to 1.39x by L5; the last RSS lever is an f16 keep-as-is compute path, not this loader. **THEN the GDN split-projection orientation LANDED (2026-07-23, `CLAIM-CPU-GDN-ORIENT-1`, [GDN projection orientation](specs/cpu-gdn-proj-orientation-2026-07-23.md)):** a fresh op-dispatch profile of the current binary (warm prefill, `vt::GetOp` hook + per-GEMM shape histogram, reverted before binding) found the four GDN input projections (`in_proj_qkv/z/b/a`, 72 GEMMs, **17.9%** of prefill: `kMatmulBTQuant 50.7% / kMatmul 17.9% / kMatmulBT 14.9%`) were the LAST weight family `LoadGdnGguf` still transposed into [K,N] (nk=false → the N-striding `kMatmul`, no M-blocking) after G4's `expand_nk` gave every other expanded weight the file's own [N,K] order. New `GgufLoadPolicy::gdn_expand_nk` + `MakeGdnProj` keep them [N,K] nk=true → the M-blocked `kMatmulBT`; **BYTE-IDENTICAL** (same sequential f32 K-reduction, only the weight offset differs — token md5 `d235db12f2cd304007530286a1755c95` unchanged across on/`VT_GGUF_GDN_NK=0`/`VT_CPU_REF=1` and threads 1/4/20), `test_qwen36_gguf_engine` 2/2·28/28·16/16 on APEX. **Binding same-binary prefill 1.090× / decode 1.09× (44.1→40.4 ms TPOT = 1.01× llama tg32, at parity), 1.44×→1.32× behind llama.cpp pp128, RSS unchanged.** Fresh post-change profile: `kMatmul` **17.9%→0% (72→0 calls, ELIMINATED)**, absorbed into `kMatmulBT` (14.9%→27.7%); **next CPU prefill lever = the quant GEMM (kMatmulBTQuant, now 55%): G7 repack-at-load.** **G7 LANDED 2026-07-23 (`CLAIM-QUANT-GGUF-CIQ-G7-1`, [compute-in-quant GEMM](specs/gguf-compute-in-quant-gemm.md) G7):** q8_0 repacked once at load into the i8mm `block_q8_0x4` interleave (ported llama.cpp `repack.cpp` `q8_0_4x8`), `kMatmulBTQuant` dispatches a pre-shuffled gemm/gemv with no per-block register shuffles. **BIT-IDENTICAL** (byte-permute weight + non-fused `vmlaq_f32`, 305-assertion memcmp across decode/prefill/bf16-out/threads, token md5 `d235db12f2cd304007530286a1755c95` unchanged on/`VT_CPU_QUANT_REPACK=0`/`VT_CPU_REF=1`; a `ResidentWeight`/`MakeTensor` flag-drop that produced all-zero tokens was caught by the E2E gate and fixed). Op-level q8_0 **3.7–5.9×** (518→2401 / 583→3456 / 514→1902 GFLOP/s); **E2E prefill 1.92× same-binary (1096→572 ms), 223.8 t/s vs llama.cpp pp128 177.3 = 1.26× — AT/BEYOND PARITY** (was ~1.5× behind), decode at parity, RSS 3.884 GiB unchanged. Fresh profile: q8_0 GEMM 55%→~21%; **the CPU prefill-lever search is CLOSED — the sole remaining gap to llama.cpp is peak RSS (1.39×), not prefill.** CUDA-inert (gated off any non-CPU-quant device), CUDA `-Werror` 0-warn, regression set UNCHANGED. **THE RSS GAP IS NOW CORRECTLY ATTRIBUTED (2026-07-23, `CLAIM-QUANT-GGUF-KEEPF16-L6-1`, [keep-quant loader](specs/gguf-keep-quant-loader.md) L6): it is NOT the f16 expansion.** L6 implemented keep-f16 residency (keep the file's 56 F16 weights + tied head resident as F16 and compute on them, mirroring llama.cpp `ggml_vec_dot_f16`) and MEASURED it **RSS-NEUTRAL** (3.884 → 3.832 GiB, −52 MB) and prefill-regressive (TTFT 577 → ~1000 ms, from 1.25× ahead of llama.cpp to 0.72× behind) — because L5's page-release had ALREADY dropped the f16 file pages, so keep-f16 only swaps an anonymous bf16 buffer for equal-size file-backed f16 pages. smaps attribution proves our weight residency is at llama.cpp parity (file-backed 2.63 ≈ 2.68 GiB); **the residual ~1.08 GiB is the engine's ANONYMOUS activation/KV workspace, not weights — the real, separate CPU RSS lever.** keep-f16 ships DEFAULT OFF (`VT_GGUF_KEEP_F16=1` opt-in), tokens byte-identical (md5 `d235db1…`), `test_gguf_keep_quant` 35/35 (x86+aarch64), regressions UNCHANGED (27B 235/235, 35B 315/315, Coder 138, dense 184, OPT 63, DeepSeek 223, Llama 92, GGUF engine 28/28). **NEXT CPU RSS lever: profile + shrink the engine's activation/KV working set, NOT the weight loader** | | 5 | `ROAD-V1-C5` | Sliding window, local attention, YaRN/long context | [engine matrix](engine-matrix.md), [coverage view §§2,11](feature-matrix.md#2-kv-cache--memory), [joint spike](specs/sliding-local-yarn-long-context.md) | **CUDA GPU CLOSURE 2026-07-27 (`CLAIM-ROADMAP-C5`, dgx GB10 sm_121a, clean build of `489f7771`, oracle vLLM 0.26.0.dev0):** the shared scaled-RoPE + local-mask CUDA path COMPILES `-Werror`-clean and RUNS on GB10; the C5 feature-positive correctness gates that were the stated `GATING` blocker now PASS — SWA (Gemma-2/Gemma-3 48/48), LongRoPE (Phi-4-mini 16/16, RED-first), llama3-rope (Llama-3.2-1B 16/16), dynamic-NTK (InternLM2 16/16); both RoPE 0.26-oracle recaptures BIT-IDENTICAL to goldens (zero drift). Leaves `ATTN-SLIDING-WINDOW`/`ATTN-ROPE-{LLAMA3,LONGROPE,DYNAMIC-NTK}`/`ATTN-YARN` → `ACTIVE` | `PARTIAL` | (RI) **Honest residual (vehicle-blocked, not skipped):** YaRN model e2e (no cached Nomic/gpt-oss consumer) + chunked-local model e2e (no Llama4 row) are REACHABLE-BLOCKED — operator/formula stay GPU/G3-gated; long-context positive-mask (prompt > W) SWA model e2e + the KV memory-optimization G8; and the roadmap-wide every-axis SPEED tail (all C5 leaves correctness-complete, speed-pending, mirroring their model consumers). Not row-DONE until speed + the blocked vehicles close | diff --git a/.agents/specs/sm120-qwen35-pareto-2026-08-09.md b/.agents/specs/sm120-qwen35-pareto-2026-08-09.md index 2807967f8..30ef6215c 100644 --- a/.agents/specs/sm120-qwen35-pareto-2026-08-09.md +++ b/.agents/specs/sm120-qwen35-pareto-2026-08-09.md @@ -7,7 +7,7 @@ VRAM gaps vs vLLM](https://github.com/mudler/vllm.cpp/issues/206) `ROAD-V1-C2-LOCAL-BF16` **Lifecycle:** combined prefill opt-in accepted locally; frontend timing parity -correction specified before the next cross-engine claim +correction implemented and CPU-gated, with fresh review and GPU A/B pending **Target:** local NVIDIA GeForce RTX 5070 Ti (`sm_120`), Qwen3.5-4B plain BF16. This vehicle does not substitute for the hardware-unavailable 27B/35B @@ -131,6 +131,17 @@ under issue #206 must satisfy this contract: Only that rerun may restore cross-engine throughput or TTFT ratios. Retain `VT_BENCH_PRETOKENIZE=0` as the exact same-binary rollback. +### Frontend implementation checkpoint (2026-08-09) + +The RED-first focused build failed because the production dispatch seam and +auditable result fields did not exist. The implementation now pre-encodes all +prompts with `EncodeWithSpecialTokens` before `t0`, moves them through the +existing token-ID overload by default, and keeps exact `=0` string admission. +The synthetic CPU gate passes 6/6 tests and 57 assertions; default and rollback +produce identical prompt and output token IDs in submission order. This is not +a performance acceptance: fresh mutation review, operator verification and the +counterbalanced production GPU A/B remain pending. + ## Correctness, tests, review and acceptance Before product code, a fresh implementer ports or restores the smallest tests diff --git a/docs/BENCHMARKS.md b/docs/BENCHMARKS.md index aed942b9c..ff72a8cf2 100644 --- a/docs/BENCHMARKS.md +++ b/docs/BENCHMARKS.md @@ -41,7 +41,7 @@ The binding comparison. vLLM runs its **production graphed config**, never | Qwen3.6-27B | NVFP4 (`nvidia` @`0893e160`, ModelOpt `modelopt_mixed`) | 0.25.0 | 0/4 | **BEHIND, uniformly 0.85x** on decode throughput (was 0.72x before the FP8 tower fix); greedy continuation IDENTICAL to vLLM. A different model from the `unsloth` row (NVFP4 MLP + FP8 W8A8 GDN/attn tower) | | Qwen3.6-35B-A3B | NVFP4 `modelopt_mixed` | 0.25.0 | 2/18 | 3-rep grid 2026-08-05 @`1ea26427`: 0.93-1.03x (c4 wins), c16 0.93x. Both c16 levers A/B'd NEG: drain event -1.9%, mirror 0.999x. ★ probe found a prod async batch-1 greedy DEGENERATION bug the mirror fixes | | DeepSeek-V2-Lite | bf16 MLA | 0.25.0 | 4/25 | Attributed miss, row stays `ACTIVE` | -| Qwen3.5-4B | bf16 direct-load | 0.26.0.dev0 | host PSS; frontend correction pending | Combined local A/B accepted: +0.718% tput, -1.340% TTFT, -0.556% TPOT; IDs exact. Cross-engine tput/TTFT **VOID/PENDING**; TPOT 1.025x and VRAM 1.018x OPEN ([data](bench-evidence/qwen35-4b-sm120-main-20260807.md)) | +| Qwen3.5-4B | bf16 direct-load | 0.26.0.dev0 | host PSS; frontend GPU measurement pending | #206 pretoken default/rollback CPU 6/6; exact IDs; GPU pending. Combined A/B +0.718% tput/-1.340% TTFT/-0.556% TPOT; cross-engine tput/TTFT VOID; TPOT/VRAM OPEN ([data](bench-evidence/qwen35-4b-sm120-main-20260807.md)) | ### GDN prefill causal-convolution by GPU diff --git a/docs/ENVIRONMENT.md b/docs/ENVIRONMENT.md index 6068401f5..b6d5e0969 100644 --- a/docs/ENVIRONMENT.md +++ b/docs/ENVIRONMENT.md @@ -26,6 +26,7 @@ These change how the engine runs and have no CLI flag (or complement one). | `VT_LMCACHE_HASH_ALGO` | `blake3` | Default LMCache key-derivation algorithm. Set `vllm` (alias `sha256_cbor`) for byte-for-byte interop with a real vLLM + LMCache peer. The `kv_connector_extra_config.hash_algo` key overrides it | | `VT_SERVER_MAX_PROMPT_CHARS` | `200000` characters | Rejects chat-completion prompts larger than this many characters. Set `0` to disable the prompt-size guard | | `VT_SERVER_MAX_NEW_TOKENS` | `4096` | Clamps the requested generation length to this many new tokens. Set `0` to disable the cap | +| `VT_BENCH_PRETOKENIZE` | `1` (on) | Makes `vllm-bench` encode every prompt before its benchmark clock and admit token IDs, matching the pinned vLLM comparison frontend. Exact `0` restores timed string admission for same-binary A/B; unset, `1`, and invalid spellings keep the safe default-on behavior | | `VT_VULKAN_DEVICE` | first suitable device | Forces the Vulkan physical device index. Required on a multi-GPU host to pin the intended device | | `VT_KV_CACHE_F32` | off (native KV dtype) | Forces the KV cache to fp32. A precision/diagnostic lever, at the cost of double the KV memory | | `VT_ENABLE_JUMP_FORWARD` | off | Opt-in to jump-forward constrained decoding (SGLang parity SW3): when a grammar/structured-output request reaches a state with exactly one valid next token, that token is emitted without a model step. Currently drives only the standalone driver (`DrainForcedTokens`); output-identical by construction (it fires only where the constrained sampler already has a single valid token), so it changes speed, never tokens. Off by default until the production scheduler splice (jumped-token KV recompute) lands. Set `1`/`true`/`on` to enable | diff --git a/docs/STATUS.md b/docs/STATUS.md index 36c068e7d..384abee43 100644 --- a/docs/STATUS.md +++ b/docs/STATUS.md @@ -78,7 +78,7 @@ token-for-token correctness against the pinned oracle. | Qwen3.6-27B (NVFP4) text generation | Correctness-complete; speed is CHECKPOINT-dependent | Token-exact GB10 on both. `unsloth` @`890bdef7` beats vLLM 0.25.0 every c (1.007-1.045x), 115/124; `nvidia` @`0893e160` (ModelOpt FP8 tower) is **0.85x BEHIND**, decode ~100% GPU-busy | | Qwen3.6-35B-A3B (NVFP4, GDN MoE) | Correctness-complete; binding grid @`a0fa12c7` FLAT 0.935x-0.979x over c1-c32 (CoV <0.81%); the prior 0.87x c2 / 0.92x c8 "weak cells" were harness mismatch, not code; memory PSS 3.81x, GPU 1.40x | Token-exact SYNC+ASYNC; `VT_ASYNC_DEVICE_MIRROR` ON fixes async batch-1 token-0 degeneration; `VT_ASYNC_EXECUTOR` Option A NEUTRAL → OFF | | Qwen3 / Qwen2 dense (BF16) | Correctness-complete, speed-pending. Async-serving P0 FIXED (`ROW-SERVE-ASYNC-DENSE-MIRROR`): classic-dense `Qwen3ForCausalLM` now honors the async device token-ids mirror; CPU-only -Werror test-guard fixes x2 | Near-tie-robust token-exact vs vLLM (Qwen3-0.6B, Qwen3-4B); c1 effective parity, c8 decode residual. **Async device-mirror (`ROW-SERVE-ASYNC-DENSE-MIRROR`, `f9c969ae`): the #31 fix ported to the classic dense family, dgx-VERIFIED.** The shared dense `EmbedInto` (qwen3.cpp) raced the async combine's device input-ids write against a stale host upload → token-0 degeneration on the depth-2 AsyncLLM serving path (quant-independent). `EmbedInto` now consumes the device override published by `ForwardQwen3ForCausalLM`'s `DeviceTokenIdsScope` (27B-dense template); gate `test_qwen3_dense_async_serving` RED on `VT_ASYNC_DEVICE_MIRROR=0`, GREEN default, byte-identical mirror-off. dgx GB10: async gate RED→GREEN 0.6B+4B, SACRED 0.6B+4B 184/184 unchanged (byte-neutral sync path), memcheck 0 errors; Yi30/Qwen3-8B-MXFP4 default-config e2e coherent + 3/4 token-exact (p2 = oracle-ratified near-tie, gap 0.0000), closing the QUANT-CT-MXFP4 async-default residual. RESIDUAL: sibling InternLM2/Mistral/Llama scope one-liner; W4 bench RAN; FA2 GQA-swap default-ON, c2-c8 <1.0x. `FLASH-PTXAS` #82: codegen at PARITY (no ptxas lever); gap=engine context. **D1 (2026-07-31, `CLAIM-D1-BF16-MERGED-QKV`): the bf16 merged-QKV path (`Qwen3QkvMergeEnabled`/`VT_QWEN3_QKV_MERGE`) is now default-ON** — one `vt::MatmulBT` over the merged `[qdim+2kdim,H]` owner + a contiguous `vt::QkvSplit` (OLMo-2 exemplar), replacing three per-shard GEMMs. Bit-exact GEMM math (A/B unit `test_ops_qkv_merge` byte-identical, RED-first); the wider-N cuBLASLt K-reduction flips the 0.6B genuine bf16 near-tie so the SACRED 0.6B golden was regenerated (all tokens within the near-tie band, max 0.125 nats), while Qwen3-4B is byte-neutral (0 diffs, stays STRICT). Re-gated 0.6B 16/16 + 4B 16/16; consistency/launch-count fold (measured NEUTRAL on 4B decode), no new throughput owed | -| Qwen3.5-4B plain BF16 direct loading on discrete CUDA | Correctness-complete; frontend correction pending; TPOT/VRAM open | Combined local A/B: +0.718% tput, -1.340% TTFT, -0.556% TPOT; IDs exact. Cross-engine tput/TTFT VOID pending pretokenized timing; TPOT 1.025x, VRAM +233 MiB OPEN ([data](bench-evidence/qwen35-4b-sm120-main-20260807.md)) | +| Qwen3.5-4B plain BF16 direct loading on discrete CUDA | Correctness-complete; frontend CPU-green/GPU-pending; TPOT/VRAM open | #206 pretoken default and `=0` string rollback preserve synthetic IDs (6/6). Cross-engine tput/TTFT VOID pending review/GPU; TPOT 1.025x, VRAM +233 MiB OPEN ([data](bench-evidence/qwen35-4b-sm120-main-20260807.md)) | | Qwen3-Coder-30B-A3B MoE (BF16) | Correctness-complete, speed-pending | Near-tie-robust token-exact 6/6; 11 of 16 binding grid cells at or above vLLM. **D1 (2026-07-31): inherits the default-ON bf16 merged-QKV via the shared dense `AttnBlock` — byte-neutral (0 token diffs, golden UNCHANGED); re-gated 6/6** | | Llama-3.x dense (BF16) | Correctness-complete, speed-pending | Near-tie-robust token-exact 16/16 (Llama-3.2-1B); llama3 RoPE scaling | | Mistral dense (BF16) | Correctness-complete, speed-pending | Paged-engine token-exact 16/16 (Mistral-7B-v0.3) | diff --git a/examples/bench/bench_core.h b/examples/bench/bench_core.h index 274e3a3a5..0f6bb70b2 100644 --- a/examples/bench/bench_core.h +++ b/examples/bench/bench_core.h @@ -36,6 +36,7 @@ #include #include #include +#include #include #include #include @@ -44,7 +45,9 @@ #include #include #include +#include #include +#include #include #include @@ -108,6 +111,7 @@ struct RequestRecord { double last_token_s = 0.0; // running: previous token arrival (for ITL). int prompt_tokens = 0; int output_tokens = 0; + std::vector prompt_token_ids; std::vector output_token_ids; std::vector itls; // inter-token latencies (s), one per chunk>1st. bool finished = false; @@ -120,6 +124,10 @@ struct BenchResult { // queue depth is one), matching vLLM's frontend across the ON/OFF control. bool async_frontend = false; bool async_scheduling_enabled = false; + // True when prompt tokenization was completed before the benchmark clock and + // timed admission used AsyncLLM's TokensPrompt overload. Auditable in every + // result so a benchmark artifact cannot silently mix frontend modes. + bool pretokenized_admission = false; int max_concurrent_batches = 1; int completed = 0; double duration_s = 0.0; @@ -127,6 +135,7 @@ struct BenchResult { int64_t total_output = 0; // Per-request generated IDs in submission order. This makes the benchmark // workload usable as a token-for-token correctness gate before timing it. + std::vector> prompt_token_ids; std::vector> output_token_ids; double request_throughput = 0.0; // req/s double output_throughput = 0.0; // tok/s (decode) @@ -174,6 +183,27 @@ inline double Mean(const std::vector& v) { return s / static_cast(v.size()); } +// Benchmark frontend parity selector. The pinned vLLM comparison tokenizes all +// prompts before starting its closed-loop clock and submits TokensPrompt IDs. +// Keep that production-parity path as the safe default: only exact `0` selects +// the timed string-admission rollback; invalid spellings stay default-ON. +inline bool ResolveBenchPretokenizedAdmission(const char* env_value) { + return env_value == nullptr || std::string_view(env_value) != "0"; +} + +// Pure production dispatch seam: exactly one callback is invoked. Keeping the +// selector separate from the benchmark loop lets the CPU contract test detect +// parser inversion and accidental double admission without hot-path counters. +template +inline decltype(auto) DispatchBenchPromptAdmission( + const char* env_value, PretokenizedCallback&& pretokenized_callback, + TimedStringCallback&& timed_string_callback) { + if (ResolveBenchPretokenizedAdmission(env_value)) { + return std::forward(pretokenized_callback)(); + } + return std::forward(timed_string_callback)(); +} + // ────────────────────────────── Synthetic model ─────────────────────────────── // A tiny hybrid-MoE Qwen3.6 (mirrors tests/vllm/v1/test_llm_engine.cpp) so the // harness runs end-to-end on the CPU box with no checkpoint. The NUMBERS from @@ -363,11 +393,12 @@ inline tok::Tokenizer BuildSyntheticTokenizer() { return t; } -// Build a prompt string that tokenizes to ~target tokens under `tok`. We can't -// hand the engine raw token ids (add_request takes text), so vLLM's exact -// random-token input is approximated with a repeated-word filler; the harness -// reports the MEASURED tokenized counts (from RequestOutput.prompt_token_ids), -// so throughput stays honest regardless of the small over/undershoot. +// Build a prompt string that tokenizes to ~target tokens under `tok`. The +// generated-workload mode intentionally remains human-readable rather than +// sampling arbitrary IDs; default admission pre-encodes this string before the +// clock and the rollback tokenizes the same string inside add_request. The +// harness reports the measured tokenized counts, so throughput stays honest +// regardless of the small over/undershoot. inline std::string BuildPrompt(const tok::Tokenizer& t, int target, uint64_t seed) { // NOTE: for the SYNTHETIC engine, only bytes in the tiny fixture's alphabet @@ -507,6 +538,23 @@ inline BenchResult RunBench(const BenchConfig& cfg) { vllm::v1::AsyncLLM& engine = loaded->async_engine(); + // Match the pinned vLLM comparison frontend: tokenize the complete workload + // before t0, preserving submission order and the string path's special-token + // processing, then move those IDs through AsyncLLM's TokensPrompt overload. + // Exact `VT_BENCH_PRETOKENIZE=0` retains the previous timed string overload + // and does not allocate this precomputed workload. + const char* const pretokenize_env = std::getenv("VT_BENCH_PRETOKENIZE"); + const bool pretokenized_admission = + ResolveBenchPretokenizedAdmission(pretokenize_env); + std::vector> pretokenized_prompts; + if (pretokenized_admission) { + pretokenized_prompts.reserve(prompts.size()); + for (const std::string& prompt : prompts) { + pretokenized_prompts.push_back( + loaded->tokenizer().EncodeWithSpecialTokens(prompt)); + } + } + std::map records; std::map active; const Clock::time_point t0 = Clock::now(); @@ -524,9 +572,19 @@ inline BenchResult RunBench(const BenchConfig& cfg) { RequestRecord rec; rec.arrival_s = now_s(); records[rid] = rec; + const size_t prompt_index = static_cast(next); active.emplace( - rid, engine.add_request(rid, prompts[static_cast(next)], - MakeSampling(cfg, next))); + rid, DispatchBenchPromptAdmission( + pretokenize_env, + [&]() { + return engine.add_request( + rid, std::move(pretokenized_prompts[prompt_index]), + MakeSampling(cfg, next)); + }, + [&]() { + return engine.add_request(rid, prompts[prompt_index], + MakeSampling(cfg, next)); + })); ++next; ++in_flight; } @@ -546,6 +604,7 @@ inline BenchResult RunBench(const BenchConfig& cfg) { RequestRecord& rec = records[out.request_id]; if (rec.prompt_tokens == 0 && !out.prompt_token_ids.empty()) { rec.prompt_tokens = static_cast(out.prompt_token_ids.size()); + rec.prompt_token_ids = out.prompt_token_ids; } if (!out.outputs.empty() && !out.outputs[0].token_ids.empty()) { const double t = now_s(); @@ -607,7 +666,9 @@ inline BenchResult RunBench(const BenchConfig& cfg) { BenchResult res; res.async_frontend = true; res.async_scheduling_enabled = loaded->async_scheduling_enabled(); + res.pretokenized_admission = pretokenized_admission; res.max_concurrent_batches = loaded->max_concurrent_batches(); + res.prompt_token_ids.resize(static_cast(cfg.num_prompts)); res.output_token_ids.resize(static_cast(cfg.num_prompts)); res.completed = done; res.duration_s = dur_s; @@ -646,6 +707,7 @@ inline BenchResult RunBench(const BenchConfig& cfg) { if (request_index >= res.output_token_ids.size()) { throw std::runtime_error("benchmark request id is out of range"); } + res.prompt_token_ids[request_index] = kv.second.prompt_token_ids; res.output_token_ids[request_index] = kv.second.output_token_ids; } // Speculative-decoding acceptance telemetry (real-checkpoint spec-ON only; @@ -686,6 +748,7 @@ inline void PrintReport(const BenchConfig& cfg, const BenchResult& r, std::fprintf(out, "%-42s %-12s\n", "Engine frontend:", r.async_frontend ? "AsyncLLM" : "LLMEngine"); line_i("Async scheduling enabled:", r.async_scheduling_enabled ? 1 : 0); + line_i("Pretokenized prompt admission:", r.pretokenized_admission ? 1 : 0); line_i("Maximum concurrent batches:", r.max_concurrent_batches); line_i("Successful requests:", r.completed); line_i("Maximum request concurrency:", cfg.concurrency); diff --git a/tests/examples/test_bench.cpp b/tests/examples/test_bench.cpp index 72cbbbc65..35ceea619 100644 --- a/tests/examples/test_bench.cpp +++ b/tests/examples/test_bench.cpp @@ -8,13 +8,93 @@ #include +#include #include #include +#include +#include using vllm::bench::BenchConfig; using vllm::bench::BenchResult; +using vllm::bench::DispatchBenchPromptAdmission; using vllm::bench::RunBench; +namespace { + +class ScopedEnv { + public: + ScopedEnv(const char* name, std::optional value) : name_(name) { + if (const char* previous = std::getenv(name)) previous_ = previous; + if (value.has_value()) { + REQUIRE(::setenv(name, value->c_str(), /*overwrite=*/1) == 0); + } else { + REQUIRE(::unsetenv(name) == 0); + } + } + + ~ScopedEnv() { + if (previous_.has_value()) { + (void)::setenv(name_.c_str(), previous_->c_str(), /*overwrite=*/1); + } else { + (void)::unsetenv(name_.c_str()); + } + } + + ScopedEnv(const ScopedEnv&) = delete; + ScopedEnv& operator=(const ScopedEnv&) = delete; + + private: + std::string name_; + std::optional previous_; +}; + +struct AdmissionObservation { + int result = 0; + int pretokenized_calls = 0; + int string_calls = 0; +}; + +AdmissionObservation ObserveAdmission(const char* env_value) { + AdmissionObservation observed; + observed.result = DispatchBenchPromptAdmission( + env_value, + [&]() { + ++observed.pretokenized_calls; + return 17; + }, + [&]() { + ++observed.string_calls; + return 29; + }); + return observed; +} + +} // namespace + +TEST_CASE("bench: pretokenized admission dispatch is default-on with exact rollback") { + const AdmissionObservation unset = ObserveAdmission(nullptr); + CHECK(unset.result == 17); + CHECK(unset.pretokenized_calls == 1); + CHECK(unset.string_calls == 0); + + const AdmissionObservation enabled = ObserveAdmission("1"); + CHECK(enabled.result == 17); + CHECK(enabled.pretokenized_calls == 1); + CHECK(enabled.string_calls == 0); + + const AdmissionObservation rollback = ObserveAdmission("0"); + CHECK(rollback.result == 29); + CHECK(rollback.pretokenized_calls == 0); + CHECK(rollback.string_calls == 1); + + // Invalid spellings keep the safe, production-parity default rather than + // silently restoring timed tokenization. + const AdmissionObservation invalid = ObserveAdmission("banana"); + CHECK(invalid.result == 17); + CHECK(invalid.pretokenized_calls == 1); + CHECK(invalid.string_calls == 0); +} + TEST_CASE("bench: synthetic engine completes all requests with sane metrics") { BenchConfig cfg; cfg.num_prompts = 8; @@ -115,3 +195,36 @@ TEST_CASE("bench: ShareGPT dataset supplies exact prompts") { CHECK(r.total_output == 8); CHECK(r.total_input > 0); } + +TEST_CASE("bench: pretokenized default preserves prompt and output token IDs") { + BenchConfig cfg; + cfg.num_prompts = 5; + cfg.input_len = 12; + cfg.output_len = 6; + cfg.concurrency = 3; + cfg.seed = 31; + cfg.temperature = 0.0; + + BenchResult pretokenized; + { + ScopedEnv env("VT_BENCH_PRETOKENIZE", std::nullopt); + pretokenized = RunBench(cfg); + } + + BenchResult timed_string; + { + ScopedEnv env("VT_BENCH_PRETOKENIZE", std::string("0")); + timed_string = RunBench(cfg); + } + + CHECK(pretokenized.pretokenized_admission); + CHECK_FALSE(timed_string.pretokenized_admission); + REQUIRE(pretokenized.prompt_token_ids.size() == + static_cast(cfg.num_prompts)); + REQUIRE(timed_string.prompt_token_ids.size() == + static_cast(cfg.num_prompts)); + CHECK(pretokenized.prompt_token_ids == timed_string.prompt_token_ids); + CHECK(pretokenized.output_token_ids == timed_string.output_token_ids); + CHECK(pretokenized.total_input == timed_string.total_input); + CHECK(pretokenized.total_output == timed_string.total_output); +} From 6b98621f1f7918e52705bb036e69696765c67cbe Mon Sep 17 00:00:00 2001 From: Richard Palethorpe Date: Sun, 9 Aug 2026 20:17:11 +0100 Subject: [PATCH 06/17] test(bench): harden pretokenization contract (#206) Make benchmark preparation own special-token encoding before clock start, and strengthen the synthetic fixture so raw encoding cannot masquerade as InputProcessor parity. Capture both resolved report modes. Focused gate: test_bench 10/10 tests, 70/70 assertions. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: Codex:gpt-5.6-sol [Codex] --- .agents/NOW.md | 2 +- .agents/engine-matrix.md | 2 +- .../specs/sm120-qwen35-pareto-2026-08-09.md | 12 +- docs/BENCHMARKS.md | 2 +- docs/STATUS.md | 2 +- examples/bench/bench_core.h | 47 +++++-- tests/examples/test_bench.cpp | 117 ++++++++++++++++++ 7 files changed, 167 insertions(+), 17 deletions(-) diff --git a/.agents/NOW.md b/.agents/NOW.md index 8e0ebb162..0fb0b9dac 100644 --- a/.agents/NOW.md +++ b/.agents/NOW.md @@ -21,7 +21,7 @@ Work: exact-chunks on main `1ce0d662b`; sm_120 measured at `3d2581551`. | MiniMax-H3 | **PRUNED ckpts RUN (#241): Q8_0 renders, seam 0.9941** | same-binary A/B | | Kimi-Linear-48B | 122/128 held; grouped router parallelised, e2e NOT ESTABLISHED | ckpt is tiktoken-only: no warm server | | 35B binding grid | @`a0fa12c7`: **flat 0.935-0.979x, NO c2/c8 weak cell** (CoV <0.81%); mem PSS 3.81x | Attribute the flat ~5% mid-band; TTFT c2 0.872x | -| Qwen3.5-4B sm_120 | #206 pretoken CPU green; IDs exact vs rollback. Cross-engine tput/TTFT VOID; TPOT/VRAM OPEN | Review; GPU A/B | +| Qwen3.5-4B sm_120 | #206 pretoken 10/10; special IDs/pre-clock/report gated. Cross-engine tput/TTFT VOID; TPOT/VRAM OPEN | Re-review; GPU A/B | | RPi5 A76 CPU | **R5 asm GREEN; llama NOT MET**: 0.461x pf, 0.653x dec | W6: BF16 GEMM | | MXFP4 parity | c1 1.020, c2-c8 0.962-0.969. **#82 CLOSED: ptxas-lineage REFUTED** | TERMINAL: at parity | | SERVE-ASYNC-DENSE-MIRROR | **LANDED+VERIFIED** (`f9c969ae`): async mirror, dense Qwen3; SACRED 184/184 | Sibling scope one-liner | diff --git a/.agents/engine-matrix.md b/.agents/engine-matrix.md index 4a2cdc939..5e0abeab5 100644 --- a/.agents/engine-matrix.md +++ b/.agents/engine-matrix.md @@ -204,7 +204,7 @@ claims it. | `SERVE-HTTP-TRANSPORT` | Serving-socket transport parity: mirror vLLM's uvicorn/asyncio default `TCP_NODELAY` on every accepted SSE socket so per-token stream frames are not held by Nagle against the peer's delayed ACK. Implemented + CPU-tested; the non-binding localhost A/B sizing is COMPLETE and NEUTRAL within noise on c1/c2 ITL/TPOT/throughput (loopback ACKs are instant, so Nagle never coalesces ~100 ms-cadence token frames) — no gate-axis credit expected; the mirror stays for real-network parity. Future keep-alive / read-write-timeout / listening-socket option parity noted, not done | T0 | vLLM serves via uvicorn over asyncio `vllm/entrypoints/launcher.py:71,76`, `vllm/entrypoints/openai/api_server.py:591,630`; asyncio disables Nagle per accepted TCP stream socket `asyncio/base_events.py:192-197` (`_set_nodelay`) called from `asyncio/selector_events.py:950`; cpp-httplib default-off `third_party/httplib/httplib.h:142`, applied on accept only when set `third_party/httplib/httplib.h:12083` | `src/vllm/entrypoints/openai/api_server.cpp:69` (`set_tcp_nodelay(true)` in the ApiServer setup) | behavioral accepted-socket `getsockopt(TCP_NODELAY)` case `tests/vllm/entrypoints/openai/test_api_server.cpp:1076` (helper `:380`); RED accepted `TCP_NODELAY` 0 → GREEN 1, full `test_openai_api_server` **22/22 cases / 242 assertions**; non-binding sizing root `~/work/vllm.cpp-tcpnodelay-sizing/ff915e8…` (raw-set SHA `f5b52900…2128`) neutral within noise; closure [ledger](parity-ledger.md#L451) | [serve-tcp-nodelay.md](specs/serve-tcp-nodelay.md) | `DONE` | `ff915e8` | | `SERVE-C-ABI` | Stable LocalAI-style C FFI (**19** exported `VLLM_API` symbols at `VLLM_ABI_VERSION 10`; blocking and nonblocking request handles. Count corrected 2026-07-24 from a stale `17`, which predated ABI v4/v5 adding `tool_parser`/`reasoning_parser` and the chat entry points; `include/vllm.h` is the source of truth and README:231 already said 19). **ABI v9 2026-07-28 (`CLAIM-CAPI-ENGINE-CONFIG-V9`): the ABI carried strictly LESS engine config than `EngineParams` does** - `max_num_batched_tokens`, the scheduler `scheduling_policy` (`fcfs` / `priority` / `lpm`), and `kv_transfer_config` (the external KV connector / LMCache JSON) were reachable from the bundled server's flags and from NO embedder. All three added, inert at their defaults (zero-filled v8 growth == byte-identical pre-v9 engine); the connector NAME is validated against `KVConnectorFactory` at load, mirroring the server's startup check. `tokenizer_config_path` stopped being a declared-since-v1 no-op and now selects the chat template's source file. Malformed `speculative_config`/`kv_transfer_config` documents now report `VLLM_ERR_INVALID_ARGUMENT` (the contract vllm.h documented since v6) instead of `VLLM_ERR_MODEL_LOAD`, via a catch scoped to the parse block so a real `FromModelDir` failure still reports MODEL_LOAD. Driver: the LocalAI vllm-cpp backend could not expose LMCache or the prefill budget in a model config) | T0 | Original project ABI; pinned vLLM has no C ABI | `include/vllm.h:143,181,207`; `src/capi/vllm_c.cpp:229,264,327,391` | `tests/capi/test_capi.cpp:320,428,505,574,606,640`; `tests/capi/test_dlopen.cpp:77,86`; `tests/capi/c_header_compile.c:1` | [c-api-library.md](specs/c-api-library.md) | `ANCHOR-BACKFILL` | `CLAIM-SERVE-C-ABI-SPIKE` | | `SERVE-CPP-API` | Rich `LLM` and `AsyncLLM` C++ API | T1 | `vllm/entrypoints/llm.py:66,422`; `vllm/v1/engine/async_llm.py:70` | - | - | `planned: specs/cpp-api.md` | `INVENTORIED` | - | -| `SERVE-CLI-BENCH` | Serve and latency/throughput/serve benchmark modes | T0 | `vllm/entrypoints/cli/serve.py:44`; `vllm/entrypoints/cli/benchmark/main.py:29`; production queue `vllm/v1/engine/core.py:200-231,622-669`; pinned comparison pretokenizes before timing `tools/bench/vllm_closed_loop_metrics.py:137-167` | separate binaries + explicit scheduler-capacity flags `examples/server/main.cpp:63,96,116,170`; production `AsyncLLM` benchmark frontend + auditable scheduler depth; #206 default pre-encodes every prompt before `t0` and admits token IDs, exact `VT_BENCH_PRETOKENIZE=0` retains timed-string admission, and `BenchResult` reports the resolved path (`examples/bench/bench_core.h:190,537-592,667,753`) | server-help/production-frontend/metric assertions plus #206 parser/callback-identity and synthetic default/rollback exact prompt/output-ID A/B: `tests/examples/test_bench.cpp:74,199` 6/6, 57 assertions CPU-GREEN; fresh mutation review and production GPU A/B remain pending | [CLI/serve/benchmark spike](specs/cli-serve-bench.md); [#206 campaign contract](specs/sm120-qwen35-pareto-2026-08-09.md) | `PARTIAL` | - | +| `SERVE-CLI-BENCH` | Serve and latency/throughput/serve benchmark modes | T0 | `vllm/entrypoints/cli/serve.py:44`; `vllm/entrypoints/cli/benchmark/main.py:29`; production queue `vllm/v1/engine/core.py:200-231,622-669`; pinned comparison pretokenizes before timing `tools/bench/vllm_closed_loop_metrics.py:137-167` | separate binaries + explicit scheduler-capacity flags `examples/server/main.cpp:63,96,116,170`; production `AsyncLLM` benchmark frontend + auditable scheduler depth; #206 default pre-encodes every prompt before `t0` and admits token IDs, exact `VT_BENCH_PRETOKENIZE=0` retains timed-string admission, and `BenchResult` reports the resolved path (`examples/bench/bench_core.h:190,212,582-586,698,780`) | server-help/production-frontend/metric assertions plus #206 parser/callback identity, preparation-before-clock, special-token/InputProcessor parity, report-mode capture and synthetic default/rollback exact prompt/output-ID A/B: `tests/examples/test_bench.cpp:123-212,316` 10/10, 70 assertions CPU-GREEN; fresh mutation re-review and production GPU A/B remain pending | [CLI/serve/benchmark spike](specs/cli-serve-bench.md); [#206 campaign contract](specs/sm120-qwen35-pareto-2026-08-09.md) | `PARTIAL` | - | | `SERVE-GATE-ONLINE` | Same-corpus online correctness, TTFT/TPOT/ITL, throughput and peak-memory gate vs vLLM v0.25.0 | T0 | `vllm/benchmarks/serve.py:1,581-615`; [v0.25 audit](sync/2026-07-12-702f481.md); `tests/benchmarks/test_serve_cli.py:1` | Schema-v5 harness plus [trace controller](../include/vt/cuda/cuda_profiler_control.h#L13), [production component driver](../scripts/dgx-gdn-packed-component.sh), and fail-closed [component finalizer](../tools/bench/gdn_packed_component.py) | **BINDING `9ecd9d0`: 114/124** (async default ON; mem 4/4, c1 20/20, c2 20/20, c16 19/20, c4 & c32 18/20, c8 15/20; `benchmark_binding` refers here, superseding `3f256ab` 55/124 and `246a23c` 49/124, both retained immutable). Two-grid totality with `f0fb727` (111/124) is 115/124 effective parity vs vLLM 0.25.0 (27B). Async CLOSED the c16/c32 ITL tails (ours now BEATS vLLM: c16 p99 1.055, c32 p90 1.034/p99 1.078) and leaves a stable c8 `p99_itl` ~0.86 residual, ROOT-CAUSED (2026-07-18, `CLAIM-C8-P99-TAIL-1`, [spec](specs/c8-p99-itl-tail-2026-07-18.md)) as IRREDUCIBLE-AS-MIRRORED: our deterministic synchronous forward keeps co-admitted c8 requests in byte-identical lockstep where vLLM's async-future jitter de-phases them; the c16/c32 INVERSION proves this is the trailing edge of the per-step determinism that wins c16/c32 + throughput, not a capability gap (scheduler + async placeholder byte-identical, `tests/vllm/v1/test_scheduler_wave.cpp:265`, [tail spec](specs/tail-stall-analysis-2026-07-16.md)). Full grid + per-binding forensics: roadmap_v1.md + parity ledger; no packed speed credit | [online serving gate](specs/cuda-online-serving-gate.md); [merged GDN projections](specs/gdn-merged-input-projections.md); [packed decode](specs/gdn-packed-decode.md) | `ANCHOR-BACKFILL` | CLAIM-SERVE-GATE-1 | | `SERVE-E2E-NIGHTLY` | Server conformance and real-model nightly suites for all release gates | T0 | `tests/entrypoints/openai/`; `tests/v1/e2e/`; `.buildkite/test-pipeline.yaml` | current unit/conformance tests only; no scheduled DGX suite | `tests/vllm/entrypoints/openai/test_conformance.cpp:1`; `tests/parity/test_qwen36_paged_engine.cpp:78`; `tests/parity/test_qwen27_paged_engine.cpp:110` | `planned: specs/server-e2e-nightly.md` | `INVENTORIED` | - | | `ENG-RELEASE-BINARIES` | Downloadable host-ABI-specific `vllm-server` bundles: adaptive CPU and fat CUDA primary artifacts, optional per-SM diagnostics, and literal-static feasibility boundary | T0 | vLLM release lanes `.buildkite/release-pipeline.yaml:1-18,34-170` @ `555967922`; release-image dependency boundary `docker/Dockerfile.cpu:262-290` | Required W1-W11/W13 implementation is complete in draft PR #196: ten-SM gencode/AOT, adaptive CPU tiers, extracted-archive validation and supply chain, least-privilege immutable handoff, eight primary bundles, byte-derived indexes, attestation, and exact-file publication; W12 remains optional/non-primary | Local manifest/archive/workflow/mutation gates, full adaptive x86 tier execution, clean CPU archive, and Vulkan 35/35 backend + 11/11 cross-device archive run are green; hosted ten-SM completion, full eight-tuple dry run, matching-hardware gates, and tagged publication remain pending, so no published-binary claim exists | [release-binary-matrix.md](specs/release-binary-matrix.md) | `ACTIVE` | `CLAIM-ENG-RELEASE-BINARIES-W1-W13` | diff --git a/.agents/specs/sm120-qwen35-pareto-2026-08-09.md b/.agents/specs/sm120-qwen35-pareto-2026-08-09.md index 30ef6215c..d21a7824a 100644 --- a/.agents/specs/sm120-qwen35-pareto-2026-08-09.md +++ b/.agents/specs/sm120-qwen35-pareto-2026-08-09.md @@ -137,10 +137,14 @@ The RED-first focused build failed because the production dispatch seam and auditable result fields did not exist. The implementation now pre-encodes all prompts with `EncodeWithSpecialTokens` before `t0`, moves them through the existing token-ID overload by default, and keeps exact `=0` string admission. -The synthetic CPU gate passes 6/6 tests and 57 assertions; default and rollback -produce identical prompt and output token IDs in submission order. This is not -a performance acceptance: fresh mutation review, operator verification and the -counterbalanced production GPU A/B remain pending. +The synthetic CPU gate passes 10/10 tests and 70 assertions; default and +rollback produce identical prompt and output token IDs in submission order. A +real `TemplateProcessing` BOS fixture proves raw `Encode` is not an equivalent +substitute, a pure preparation/clock seam proves all default tokenization ends +before timing begins, and report capture proves the resolved mode remains +auditable. This is not a performance acceptance: fresh mutation re-review, +operator verification and the counterbalanced production GPU A/B remain +pending. ## Correctness, tests, review and acceptance diff --git a/docs/BENCHMARKS.md b/docs/BENCHMARKS.md index ff72a8cf2..caab28b2b 100644 --- a/docs/BENCHMARKS.md +++ b/docs/BENCHMARKS.md @@ -41,7 +41,7 @@ The binding comparison. vLLM runs its **production graphed config**, never | Qwen3.6-27B | NVFP4 (`nvidia` @`0893e160`, ModelOpt `modelopt_mixed`) | 0.25.0 | 0/4 | **BEHIND, uniformly 0.85x** on decode throughput (was 0.72x before the FP8 tower fix); greedy continuation IDENTICAL to vLLM. A different model from the `unsloth` row (NVFP4 MLP + FP8 W8A8 GDN/attn tower) | | Qwen3.6-35B-A3B | NVFP4 `modelopt_mixed` | 0.25.0 | 2/18 | 3-rep grid 2026-08-05 @`1ea26427`: 0.93-1.03x (c4 wins), c16 0.93x. Both c16 levers A/B'd NEG: drain event -1.9%, mirror 0.999x. ★ probe found a prod async batch-1 greedy DEGENERATION bug the mirror fixes | | DeepSeek-V2-Lite | bf16 MLA | 0.25.0 | 4/25 | Attributed miss, row stays `ACTIVE` | -| Qwen3.5-4B | bf16 direct-load | 0.26.0.dev0 | host PSS; frontend GPU measurement pending | #206 pretoken default/rollback CPU 6/6; exact IDs; GPU pending. Combined A/B +0.718% tput/-1.340% TTFT/-0.556% TPOT; cross-engine tput/TTFT VOID; TPOT/VRAM OPEN ([data](bench-evidence/qwen35-4b-sm120-main-20260807.md)) | +| Qwen3.5-4B | bf16 direct-load | 0.26.0.dev0 | host PSS; frontend GPU measurement pending | #206 CPU 10/10 (special IDs/pre-clock/report); GPU pending. Combined +0.718% tput/-1.340% TTFT/-0.556% TPOT; cross-engine tput/TTFT VOID; TPOT/VRAM OPEN ([data](bench-evidence/qwen35-4b-sm120-main-20260807.md)) | ### GDN prefill causal-convolution by GPU diff --git a/docs/STATUS.md b/docs/STATUS.md index 384abee43..e9dfc9f44 100644 --- a/docs/STATUS.md +++ b/docs/STATUS.md @@ -78,7 +78,7 @@ token-for-token correctness against the pinned oracle. | Qwen3.6-27B (NVFP4) text generation | Correctness-complete; speed is CHECKPOINT-dependent | Token-exact GB10 on both. `unsloth` @`890bdef7` beats vLLM 0.25.0 every c (1.007-1.045x), 115/124; `nvidia` @`0893e160` (ModelOpt FP8 tower) is **0.85x BEHIND**, decode ~100% GPU-busy | | Qwen3.6-35B-A3B (NVFP4, GDN MoE) | Correctness-complete; binding grid @`a0fa12c7` FLAT 0.935x-0.979x over c1-c32 (CoV <0.81%); the prior 0.87x c2 / 0.92x c8 "weak cells" were harness mismatch, not code; memory PSS 3.81x, GPU 1.40x | Token-exact SYNC+ASYNC; `VT_ASYNC_DEVICE_MIRROR` ON fixes async batch-1 token-0 degeneration; `VT_ASYNC_EXECUTOR` Option A NEUTRAL → OFF | | Qwen3 / Qwen2 dense (BF16) | Correctness-complete, speed-pending. Async-serving P0 FIXED (`ROW-SERVE-ASYNC-DENSE-MIRROR`): classic-dense `Qwen3ForCausalLM` now honors the async device token-ids mirror; CPU-only -Werror test-guard fixes x2 | Near-tie-robust token-exact vs vLLM (Qwen3-0.6B, Qwen3-4B); c1 effective parity, c8 decode residual. **Async device-mirror (`ROW-SERVE-ASYNC-DENSE-MIRROR`, `f9c969ae`): the #31 fix ported to the classic dense family, dgx-VERIFIED.** The shared dense `EmbedInto` (qwen3.cpp) raced the async combine's device input-ids write against a stale host upload → token-0 degeneration on the depth-2 AsyncLLM serving path (quant-independent). `EmbedInto` now consumes the device override published by `ForwardQwen3ForCausalLM`'s `DeviceTokenIdsScope` (27B-dense template); gate `test_qwen3_dense_async_serving` RED on `VT_ASYNC_DEVICE_MIRROR=0`, GREEN default, byte-identical mirror-off. dgx GB10: async gate RED→GREEN 0.6B+4B, SACRED 0.6B+4B 184/184 unchanged (byte-neutral sync path), memcheck 0 errors; Yi30/Qwen3-8B-MXFP4 default-config e2e coherent + 3/4 token-exact (p2 = oracle-ratified near-tie, gap 0.0000), closing the QUANT-CT-MXFP4 async-default residual. RESIDUAL: sibling InternLM2/Mistral/Llama scope one-liner; W4 bench RAN; FA2 GQA-swap default-ON, c2-c8 <1.0x. `FLASH-PTXAS` #82: codegen at PARITY (no ptxas lever); gap=engine context. **D1 (2026-07-31, `CLAIM-D1-BF16-MERGED-QKV`): the bf16 merged-QKV path (`Qwen3QkvMergeEnabled`/`VT_QWEN3_QKV_MERGE`) is now default-ON** — one `vt::MatmulBT` over the merged `[qdim+2kdim,H]` owner + a contiguous `vt::QkvSplit` (OLMo-2 exemplar), replacing three per-shard GEMMs. Bit-exact GEMM math (A/B unit `test_ops_qkv_merge` byte-identical, RED-first); the wider-N cuBLASLt K-reduction flips the 0.6B genuine bf16 near-tie so the SACRED 0.6B golden was regenerated (all tokens within the near-tie band, max 0.125 nats), while Qwen3-4B is byte-neutral (0 diffs, stays STRICT). Re-gated 0.6B 16/16 + 4B 16/16; consistency/launch-count fold (measured NEUTRAL on 4B decode), no new throughput owed | -| Qwen3.5-4B plain BF16 direct loading on discrete CUDA | Correctness-complete; frontend CPU-green/GPU-pending; TPOT/VRAM open | #206 pretoken default and `=0` string rollback preserve synthetic IDs (6/6). Cross-engine tput/TTFT VOID pending review/GPU; TPOT 1.025x, VRAM +233 MiB OPEN ([data](bench-evidence/qwen35-4b-sm120-main-20260807.md)) | +| Qwen3.5-4B plain BF16 direct loading on discrete CUDA | Correctness-complete; frontend CPU-green/GPU-pending; TPOT/VRAM open | #206 frontend CPU 10/10: special IDs, pre-clock timing, mode report. Cross-engine tput/TTFT VOID pending GPU; TPOT 1.025x, VRAM +233 MiB OPEN ([data](bench-evidence/qwen35-4b-sm120-main-20260807.md)) | | Qwen3-Coder-30B-A3B MoE (BF16) | Correctness-complete, speed-pending | Near-tie-robust token-exact 6/6; 11 of 16 binding grid cells at or above vLLM. **D1 (2026-07-31): inherits the default-ON bf16 merged-QKV via the shared dense `AttnBlock` — byte-neutral (0 token diffs, golden UNCHANGED); re-gated 6/6** | | Llama-3.x dense (BF16) | Correctness-complete, speed-pending | Near-tie-robust token-exact 16/16 (Llama-3.2-1B); llama3 RoPE scaling | | Mistral dense (BF16) | Correctness-complete, speed-pending | Paged-engine token-exact 16/16 (Mistral-7B-v0.3) | diff --git a/examples/bench/bench_core.h b/examples/bench/bench_core.h index 0f6bb70b2..f3937c24b 100644 --- a/examples/bench/bench_core.h +++ b/examples/bench/bench_core.h @@ -204,6 +204,27 @@ inline decltype(auto) DispatchBenchPromptAdmission( return std::forward(timed_string_callback)(); } +// Own the ordering boundary between workload preparation and measurement. +// The clock callback is invoked exactly once and only after every default-path +// prompt has been encoded with the same special-token policy as InputProcessor. +// Supplying the clock keeps this pure host contract deterministic in tests. +template +inline auto PretokenizeBenchPromptsThenStartClock( + bool pretokenized_admission, const TokenizerLike& tokenizer, + const std::vector& prompts, + StartClockCallback&& start_clock_callback) { + std::vector> pretokenized_prompts; + if (pretokenized_admission) { + pretokenized_prompts.reserve(prompts.size()); + for (const std::string& prompt : prompts) { + pretokenized_prompts.push_back( + tokenizer.EncodeWithSpecialTokens(prompt)); + } + } + auto start = std::forward(start_clock_callback)(); + return std::make_pair(std::move(pretokenized_prompts), std::move(start)); +} + // ────────────────────────────── Synthetic model ─────────────────────────────── // A tiny hybrid-MoE Qwen3.6 (mirrors tests/vllm/v1/test_llm_engine.cpp) so the // harness runs end-to-end on the CPU box with no checkpoint. The NUMBERS from @@ -367,6 +388,20 @@ inline tok::Tokenizer BuildSyntheticTokenizer() { {"add_prefix_space", false}, {"trim_offsets", false}, {"use_regex", false}}})}}; + // Keep a real TemplateProcessing distinction in the benchmark fixture. The + // production string path applies this post-processor through InputProcessor; + // the pretokenized path must therefore use EncodeWithSpecialTokens too. + doc["post_processor"] = nlohmann::json::parse(R"json({ + "type": "TemplateProcessing", + "single": [ + {"SpecialToken": {"id": "", "type_id": 0}}, + {"Sequence": {"id": "A", "type_id": 0}} + ], + "pair": [], + "special_tokens": { + "": {"id": "", "ids": [20], "tokens": [""]} + } + })json"); nlohmann::json vocab = { {"h", 0}, {"e", 1}, {"l", 2}, {"o", 3}, {"w", 4}, {"r", 5}, {"d", 6}, {"Ġ", 7}, {"1", 8}, {"2", 9}, @@ -546,18 +581,12 @@ inline BenchResult RunBench(const BenchConfig& cfg) { const char* const pretokenize_env = std::getenv("VT_BENCH_PRETOKENIZE"); const bool pretokenized_admission = ResolveBenchPretokenizedAdmission(pretokenize_env); - std::vector> pretokenized_prompts; - if (pretokenized_admission) { - pretokenized_prompts.reserve(prompts.size()); - for (const std::string& prompt : prompts) { - pretokenized_prompts.push_back( - loaded->tokenizer().EncodeWithSpecialTokens(prompt)); - } - } + auto [pretokenized_prompts, t0] = PretokenizeBenchPromptsThenStartClock( + pretokenized_admission, loaded->tokenizer(), prompts, + []() { return Clock::now(); }); std::map records; std::map active; - const Clock::time_point t0 = Clock::now(); auto now_s = [&]() { return std::chrono::duration(Clock::now() - t0).count(); }; diff --git a/tests/examples/test_bench.cpp b/tests/examples/test_bench.cpp index 35ceea619..3b3a4c264 100644 --- a/tests/examples/test_bench.cpp +++ b/tests/examples/test_bench.cpp @@ -13,10 +13,15 @@ #include #include #include +#include +#include + +#include "vllm/v1/engine/input_processor.h" using vllm::bench::BenchConfig; using vllm::bench::BenchResult; using vllm::bench::DispatchBenchPromptAdmission; +using vllm::bench::PretokenizeBenchPromptsThenStartClock; using vllm::bench::RunBench; namespace { @@ -69,6 +74,50 @@ AdmissionObservation ObserveAdmission(const char* env_value) { return observed; } +std::string RenderReport(const BenchResult& result) { + std::FILE* const file = std::tmpfile(); + if (file == nullptr) throw std::runtime_error("tmpfile failed"); + vllm::bench::PrintReport(BenchConfig{}, result, file); + std::fflush(file); + std::rewind(file); + std::string report; + char buffer[1024]; + while (const size_t count = std::fread(buffer, 1, sizeof(buffer), file)) { + report.append(buffer, count); + } + std::fclose(file); + return report; +} + +int ReportedPretokenizedAdmission(const BenchResult& result) { + const std::string report = RenderReport(result); + const std::string label = "Pretokenized prompt admission:"; + const size_t begin = report.find(label); + if (begin == std::string::npos) return -1; + const size_t end = report.find('\n', begin); + const std::string line = report.substr(begin, end - begin); + int value = -1; + if (std::sscanf(line.c_str(), "Pretokenized prompt admission: %d", &value) != + 1) { + return -1; + } + return value; +} + +struct RecordingTokenizer { + std::vector* events = nullptr; + + std::vector EncodeWithSpecialTokens(std::string_view prompt) const { + events->push_back("special:" + std::string(prompt)); + return {20, static_cast(prompt.size())}; + } + + std::vector Encode(std::string_view prompt) const { + events->push_back("raw:" + std::string(prompt)); + return {static_cast(prompt.size())}; + } +}; + } // namespace TEST_CASE("bench: pretokenized admission dispatch is default-on with exact rollback") { @@ -95,6 +144,74 @@ TEST_CASE("bench: pretokenized admission dispatch is default-on with exact rollb CHECK(invalid.string_calls == 0); } +TEST_CASE("bench: report exposes the resolved pretokenized admission mode") { + BenchResult result; + result.pretokenized_admission = true; + CHECK(ReportedPretokenizedAdmission(result) == 1); + + result.pretokenized_admission = false; + CHECK(ReportedPretokenizedAdmission(result) == 0); +} + +TEST_CASE("bench: synthetic tokenizer distinguishes prompt special tokens") { + const vllm::tok::Tokenizer tokenizer = + vllm::bench::detail::BuildSyntheticTokenizer(); + const std::string prompt = "hello world"; + + CHECK(tokenizer.EncodeWithSpecialTokens(prompt) != tokenizer.Encode(prompt)); +} + +TEST_CASE("bench: default pretokenization completes before the clock") { + const std::vector prompts = {"hello", "world"}; + std::vector events; + const RecordingTokenizer tokenizer{&events}; + + auto [prepared, clock_value] = PretokenizeBenchPromptsThenStartClock( + vllm::bench::ResolveBenchPretokenizedAdmission(nullptr), tokenizer, + prompts, [&]() { + events.push_back("clock"); + return 73; + }); + + CHECK(events == std::vector{"special:hello", "special:world", + "clock"}); + CHECK(prepared == + std::vector>{{20, 5}, {20, 5}}); + CHECK(clock_value == 73); + + events.clear(); + auto [rollback, rollback_clock] = PretokenizeBenchPromptsThenStartClock( + /*pretokenized_admission=*/false, tokenizer, prompts, [&]() { + events.push_back("clock"); + return 91; + }); + CHECK(rollback.empty()); + CHECK(events == std::vector{"clock"}); + CHECK(rollback_clock == 91); +} + +TEST_CASE("bench: pretokenized vectors match timed-string InputProcessor") { + const vllm::tok::Tokenizer tokenizer = + vllm::bench::detail::BuildSyntheticTokenizer(); + const std::vector prompts = {"hello", "hello world"}; + auto [prepared, clock_value] = PretokenizeBenchPromptsThenStartClock( + vllm::bench::ResolveBenchPretokenizedAdmission(nullptr), tokenizer, + prompts, []() { return 17; }); + + const vllm::HfConfig config = + vllm::bench::detail::MakeSyntheticConfig(/*max_model_len=*/128); + const vllm::v1::InputProcessor input_processor(tokenizer, config); + REQUIRE(prepared.size() == prompts.size()); + for (size_t i = 0; i < prompts.size(); ++i) { + const vllm::v1::EngineCoreRequest timed_string = + input_processor.process_inputs(std::to_string(i), prompts[i], + vllm::SamplingParams{}, + /*arrival_time=*/0.0); + CHECK(prepared[i] == timed_string.prompt_token_ids); + } + CHECK(clock_value == 17); +} + TEST_CASE("bench: synthetic engine completes all requests with sane metrics") { BenchConfig cfg; cfg.num_prompts = 8; From 0473cd1e8bfb7deb21d1b61a57c6025aa391b332 Mon Sep 17 00:00:00 2001 From: Richard Palethorpe Date: Sun, 9 Aug 2026 20:37:33 +0100 Subject: [PATCH 07/17] spec(engine): require atomic benchmark wave admission (#206) Record the failed real pretokenized A/B as correctness-void and bind the next repair to atomic AsyncLLM/core-queue wave publication. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: Codex:gpt-5.6-sol [Codex] --- .agents/NOW.md | 4 +- .agents/benchmark-record.md | 48 +++++++++++ .agents/engine-matrix.md | 2 +- .agents/roadmap_v1.md | 2 +- .../specs/sm120-qwen35-pareto-2026-08-09.md | 83 +++++++++++++++++-- docs/BENCHMARKS.md | 2 +- docs/STATUS.md | 2 +- 7 files changed, 132 insertions(+), 11 deletions(-) diff --git a/.agents/NOW.md b/.agents/NOW.md index 0fb0b9dac..f32f071aa 100644 --- a/.agents/NOW.md +++ b/.agents/NOW.md @@ -21,7 +21,7 @@ Work: exact-chunks on main `1ce0d662b`; sm_120 measured at `3d2581551`. | MiniMax-H3 | **PRUNED ckpts RUN (#241): Q8_0 renders, seam 0.9941** | same-binary A/B | | Kimi-Linear-48B | 122/128 held; grouped router parallelised, e2e NOT ESTABLISHED | ckpt is tiktoken-only: no warm server | | 35B binding grid | @`a0fa12c7`: **flat 0.935-0.979x, NO c2/c8 weak cell** (CoV <0.81%); mem PSS 3.81x | Attribute the flat ~5% mid-band; TTFT c2 0.872x | -| Qwen3.5-4B sm_120 | #206 pretoken 10/10; special IDs/pre-clock/report gated. Cross-engine tput/TTFT VOID; TPOT/VRAM OPEN | Re-review; GPU A/B | +| Qwen3.5-4B sm_120 | #206 pretoken A/B FAIL: 98/128 exact; timings VOID. Sequential publish changes wave | Atomic admission; rerun | | RPi5 A76 CPU | **R5 asm GREEN; llama NOT MET**: 0.461x pf, 0.653x dec | W6: BF16 GEMM | | MXFP4 parity | c1 1.020, c2-c8 0.962-0.969. **#82 CLOSED: ptxas-lineage REFUTED** | TERMINAL: at parity | | SERVE-ASYNC-DENSE-MIRROR | **LANDED+VERIFIED** (`f9c969ae`): async mirror, dense Qwen3; SACRED 184/184 | Sibling scope one-liner | @@ -53,7 +53,7 @@ latency/memory on every axis, both gate models, reproduced 2–3x idle. See of roof. Dense-marlin +0.5%; Triton-AOT GDN a WASH. 2. **Spike the Parakeet encoder row** (vLLM: `nano_nemotron_vl.py`; the transducer half is NOT in vLLM: separate call). -3. **Qwen3.5-4B #206:** pretoken green; review/operator; GPU/vLLM rerun. +3. **Qwen3.5-4B #206:** atomic admission, review, then GPU/vLLM rerun. 2. **Merge the invocation-parity prevention** (CI guard + AGENTS.md checklist); CUDA build-verify the byte-exact `kGemvHeuristicAlgos` refactor on dgx. 4. **Restore `local-ai-worker`** on dgx at campaign end (`--restart=always`). diff --git a/.agents/benchmark-record.md b/.agents/benchmark-record.md index 2d5dbe0cc..037402374 100644 --- a/.agents/benchmark-record.md +++ b/.agents/benchmark-record.md @@ -18117,3 +18117,51 @@ Diagnostic evidence hashes: local log/token and `/tmp/qwen35-vllm-ttft-split-pin.*`. The next gate pre-tokenizes every local prompt before `t0`, submits the existing token-ID overload, then repeats counterbalanced same-binary rollback and fresh local/vLLM comparisons. + +## 2026-08-09 — pretokenized frontend real A/B failed token identity; timings void + +**Issue:** [#206](https://github.com/mudler/vllm.cpp/issues/206). +The first real-GPU pair used clean implementation/review head +`a33993a7cd8fdcf33b2b91112ef7a172b8f63fe1`, binary SHA-256 +`50f509cbde5865a6b0fe3b2bce0a05b68bf716337d41a90cb88d66480f1f8d02` +and the accepted K4 causal-conv plus post-conv-token-tile flags in both arms. +STRING-r1 set `VT_BENCH_PRETOKENIZE=0`; TOKENS-r1 set +`VT_BENCH_PRETOKENIZE=1`. The ShareGPT corpus SHA-256 remained +`9ea13603767c62c267e3f381fbccf42d0c9ca0c393655c37533eadca7aefca0c`. + +The hard correctness gate failed and stopped the series after that pair. Both +arms reported 131,784 input tokens and produced 128 requests x 128 output +tokens, but only **98/128 requests** and **15,507/16,384 positions** matched. +The differing request IDs were +`[1,7,10,23,30,41,46,50,51,60,61,68,70,74,77,80,82,83,85,86,88,92,95,110,113,116,122,124,126,127]`. +STRING-r1 output SHA-256 was +`83fcdc45f79ddb06a634c7d7d95eba3384543b3cd781a45a8db1fc4e2a453545`; +TOKENS-r1 was +`be20ffbceb61f0264ca21d972bfc5fc51e855e64f2b945de71669cae666aa702`. +Raw evidence is `/tmp/qwen35-ab-pretoken-a33993a7/`; the two log SHA-256 +values are `02c4381844c0a62aa2cf20023750ada2042d54f11f65189d7cf581060a6fe197` +and `3b0efb96c8b5ede5565bb12081d0d9010ccf10957a6fcb33ad7203ad766bc8d0`. + +| Invalid one-pair observation | total tok/s | output tok/s | mean TTFT | mean TPOT / ITL | mean E2E | +|---|---:|---:|---:|---:|---:| +| STRING-r1 | 6803.47 | 752.31 | 1012.01 ms | 34.65 ms | 5412.13 ms | +| TOKENS-r1 | 6839.78 | 756.32 | 1012.97 ms | 34.45 ms | 5388.73 ms | + +These values are **VOID**: changed output means changed work, so none receives +performance credit and the missing second/third repetitions are not a reason to +continue an invalid series. + +The mechanism is the admission boundary. Local `AsyncLLM` publishes and +notifies one request at a time, and its background `EngineCoreProc` blocks on +the first item then drains only items immediately available. String admission +tokenizes between publishes; pretokenized admission publishes fast enough to +change which requests are present for the first scheduling step. The pinned +wrapper synchronously adds every request in the initial concurrency before its +first `engine.step`, then publishes each refill group between steps. The next +discriminator is therefore an additive **atomic wave admission** API: prepare +all inputs and collectors first, register them under the shutdown lock, append +every core ADD under one queue lock, and notify once. The benchmark will record +arrivals for a whole refill wave and publish it once, so string rollback still +times tokenization while both arms execute identical batches. The full +RED-first, rollback and mutation gates are binding in the +[campaign spec](specs/sm120-qwen35-pareto-2026-08-09.md). diff --git a/.agents/engine-matrix.md b/.agents/engine-matrix.md index 5e0abeab5..fb8b32f87 100644 --- a/.agents/engine-matrix.md +++ b/.agents/engine-matrix.md @@ -204,7 +204,7 @@ claims it. | `SERVE-HTTP-TRANSPORT` | Serving-socket transport parity: mirror vLLM's uvicorn/asyncio default `TCP_NODELAY` on every accepted SSE socket so per-token stream frames are not held by Nagle against the peer's delayed ACK. Implemented + CPU-tested; the non-binding localhost A/B sizing is COMPLETE and NEUTRAL within noise on c1/c2 ITL/TPOT/throughput (loopback ACKs are instant, so Nagle never coalesces ~100 ms-cadence token frames) — no gate-axis credit expected; the mirror stays for real-network parity. Future keep-alive / read-write-timeout / listening-socket option parity noted, not done | T0 | vLLM serves via uvicorn over asyncio `vllm/entrypoints/launcher.py:71,76`, `vllm/entrypoints/openai/api_server.py:591,630`; asyncio disables Nagle per accepted TCP stream socket `asyncio/base_events.py:192-197` (`_set_nodelay`) called from `asyncio/selector_events.py:950`; cpp-httplib default-off `third_party/httplib/httplib.h:142`, applied on accept only when set `third_party/httplib/httplib.h:12083` | `src/vllm/entrypoints/openai/api_server.cpp:69` (`set_tcp_nodelay(true)` in the ApiServer setup) | behavioral accepted-socket `getsockopt(TCP_NODELAY)` case `tests/vllm/entrypoints/openai/test_api_server.cpp:1076` (helper `:380`); RED accepted `TCP_NODELAY` 0 → GREEN 1, full `test_openai_api_server` **22/22 cases / 242 assertions**; non-binding sizing root `~/work/vllm.cpp-tcpnodelay-sizing/ff915e8…` (raw-set SHA `f5b52900…2128`) neutral within noise; closure [ledger](parity-ledger.md#L451) | [serve-tcp-nodelay.md](specs/serve-tcp-nodelay.md) | `DONE` | `ff915e8` | | `SERVE-C-ABI` | Stable LocalAI-style C FFI (**19** exported `VLLM_API` symbols at `VLLM_ABI_VERSION 10`; blocking and nonblocking request handles. Count corrected 2026-07-24 from a stale `17`, which predated ABI v4/v5 adding `tool_parser`/`reasoning_parser` and the chat entry points; `include/vllm.h` is the source of truth and README:231 already said 19). **ABI v9 2026-07-28 (`CLAIM-CAPI-ENGINE-CONFIG-V9`): the ABI carried strictly LESS engine config than `EngineParams` does** - `max_num_batched_tokens`, the scheduler `scheduling_policy` (`fcfs` / `priority` / `lpm`), and `kv_transfer_config` (the external KV connector / LMCache JSON) were reachable from the bundled server's flags and from NO embedder. All three added, inert at their defaults (zero-filled v8 growth == byte-identical pre-v9 engine); the connector NAME is validated against `KVConnectorFactory` at load, mirroring the server's startup check. `tokenizer_config_path` stopped being a declared-since-v1 no-op and now selects the chat template's source file. Malformed `speculative_config`/`kv_transfer_config` documents now report `VLLM_ERR_INVALID_ARGUMENT` (the contract vllm.h documented since v6) instead of `VLLM_ERR_MODEL_LOAD`, via a catch scoped to the parse block so a real `FromModelDir` failure still reports MODEL_LOAD. Driver: the LocalAI vllm-cpp backend could not expose LMCache or the prefill budget in a model config) | T0 | Original project ABI; pinned vLLM has no C ABI | `include/vllm.h:143,181,207`; `src/capi/vllm_c.cpp:229,264,327,391` | `tests/capi/test_capi.cpp:320,428,505,574,606,640`; `tests/capi/test_dlopen.cpp:77,86`; `tests/capi/c_header_compile.c:1` | [c-api-library.md](specs/c-api-library.md) | `ANCHOR-BACKFILL` | `CLAIM-SERVE-C-ABI-SPIKE` | | `SERVE-CPP-API` | Rich `LLM` and `AsyncLLM` C++ API | T1 | `vllm/entrypoints/llm.py:66,422`; `vllm/v1/engine/async_llm.py:70` | - | - | `planned: specs/cpp-api.md` | `INVENTORIED` | - | -| `SERVE-CLI-BENCH` | Serve and latency/throughput/serve benchmark modes | T0 | `vllm/entrypoints/cli/serve.py:44`; `vllm/entrypoints/cli/benchmark/main.py:29`; production queue `vllm/v1/engine/core.py:200-231,622-669`; pinned comparison pretokenizes before timing `tools/bench/vllm_closed_loop_metrics.py:137-167` | separate binaries + explicit scheduler-capacity flags `examples/server/main.cpp:63,96,116,170`; production `AsyncLLM` benchmark frontend + auditable scheduler depth; #206 default pre-encodes every prompt before `t0` and admits token IDs, exact `VT_BENCH_PRETOKENIZE=0` retains timed-string admission, and `BenchResult` reports the resolved path (`examples/bench/bench_core.h:190,212,582-586,698,780`) | server-help/production-frontend/metric assertions plus #206 parser/callback identity, preparation-before-clock, special-token/InputProcessor parity, report-mode capture and synthetic default/rollback exact prompt/output-ID A/B: `tests/examples/test_bench.cpp:123-212,316` 10/10, 70 assertions CPU-GREEN; fresh mutation re-review and production GPU A/B remain pending | [CLI/serve/benchmark spike](specs/cli-serve-bench.md); [#206 campaign contract](specs/sm120-qwen35-pareto-2026-08-09.md) | `PARTIAL` | - | +| `SERVE-CLI-BENCH` | Serve and latency/throughput/serve benchmark modes | T0 | `vllm/entrypoints/cli/serve.py:44`; `vllm/entrypoints/cli/benchmark/main.py:29`; production queue `vllm/v1/engine/core.py:200-231,622-669`; pinned comparison pretokenizes before timing and synchronously admits a complete concurrency wave before each explicit step `tools/bench/vllm_closed_loop_metrics.py:57-102,137-167` | separate binaries + explicit scheduler-capacity flags `examples/server/main.cpp:63,96,116,170`; production `AsyncLLM` benchmark frontend + auditable scheduler depth; #206 default pre-encodes every prompt before `t0` and admits token IDs, exact `VT_BENCH_PRETOKENIZE=0` retains timed-string admission, and `BenchResult` reports the resolved path (`examples/bench/bench_core.h:190,212,582-586,698,780`). The current one-at-a-time publish changes EngineCore's first partial wave between string/token modes; additive atomic wave admission is specified but absent | server-help/production-frontend/metric assertions plus #206 parser/callback identity, preparation-before-clock, special-token/InputProcessor parity, report-mode capture and synthetic default/rollback exact prompt/output-ID A/B: `tests/examples/test_bench.cpp:123-212,316` 10/10, 70 assertions CPU-GREEN. Real `a33993a7` A/B **FAILED** token identity (98/128 requests, 15,507/16,384 positions); no timing credit. Atomic queue visibility, rollback/shutdown/duplicate, one-item, async exact-ID and real counterbalanced gates are pending | [CLI/serve/benchmark spike](specs/cli-serve-bench.md); [#206 campaign contract](specs/sm120-qwen35-pareto-2026-08-09.md) | `PARTIAL` | - | | `SERVE-GATE-ONLINE` | Same-corpus online correctness, TTFT/TPOT/ITL, throughput and peak-memory gate vs vLLM v0.25.0 | T0 | `vllm/benchmarks/serve.py:1,581-615`; [v0.25 audit](sync/2026-07-12-702f481.md); `tests/benchmarks/test_serve_cli.py:1` | Schema-v5 harness plus [trace controller](../include/vt/cuda/cuda_profiler_control.h#L13), [production component driver](../scripts/dgx-gdn-packed-component.sh), and fail-closed [component finalizer](../tools/bench/gdn_packed_component.py) | **BINDING `9ecd9d0`: 114/124** (async default ON; mem 4/4, c1 20/20, c2 20/20, c16 19/20, c4 & c32 18/20, c8 15/20; `benchmark_binding` refers here, superseding `3f256ab` 55/124 and `246a23c` 49/124, both retained immutable). Two-grid totality with `f0fb727` (111/124) is 115/124 effective parity vs vLLM 0.25.0 (27B). Async CLOSED the c16/c32 ITL tails (ours now BEATS vLLM: c16 p99 1.055, c32 p90 1.034/p99 1.078) and leaves a stable c8 `p99_itl` ~0.86 residual, ROOT-CAUSED (2026-07-18, `CLAIM-C8-P99-TAIL-1`, [spec](specs/c8-p99-itl-tail-2026-07-18.md)) as IRREDUCIBLE-AS-MIRRORED: our deterministic synchronous forward keeps co-admitted c8 requests in byte-identical lockstep where vLLM's async-future jitter de-phases them; the c16/c32 INVERSION proves this is the trailing edge of the per-step determinism that wins c16/c32 + throughput, not a capability gap (scheduler + async placeholder byte-identical, `tests/vllm/v1/test_scheduler_wave.cpp:265`, [tail spec](specs/tail-stall-analysis-2026-07-16.md)). Full grid + per-binding forensics: roadmap_v1.md + parity ledger; no packed speed credit | [online serving gate](specs/cuda-online-serving-gate.md); [merged GDN projections](specs/gdn-merged-input-projections.md); [packed decode](specs/gdn-packed-decode.md) | `ANCHOR-BACKFILL` | CLAIM-SERVE-GATE-1 | | `SERVE-E2E-NIGHTLY` | Server conformance and real-model nightly suites for all release gates | T0 | `tests/entrypoints/openai/`; `tests/v1/e2e/`; `.buildkite/test-pipeline.yaml` | current unit/conformance tests only; no scheduled DGX suite | `tests/vllm/entrypoints/openai/test_conformance.cpp:1`; `tests/parity/test_qwen36_paged_engine.cpp:78`; `tests/parity/test_qwen27_paged_engine.cpp:110` | `planned: specs/server-e2e-nightly.md` | `INVENTORIED` | - | | `ENG-RELEASE-BINARIES` | Downloadable host-ABI-specific `vllm-server` bundles: adaptive CPU and fat CUDA primary artifacts, optional per-SM diagnostics, and literal-static feasibility boundary | T0 | vLLM release lanes `.buildkite/release-pipeline.yaml:1-18,34-170` @ `555967922`; release-image dependency boundary `docker/Dockerfile.cpu:262-290` | Required W1-W11/W13 implementation is complete in draft PR #196: ten-SM gencode/AOT, adaptive CPU tiers, extracted-archive validation and supply chain, least-privilege immutable handoff, eight primary bundles, byte-derived indexes, attestation, and exact-file publication; W12 remains optional/non-primary | Local manifest/archive/workflow/mutation gates, full adaptive x86 tier execution, clean CPU archive, and Vulkan 35/35 backend + 11/11 cross-device archive run are green; hosted ten-SM completion, full eight-tuple dry run, matching-hardware gates, and tagged publication remain pending, so no published-binary claim exists | [release-binary-matrix.md](specs/release-binary-matrix.md) | `ACTIVE` | `CLAIM-ENG-RELEASE-BINARIES-W1-W13` | diff --git a/.agents/roadmap_v1.md b/.agents/roadmap_v1.md index 4aa10dba2..01fb2bed7 100644 --- a/.agents/roadmap_v1.md +++ b/.agents/roadmap_v1.md @@ -101,7 +101,7 @@ models we already ship + benchmark. Full seam map + M0–M5 W-plan: | MEM | `ROAD-V1-MEM` | **Memory budgeting: auto-size to the declared workload by default, optional total-footprint cap, pre-flight error instead of an OOM (user-directed 2026-08-06, [#83](https://github.com/mudler/vllm.cpp/issues/83))** — the user-facing wart that every operator hits before they hit any perf question: vLLM makes you compute your own VRAM budget, express it as a PERCENT, and nail it or OOM | [coverage view §2](feature-matrix.md#2-kv-cache--memory), [porting inventory](porting-inventory.md) | — **M1+M2 LANDED 2026-08-08** ([`specs/kv-sizing.md`](specs/kv-sizing.md)): the absolute `--kv-cache-memory` knob sizes the pool via a group-aware `KVBytesPerBlock` divisor, `--num-blocks` is the override, both mirrored on the C ABI at v16; `ResolveNumBlocks` precedence `num_blocks > bytes > 256`, CPU-gated (`KVBytesPerBlock` 5/5 + `test_capi` v16). M3 (the `gpu_memory_utilization` profile run) stays dgx-gated. (M0 design grounded in vLLM `config/cache.py` + `gpu_worker.py:497-599`; GB10 unified-pool caveat carried) | `M1+M2 DONE` | **Source-verified 2026-08-06 (records-only, NO code).** WE ARE CURRENTLY BEHIND vLLM ON THIS AXIS, not ahead: there is NO memory profiling at all and the KV pool is a RAW BLOCK COUNT the user types by hand — `EngineParams::num_blocks = 256` (`include/vllm/entrypoints/model_loader.h:58`, beside `block_size = 32` `:57` / `max_model_len` `:59` / `max_num_seqs = 8` `:60`), exposed verbatim as `--num-blocks N` (`examples/server/main.cpp:100,203-204,370`), carried on the C ABI as `vllm_model_params.num_blocks` at the same 256 default (`src/capi/vllm_c.cpp:429,485`), landing as `BlockPool(num_gpu_blocks, ...)` which asserts `> 0` and otherwise TRUSTS it (`include/vllm/v1/core/block_pool.h:96,223`; `src/vllm/v1/core/block_pool.cpp:51`). So a user must convert "40 GB free, 32k context, concurrency 8" into a block count themselves — strictly worse ergonomics than a percentage. UPSTREAM HAS THREE KNOBS, all `config/cache.py`, all T0, all rowed at [porting-inventory.md](porting-inventory.md) §T0: `gpu_memory_utilization` (fraction of TOTAL, default 0.9), `kv_cache_memory_bytes` (absolute KV pool) and `num_gpu_blocks_override` (exact block pin), sized off a profile run as `total x utilization - non-torch - peak activation`. **Mirroring that is necessary but does NOT solve the complaint**, for three source-grounded reasons: (1) the fraction is of TOTAL not FREE, so on any shared card the right fraction is a function of what someone else already holds — exactly the arithmetic the user is being asked to do; (2) WEIGHTS LOAD BEFORE THE KNOB ENGAGES (utilization sizes the KV pool AFTER the model is resident), so an oversized model OOMs during load and never reaches the check — which is the failure operators actually hit; (3) 0.9 is taken whether or not it is needed (a 4B model on an 80 GB card reserves 72 GB it will never touch and blocks everything else on the device). THE DESIGN (user-ratified 2026-08-06) is three modes: **Mode 1 default = SIZE TO THE DECLARED WORKLOAD** — per-allocation-class accounting BEFORE allocating anything (weights from checkpoint metadata / safetensors header / GGUF manifest, known before reading a byte of tensor data; CUDA context measured at creation; peak activation from a profile run at `max_num_batched_tokens`; KV for `max_model_len x max_num_seqs` at the resolved `block_size`/KV dtype; CUDA-graph capture-set footprint) → allocate exactly that and LEAVE THE REMAINDER OF THE DEVICE FREE. This is the SURPASS over vLLM, which takes its 90% regardless of whether the workload needs 8 GiB or 80. **Mode 2 = a cap on the TOTAL ENGINE FOOTPRINT** (weights + activations + KV + graph pools + context), NOT on the KV pool alone — load-bearing, because a KV-only cap cannot prevent the weight-load OOM; three spellings of the same cap: `--memory-limit 40GiB` (absolute, the primary form), `--gpu-memory-utilization 0.85` (vLLM's exact flag name and fraction semantics so existing vLLM launch lines port unchanged, per [[mirror-vllm-always-no-asking]]) and `--num-gpu-blocks-override N` (upstream's reproducibility escape hatch — this is where today's `--num-blocks` GOES, demoted from primary knob to explicit override), with precedence spelled out and TESTED, not left to argument order. **Mode 3 = REFUSE BEFORE ALLOCATING** with the full per-class breakdown and remedies COMPUTED FROM THE ACTUAL BUDGET (`--max-model-len N` / `--max-num-seqs M` / `--kv-cache-dtype fp8` / smaller quant, each with the GiB it recovers) — "you are 43.9 GiB over and here are the three levers that close it" is the difference between a usable error and a stack trace. **UNIFIED-MEMORY HAZARD (not hypothetical):** on GB10 the ~119 GiB pool is UNIFIED, a fraction-of-total setting reserves HOST RAM too, and `gpu_memory_utilization=0.85` has HARD-REBOOTED our DGX three separate times ([[gb10-unified-memory-oom-reboots-box]]) — hence absolute bytes is the PRIMARY form with the percentage kept only for vLLM compatibility, and hence the accounting must know whether the pool is unified, which makes free/total + an is-unified predicate a PLATFORM-SEAM question belonging behind `ROAD-V1-C1`'s abstraction (note `Platform::needs_weight_staging()` was deliberately NOT `is_unified_memory()` because the latter FLIPS GB10 — the distinction matters here) rather than a CUDA-specific branch, since discrete and unified devices need different safety margins. CORRECTNESS: pool size changes preemption/scheduling TIMING but not emitted tokens, so the SACRED token-exact gates are unaffected — and M2's gate makes that explicit by re-running them with no block flag at all. **Next gate = M0 the `specs/kv-sizing.md` spike (accounting model + precedence rules + upstream `file:line`); then M1 a `MemoryBudget` computing required bytes per class WITHOUT allocating plus the Platform free/total + is-unified seam, unit-gated predicted-vs-actual weight bytes; M2 auto-sizing as the default with `--num-blocks` demoted to `--num-gpu-blocks-override`, gated by every existing model gate running with NO block flag and staying token-exact; M3 the three caps + precedence through the server flags and the C ABI, gated by our KV pool matching vLLM's own at a matched `--gpu-memory-utilization`; M4 the pre-flight refusal, gated by a deliberately over-subscribed config exiting cleanly on GB10 (non-zero exit, no OOM, NO BOX REBOOT) covering the unified-pool path specifically; M5 (optional) a runtime guard failing the REQUEST rather than the engine.** Docs (README, [STATUS](../docs/STATUS.md)) update in the SAME change as whichever milestone shifts externally-visible behaviour, per [[keep-readme-current]] | | 1 | `ROAD-V1-C1` | **Extensibility-first (USER PRIORITY 2026-07-18):** drop-in kernel ABI + the MISSING Platform seam + model self-registration — make new GPUs/archs/models ADDITIVE (plan: [extensibility-platform-seam-2026-07-18.md](specs/extensibility-platform-seam-2026-07-18.md)) | [`BACKEND-ABI-VT`](backend-matrix.md), [kernel matrix](kernel-matrix.md) | exhaustive kernel/dependency inventory and [raw-pointer adapter ABI](specs/dropin-kernel-abi.md) accepted; additive W0 implemented and CPU 94/94. `CLAIM-BACKEND-ABI-W0-GPU-1` repaired the GCC13/doctest blocker without runtime changes; exact sm_121a all-target build, focused CUDA/ABI sanitizer, and both gate-model tests pass at `1141b79`. Cross-arch/trace/A-B and scalar-forwarder/backend-shim debts remain explicit | `PARTIAL` | **★ NEW ORDER-1 HEAD (user-directed 2026-07-19): the PORTABLE AUTOMATIC OP-FUSION FRAMEWORK (`KERNEL-FUSION-FRAMEWORK`, spike [portable-fusion-framework.md](specs/portable-fusion-framework.md), `SPIKE`).** The extensibility cornerstone: fusions DECLARED ONCE (backend-agnostic `constexpr FusedRecipe` catalog above `vt::`, transcribing vLLM's finite pattern-pass set `passes/fusion/*` @ `pass_manager.py:138-200`, mirroring the `CustomOp` `forward_native`/`forward_cuda` seam `custom_op.py:103`) and REALIZED PER-BACKEND through the existing `vt::` op table (Tier-0 composite = the CPU oracle inherited free by any backend; Tier-1 interpreter = one kernel port per backend lights up every recipe). Makes a new vLLM fusion PR a ONE-declaration port, a new GPU an additive catalog realization, a new model an additive pattern declaration — the PR-#4 remedy, composed with the Platform/attn-registry/model-registry seams below. The TDR Phase-0 skeleton is already LANDED (`fused_recipe.h`/`recipes.h` one recipe + `OpId::kFusedChain` Tier-0/1 on CPU+CUDA + byte-exact `test_ops_fused_chain.cpp`). **W0 ADOPTED 2026-07-19 (`CLAIM-FUSION-FRAMEWORK-W0`):** the seam is now used in production at ONE real site — the 35B `RunLayerPaged` post-attention layernorm routes its plain add+residual+gemma-RMSNorm through `vt::FusedChain(kFusedAddRmsNorm)` (`VT_FUSED_CHAIN_ADOPT` default-ON / `=0` rollback), behaviour-preserving + byte-identical to the prior hand-call (DGX: clean CUDA `-Werror` 0 warn, byte-exact composite==interp==golden incl. H=2048, 35B 315/315 + 27B 235/235 token-exact BOTH arms, memcheck 0 errors) — proving the declare-once/realize-per-backend seam end-to-end; the current 3-opcode/4-role POD sufficed byte-identically for the plain add+rmsnorm class, so W0 needed NO generalization. **W1 GENERALIZED the POD 2026-07-20 (`CLAIM-FUSION-FRAMEWORK-W1`, `1115648`):** full activation/norm/quant/rope opcode set + indexed operand table; all 5 quant-fused chains declared byte-exact; Tier-0 composite = ONE device-agnostic walker (kills CPU/CUDA oracle drift); infrastructure only, no call site changed (DGX: `-Werror` 0-warn, byte-exact CPU 196 + CUDA 361, memcheck 0, 27B 235/235 + 35B 315/315 both arms). **W2 MIGRATED the hand-fusions 2026-07-20 (`CLAIM-FUSION-FRAMEWORK-W2`):** the framework now OWNS the fusion dispatch — each recipe binds (new backend-agnostic `FusedRecipe.fast_op`) to its EXISTING single-launch fused kernel, so `FusedChain(recipe)` dispatches to the SAME fast kernel the model called directly pre-migration (byte-identical + perf-neutral by construction; composite is the graceful fallback + oracle). SIX call sites migrated behind `VT_FUSED_CHAIN_ADOPT` (`kSiluMulFp4Quant`/`kSigmoidGateFp4Quant`/`kRmsNormGatedQuantFp8`×2/`kRmsNormQuantFp8`/`kAttnQkNormRopeGate`×2). DGX: `-Werror` 0-warn, byte-exact CPU 228 + CUDA 420, memcheck 0, 27B 235/235 + 35B 315/315 BOTH arms. **W3 MECHANICAL-SYNC PROOF LANDED 2026-07-20 (`CLAIM-FUSION-FRAMEWORK-W3`):** ported a NEW, previously-unported vLLM fusion pass — `SiluMulFp8StaticQuantPattern` (`act_quant_fusion.py:81` → `_C.silu_and_mul_quant`, the static-per-tensor-FP8 sibling of `kSiluMulFp4Quant`) — as ONE `constexpr FusedRecipe kSiluMulQuantFp8` + its byte-exact test, touching EXACTLY 2 shared files (`recipes.h` + `test_ops_fused_chain.cpp`), NO kernel/dispatch/composite-walker/model-site edit and NO new primitive (composite = existing `vt::MoeSiluMul` + `vt::QuantFp8Static`; `fast_op=kNoFastOp`). The PR-#4 additivity test made concrete: a whole new fusion pattern = one declaration. DGX: `-Werror` 0-warn, byte-exact CUDA 432, memcheck 0, no token regression (recipe declared-only) 27B 235/235 + 35B 315/315. **W4 BACKEND-ADDITIVITY PROOF LANDED 2026-07-20 (`CLAIM-FUSION-FRAMEWORK-W4`) — the W-series proof milestone is DONE:** made the additivity claim EXECUTABLE — new test `test_fused_chain_additivity.cpp` treats the EXISTING CPU backend AS the 'second backend' relative to CUDA (no mock `DeviceType` — that would edit the core enum + every switch, ironically non-additive) and, in ONE generic loop over the WHOLE catalog (all 7 recipes), asserts each runs byte-exact on the CPU backend via the Tier-0 composite — 4 CPU-full end-to-end + 3 fp8-prefix (byte-exact prefix + the FULL composite asserted to THROW on CPU, documenting the CUDA-only static-fp8 backend-negotiated tail, §3b/§6). Additivity evidence: catalog `recipes.h` grew 1→6→7 while the composite walker stayed ONE per-OPCODE function (12 `FOp::` cases) + the CPU/CUDA `kFusedChain` registration ONE line each + `cpu_ops.cpp` never `#include`s `recipes.h` — W3's whole new recipe `kSiluMulQuantFp8` is in ZERO backend TUs, inherited free. CPU `-Werror` 0-warn, `test_fused_chain_additivity` 17/17 + `test_ops_fused_chain` 228/228; engine byte-identical (no `src/`/`include/` change) ⇒ 27B 235/235 + 35B 315/315 structurally unchanged; memcheck N/A. Honest deferred (named, non-blocking the ORDER-1 milestone): Tier-1 perf interpreter for the quant chains (composite-only today), a REAL Metal/Vulkan catalog realization (M4 HW-blocked), and per-recipe fast single-launch kernels. Honest payoff: perf ceiling ~3.5%/step compute-bound on 35B (NOT a perf lever — tasks #61/#62; W0 is perf-neutral by construction); primary value = extensibility + mechanical upstream-sync + CPU/CUDA oracle-drift elimination. Incremental W0 adopt-one **(DONE)** → W1 generalize POD **(DONE)** → W2 migrate hand-fusions **(DONE)** → W3 mechanical-sync proof **(DONE)** → W4 backend-additivity proof **(DONE)** → Wn honest re-measure (optional, off the extensibility critical path). **W-SERIES ORDER-1 PROOF MILESTONE DONE 2026-07-20 (`CLAIM-FUSION-FRAMEWORK-W4`).** **PRIOR extensibility items (all LANDED, the seams this composes with):** **#1 extensibility item — extract the Platform seam — LANDED 2026-07-18** (`BACKEND-PLATFORM` `ACTIVE`, `CLAIM-BACKEND-PLATFORM-1`): `include/vllm/platforms/interface.h` + `src/vllm/platforms/{platform,cpu,cuda}.cpp` mirror `vllm/platforms/interface.py:134-229` 1:1; `CurrentPlatform()` self-registered per `DeviceType`; the 7 memory-model/residency `device.type == kCUDA` sites (of PR #4's ~37) now route through it → new-GPU memory model is ONE additive `platforms/.cpp`. Behavior-preserving (clean CPU build + `test_platform` + full CPU CTest green; DGX 235/235 + 315/315 pending). **Item 2 residency-as-Platform-capability LANDED 2026-07-19** (`CLAIM-BACKEND-PLATFORM-2`): the host-free / load-stream / DevicePool-cap decisions in `qwen3_5.cpp` now READ `GetPlatform(.device.type).residency_policy()` (per-device) via the pure `ShouldReleaseHostWeights`/`ShouldInterleaveLoadStream` helpers + `device_pool_cap_bytes`, not an inline `device.type`/env gate; `CudaPlatform.release_host_weights_after_upload` flipped false→true (now CONSUMED ⇒ reproduces today's GB10 host-free-after-Marlin + ~4 GiB load peak EXACTLY); `MarlinMoeEnabled()` stays the orthogonal kernel-path gate. **A new (discrete) GPU sets `residency_policy()` values ⇒ ZERO model edit.** Behavior-preserving (clean CPU build + `test_platform` consumption cases 7/43 + full CPU CTest + tools 164/164 green; **DGX-CONFIRMED @ `62fc0e0`: clean CUDA `-Werror`, 27B 235/235 + 35B 315/315 token-exact, 35B VmHWM ≈ 4.0 GiB load-stream win preserved, memcheck 0 errors**). Then item 3 drop-in ABI family migration. **Item 4 attn-backend registry LANDED 2026-07-19** (`CLAIM-ATTN-REGISTRY-1`, `BACKEND-ATTN-REGISTRY`): NEW `include/vllm/v1/attention/registry.{h,cpp}` `(DeviceType,name)` registry + `SelectAttentionBackendName` selector (mirror `registry.py` self-registration + `cuda.py:361-470` `get_attn_backend_cls`/`_get_backend_priorities`); `Platform::get_attn_backend_priority()` filled (was the item-1 STUB) → capability-ordered name lists on `CudaPlatform` (major-10 vs else) + `CpuPlatform`; FLASH_ATTN/GDN self-register. **Adding a backend's attention = 1 self-registering TU + 1 priority slot, ZERO selector/model/runner edit.** Behavior-preserving — the walk returns FLASH_ATTN on CUDA+CPU (the same FA2 attention runs); clean CPU build + `test_attn_backend_registry` (8/25) + full CPU CTest, tools 164/164, checkers green; **DGX-CONFIRMED @ `2c732e7`: 27B 235/235 + 35B 315/315 token-exact (FA2 sm_121a), memcheck 0/315**. **Item 5 model self-registration LANDED 2026-07-19** (`CLAIM-MODEL-SELFREG-1`, `MODEL-FACTORY-registry`): the fixed `kRegistrations` array → `REGISTER_VLLM_MODEL(...)` static-`Registrar` idiom (`model_registry.h:167-189`) + Qwen dense/MoE arch entry points split into per-variant TUs (`qwen3_5_dense.cpp`/`qwen3_5_moe.cpp`) over shared `qwen3_5_common.{h,cpp}`, so **adding a model = 1 new TU + 1 REGISTER line, ZERO shared-array edit**; behavior-preserving (clean CPU build + `test_model_registry` extension + full CPU CTest, tools 164/164, checkers green; DGX 27B/35B token-exact pending). Deep `qwen3_5.cpp` machinery factoring deferred. Metal/MLX bring-up proves the seams (needs M4). **★ THE ARCH HALF OF THIS ITEM IS NOW PROVEN BY MEASUREMENT, NOT ARGUED (2026-07-22, `CLAIM-CUDA-SM120-BRINGUP`, [spec §W8](specs/cuda-arch-additivity.md)):** consumer-Blackwell `sm_120a` was brought up as a BUILD-supported target and required **ZERO kernel, model, runner, sampler or feature-table edits** — the additive seams (per-arch FEATURE TABLE, capability-keyed tactic registry keyed on `major == 12`, Platform auto-probe, `pageable && integrated` residency classification) already covered it, so the entire diff is build configuration, a configure-tier test and records. That is the PR-#4 additivity test passed on a real second architecture. It is deliberately NOT a runtime-support claim: no sm_120 board exists here. **★ THE MODEL/QUANT HALF NOW ADVANCES TOO (2026-07-23, `CLAIM-BACKEND-SEAM-S4-1`): the `model_executor/layers/` `LinearMethod`/`QuantizationConfig` seam the [accelerator-seam audit](specs/accelerator-seam-audit.md) §9 named ABSENT now EXISTS in part.** `S4` landed byte-identical: the dense model's projections route through a `method.Apply()` chosen ONCE from the checkpoint (retiring the per-call `IsNvfp4()` tensor-name probe), and 18 shared-layer `device==kCUDA` availability gates became `vt::OpRegistered` op-table queries — the policy(scheme)/implementation(kernel) split the audit's binding rule requires. **DSR 86 → 67**; all six SACRED gates byte-identical on dgx (27B/35B/Coder/dense/OPT/DeepSeek-V2); the fragile 27B-W4A4/fp8-recipe gates are correctly deferred to `S6` behind `S5`'s reference tier. **★ `S6` ASSESSED 2026-07-23 (`CLAIM-BACKEND-SEAM-S6-1`) → NO-OP / BLOCKED, DSR stays 67 (§11):** the deferred fp4/fp8 gates convert ZERO sites byte-identically — every one bottoms out at a **dual-registered** (CPU+CUDA) bespoke op (none CUDA-only, unlike S4's convertible gates), so `OpRegistered(op,dev)` is TRUE on `kCPU` ⇒ the class-A swap is bit-changing on the CPU reference/emulation path (two numerics per device); S5's reference tier does not change this (those CPU kernels are present natively, never a miss). No `src/`/`include/`/test byte changed, no baseline moved. The genuine byte-identical unlock is re-scoped to **`S3`** (Platform capability fields mirroring `supports_fp8`/`cutlass_fp4_supported` — the audit's own class-D fix) and **`S7`** (layer extraction); the plan's `~37` S6 target assumed the class-A `OpRegistered` swap was byte-identical, which holds only for CUDA-only ops (all taken by S4). **★ `S3` LANDED 2026-07-23 (`CLAIM-BACKEND-SEAM-S3-1`) — the byte-identical unlock S6 re-scoped:** mirrors vLLM's `Platform` capability surface (`supports_fp8`/`cutlass_fp4_supported`/`opaque_attention_op`/`is_integrated_gpu`/`support_static_graph_mode`/`is_device_capability_family`, base false in `interface.h`; CudaPlatform answers GB10 values in `cuda.cpp`/`platform.cpp`) and converts **12** deferred `qwen3_5.cpp` gates onto it (7 fp4-act `cutlass_fp4_supported`, 3 fp8-fused `supports_fp8`, 2 decode-graph `support_static_graph_mode`) — byte-identical because a capability answers the base false off CUDA, exactly what `device==kCUDA` did (where S6's `OpRegistered` was TRUE on `kCPU`), and it DECOUPLES (a future accelerator answers for itself). **DSR 67 → 55** (`kcuda` 25→13), baseline lowered same commit, ratchet + 24-case suite green; all seven SACRED gates byte-identical on dgx (27B 235/235 · 35B 315/315 · Coder · dense-32B · OPT · DeepSeek-V2 · Llama), `test_platform` CUDA-leg proves each predicate == former `device==kCUDA` on GB10, memcheck 0 errors, clean CUDA+CPU `-Werror`. Residency/stream/FA2-dtype/merged-layout sites LEFT for `S7`. **★ `S7` LANDED 2026-07-23 (`CLAIM-BACKEND-SEAM-S7-1`) — the seam campaign's TERMINAL runtime-decoupling state, the closest this extensibility work comes to a finish line:** ALL 23 remaining runtime `kCUDA`/`is_cuda()` sites in the shared model layer hoisted onto capabilities — new `Platform::needs_weight_staging()` (the CUDA device-resident staging policy, NOT `is_unified_memory()` which would FLIP GB10; covers residency/merged-GDN/packed-decode/direct-load), `Platform::supports_fa2_attention()` (FA2 dtype), `Backend::SupportsAuxStream()` (MoE aux-stream), reusing S3's `supports_fp8`/`cutlass_fp4_supported`/`support_static_graph_mode`/`is_integrated_gpu` (runner combine/scatter) and `vt::OpRegistered(kMoeGroupedGemmBf16)` (a CUDA-only op). Each returns the former `device==kCUDA` value on GB10 → byte-identical. **DSR 55 → 32 — the IRREDUCIBLE build-gate floor:** the shared model layer holds ZERO runtime device tests; the 32 residual are all `#ifdef VT_*` compile-time gates for kernels that only build on one GPU family (a kernel that only compiles on one arch is legitimately irreducible), so the audit's `<10` is NOT reachable and this is the honest answer to "how additive can the shared layer get" — every runtime device coupling is gone. baseline lowered same commit, ratchet + 24-case suite green; all seven SACRED gates byte-identical on dgx (27B 235/235 · 35B 315/315 · Coder 6/6 · dense-32B 16/16 · OPT 6/6 · DeepSeek-V2 8/8 · Llama 16/16), new `test_platform`/`test_backend`/`test_cuda_backend` cases green, memcheck 0, clean CUDA+CPU `-Werror`. The `layers/`-library physical relocation (shrinking `qwen3_5.cpp` toward `qwen3_next.py`'s 802-line shape) is a follow-on refactor; the device coupling it was to remove is already gone | | 2 | `ROAD-V1-C2` | Model families: Llama/Qwen3/Mistral, MoE, Qwen3-Next | [model matrix](model-matrix.md) | current pin has 353 static IDs; v0.25.0 adds three sync-target rows (MOSS-Transcribe-Diarize, Laguna DFlash, Bailing hybrid MTP), yielding 356 after pin advance. **FIRST ADDITIVE-MODEL BRING-UP W0-W4 LANDED 2026-07-20 — CORRECTNESS COMPLETE (0.6B + 4B gates PASS 16/16; SPEED pending)** ([first-additive-model-qwen3-dense.md](specs/first-additive-model-qwen3-dense.md), `MODEL-TEXT-qwen3-qwen3-for-causal-lm` `ACTIVE`(correctness COMPLETE, speed pending), runner generalization `ENG-RUNNER-MODELSHAPE`, `CLAIM-MODEL-QWEN3-DENSE`) **MLA CAMPAIGN SPIKED 2026-07-21** ([mla-deepseek-campaign](specs/mla-deepseek-campaign.md), `CLAIM-MLA-DEEPSEEK`): five rows `INVENTORIED` -> `SPIKE` (DeepSeek V2 / V3+V3.2 / v1-MHA, Kimi-Linear, MiniMax-M2). **KIMI-LINEAR-48B W0 DEDICATED SPIKE 2026-08-05** ([kimi-linear.md](specs/kimi-linear.md), `CLAIM-KIMI-LINEAR-W0`): full dedicated W0 spike for `MODEL-TEXT-kimi-linear-*` (stays `SPIKE` — actively claimed) — the ONE Kimi text model that FITS one GB10 (91.5 GiB, 0.77x pool) with a real e2e SACRED gate; HEAVY reuse (MLA + sigmoid/`noaux_tc` MoE + GDN family + KDA host refs landed), NET-NEW = the KDA device kernel + NoPE-MLA branch + hybrid schedule/loader; W1 implementation can start. Answers the Tier-3 "MLA = new attention, new campaign" item in [breadth-sweep-plan](specs/breadth-sweep-plan.md) §B.3. Key determinations: GB10/sm_121 selects **`TRITON_MLA`** for dense MLA decode and **`FLASH_ATTN`** for MLA prefill (`vllm/platforms/cuda.py:129-133`, `mla/prefill/selector.py:74-77`), so the sm90/sm100-only MLA kernels are out of reach and out of scope; the cross-cutting cost is the **compressed-latent KV cache** (one 576-wide latent per token, `num_kv_heads=1`, no separate V), which our allocator and `vt::ReshapeAndCache`/`vt::PagedAttention` cannot express; and **only DeepSeek-V2-Lite (~29.3 GiB bf16) fits GB10** — V3/V3.2, Kimi-K2.5, MiniMax-M2/M3 are HW-BLOCKED e2e, Kimi-Linear-48B is HW-MARGINAL. W0-W10 plan recorded; nothing implemented. **GLM + DSA + LATEST-DEEPSEEK SPIKED 2026-07-21** ([glm-dsa-latest-deepseek](specs/glm-dsa-latest-deepseek.md), `CLAIM-GLM-DSA-LATEST-DEEPSEEK`): seven rows `INVENTORIED` -> `SPIKE` (ChatGLM, Glm, Glm4, Glm4Moe, Glm4MoeLite, GlmMoeDsa, DeepSeek-V4). Answers the user's "also glm, and deepseek latest versions" priority. Headline: **`Glm4MoeLiteForCausalLM` / `zai-org/GLM-4.7-Flash` (31.2B, 58.2 GiB bf16) FITS GB10 and is a SECOND MLA gate vehicle that closes BOTH coverage gaps the MLA campaign named as unit-gated-only** (it has `q_lora_rank=768` and `noaux_tc`/`e_score_correction_bias`, which DeepSeek-V2-Lite lacks). **DSA is DOUBLY BLOCKED on GB10:** the sparse XOR filter eliminates `TRITON_MLA` for sparse models leaving `FLASHINFER_MLA_SPARSE_SM120` as the sole candidate, and that path is non-functional on flashinfer 0.6.12 (XQA backend is dense-only, discards `sparse_mla_top_k`); GLM-5 is 1404 GiB and V3.2 642 GiB regardless. DeepSeek-V4 is a NEW architecture (Sinkhorn-normalized Manifold Hyper-Connections, CSA/HCA compressor with recurrent state, hash-routed MoE) and HW-BLOCKED at 148.7 GiB — but its TOKENIZER risk is LOW (standard HF fast BPE; only the chat template needs porting, with upstream golden fixtures). Glm4/Glm need two primitives we have NONE of: partial rotary factor and sandwich norms. Nothing implemented. **NVFP4A16 (W4A16)** on the already-done dense `Qwen3ForCausalLM` (`RedHatAI/Qwen3-32B-NVFP4A16`, 64L) — the QUANT-SCHEME additivity experiment, serving user priorities #2 (models) and #4 (quants) at once. KERNEL LAYER FULLY ADDITIVE (ZERO new kernel code: vLLM FORCES Marlin for `use_a16`, OBSERVED `Using MarlinNvFp4LinearKernel`, and that is the GEMM we already vendored for the 35B). **CORRECTNESS CLOSED 2026-07-21 (W4b):** the strict gate's 4/6 was diagnosed by the ratified TEACHER-FORCING isolation — all 29 divergent positions gap <= 0.0625 nats with 28/29 EXACTLY 0.0, one root flip an EXACT bf16 tie at which vLLM's teacher-forced argmax is OURS and vLLM contradicts its own greedy. **NOT a W4A16 defect: it is the pre-existing dense-forward bf16 near-tie drift, recorded against `MODEL-TEXT-qwen3-qwen3-for-causal-lm`.** Gate closes **6/6** under the ratified near-tie-robust bar with the nats evidence committed. SPEED still pending ⇒ `ACTIVE`, not `DONE`. **GEMMA FAMILY SPIKED 2026-07-24** ([sweep-gemma](specs/sweep-gemma.md), `CLAIM-SWEEP-GEMMA`): four rows `INVENTORIED` → `SPIKE` (Gemma 1/2/3/4). Answers the user's "and then we do gemma" ("gemma 4") next-target. **The newest registered Gemma is Gemma 4** (real, public, but ALL checkpoints multimodal-wrapped `Gemma4*ForConditionalGeneration`, ≥12B, 0.25.0 oracle-support unverified, needs a PLE/YOCO/MoE/k_eq_v stack) — it leads the characterization but is gate-BLOCKED as a first vehicle. **The recent-first gate vehicle that FITS + is oracle-certain is Gemma 3** (`Gemma3ForCausalLM` on `google/gemma-3-1b-it`). Headline: Gemma reduces MOSTLY to landed infra — gemma-RMSNorm, sandwich norms (glm4 `b568d20`), SentencePiece (names "Gemma"), sliding-window (FA-2 + SlidingWindow/ChunkedLocalAttention specs), the `kAttnQkNormRopeGate` QK-norm+rope recipe, tied embeddings are ALL REUSE; the one genuinely-new compute kernel is GeGLU (`gelu_pytorch_tanh`+mul; we have only SiLU), plus the final logit soft-cap + qpas/embed-scale scalars + dual-rope routing. Per-version delta: Gemma-2 has an attn logit soft-cap, Gemma-3 removed it for QK-norm. **GEMMA-3 W0-W2 LANDED 2026-07-24 — CORRECTNESS COMPLETE, the FIRST Gemma family** (`MODEL-TEXT-gemma3-gemma3-for-causal-lm` `ACTIVE`, speed pending): `Gemma3ForCausalLM` on `google/gemma-3-1b-it`. W1 = two additive default-inert vt ops `kGeluAndMul` (GeGLU `gelu_pytorch_tanh`) + `kMulScalar` (bf16 embed-scale), CUDA+CPU, unit 12/12. W2 = `gemma3.{h,cpp}`/`gemma3_weights.cpp`/`gemma3_registry.cpp` reusing the GLM-4 sandwich-norm layout + `dense_attn_block.h` + FA-only KV: GemmaRMSNorm `(1+w)`, per-head Gemma q/k norm, dual per-layer RoPE theta, qpas scale, per-layer sliding window, GeGLU MLP, `sqrt(hidden)` embed-scale, tied lm_head. **SACRED gate STRICT token-exact 48/48** greedy vs vLLM 0.25.0 (K=5 ALL-DETERMINISTIC → STRICT; BOS-verified; tokenizer-free like Mistral's `LOAD-SENTENCEPIECE` path). Loader 340 tensors, registry 23/23, clean `-Werror` 0 warn. **GEMMA-2 + GEMMA-1 W3-W6 LANDED 2026-07-24 — CORRECTNESS COMPLETE** (`MODEL-TEXT-gemma2-gemma2-for-causal-lm` + `MODEL-TEXT-gemma-gemma-for-causal-lm` `ACTIVE`, speed pending): W3 = the logit soft-cap primitives (`vt::SoftCap` final cap + `PagedAttentionArgs.logits_soft_cap` attention cap threaded into the native/CPU/FA-2 attention, default-inert). W4 `Gemma2ForCausalLM` (gemma-2-2b-it) = the inverse of Gemma-3 (BOTH soft-caps, NO QK-norm, single rope) — **near-tie-band SACRED 48/48** (44/48 strict + 4/48 at 0.0-nat vLLM-own ties, 0 forward-divergent; soft-cap PROVEN applied by a cap-on≠cap-off A/B). W5 `GemmaForCausalLM` (gemma-2b) = the original Gemma (two fused norms, head_dim scale) — **STRICT 48/48**. W6 = Gemma-4 honesty pass (HW/DEP-BLOCKED, not registered). Regressions byte-identical (Gemma-3 48/48, Qwen3-dense 184/184, OPT 63/63, Llama 92/92, Mistral 92/92) + DeepSeek-V2 asserts-on 223/223; compute-sanitizer 0; clean `-Werror` 0 warn. Gemma-4 stays `BLOCKED`. | `PARTIAL` | **ACTIVE: the first additive-model bring-up = Qwen3 dense (`Qwen3ForCausalLM`) on `Qwen3-0.6B` BF16 — W0+W1 landed 2026-07-20.** W0 (config+registry stub: new TU `qwen3_dense.cpp`+`qwen3.h`, one `REGISTER_VLLM_MODEL`, full-attention-only KV spec, forward stub) + W1 (the RUNNER GENERALIZATION `ENG-RUNNER-MODELSHAPE`) are DONE and gated: dgx CUDA `-Werror` 0-warn, **27B 235/235 + 35B 315/315 token-exact UNCHANGED** (behaviour-preserving), new CPU runner tests RED(SIGSEGV)→GREEN, registry resolves `Qwen3ForCausalLM`, ASan/UBSan clean. The runner is now MODEL-SHAPE-AGNOSTIC (a full-attention-only KV config allocates+steps without the hybrid GDN path) → every future dense/non-hybrid arch adds new-files-only. Qwen3-0.6B is the only standard-dense arch with a checkpoint + runnable vLLM 0.25.0 oracle on dgx TODAY (no Llama/Mistral checkpoint present → Llama-first needs a download, sequenced as W-next for genuine cross-family additivity). **W2 loader + W3 forward LANDED 2026-07-20:** dense forward `qwen3.cpp` (`Qwen3DenseModel::Forward/ForwardDevice`) composed from vt:: ops + the fusion catalog (2 new byte-exact recipes: `kFusedAddRmsNormStd`, `kAttnQkNormRope`); bf16 attention numerics mirror vLLM. The first pure-dense bf16 model forced out + FIXED 2 genuine latent bugs: tokenizer `kQwen2Classic` (classic Qwen2/Qwen3 pre-tokenizer was hard-rejected) and `cuda_paged_attn.cu` WMMA prefill mistokenizing at head_dim≠256 (now gated to the validated d=256). **W4 CORRECTNESS COMPLETE 2026-07-20 — near-tie-robust gate PASSES on 0.6B AND a bigger 4B.** The 2026-07-20 razor's "vLLM greedy non-deterministic" premise was a BATCHING artifact: per-prompt (batch=1, the gate regime) vLLM 0.25.0 greedy is DETERMINISTIC (0.6B 0-multi/K=10, 4B 0-multi/K=5). Forward correctness is PROVEN by teacher-forcing vLLM on OUR exact prefix (`scripts/qwen3-neartie-gap.py`): at all-but-2 positions vLLM's own argmax given our prefix IS our token (gap 0.0000, bit-identical logprobs — our forward matches vLLM's prefill logits); residual flips are bf16 near-ties (0.6B ≤0.125 nats, 4B ≤0.25) where vLLM's own prefill argmax disagrees with its decode. Gate = our token within 0.5 nats of vLLM's teacher-forced argmax (strict where equal): **Qwen3-0.6B 16/16** (strict 12 + near-tie 4) and the **bigger-model complete-correctness proof Qwen3-4B (36L, GQA 32/8, hidden 2560, same forward code) 16/16** (strict 10 + near-tie 6). Regression 27B 235/235 + 35B 315/315 UNCHANGED, CUDA `-Werror` 0-warn, memcheck 0. Correctness-complete. **SPEED — d128 FA2 PREFILL + DECODE DEFAULT-ON 2026-07-20 (`Qwen3-4B` vs vLLM 0.25.0 production/graphed, in1024/out128) — big gap-close, still below vLLM, `MODEL-TEXT-qwen3-qwen3-for-causal-lm` stays `ACTIVE` NOT `DONE`:** implemented the dominant prefill lever (a d128 FlashAttention-2 varlen prefill — generalized the vendored FA2 launcher to head_dim 128, `VT_FA2_PREFILL_QWEN3` default-ON) and flipped the FA2 varlen d128 decode default ON (near-tie gate re-passes 16/16 on 0.6B + 4B). Total tput now 0.90× (c1)/0.62× (c8) (up from 0.80×/0.48×), c1 decode at parity (TPOT 1.04×, ITL P99 0.98× win); prefill A/B = +7%/+41% total, −55%/−48% TTFT. STILL failing TTFT median 5.85×/10.2× + total <1×: the full prefill STEP (not the attention kernel, now vLLM's FA2 family) is ~6× vLLM = non-attention glue (GEMM/MLP fusion) + host-side launch overhead (un-graphed prefill); plus c8 decode batch efficiency (TPOT 1.38×). Dominant residual lever = portable prefill-step fusion + graphed prefill (secondary = c8 split-KV decode occupancy). **RoPE cos/sin cache flipped DEFAULT-ON 2026-07-20** (`VT_QWEN3_ROPE_CACHE`): the opt-in blocker (an alleged FA2-split-KV-combine run-to-run nondeterminism) was GROUNDED + DISPROVEN — the paged engine is byte-deterministic run-to-run and the combine never launches on the gate (`num_splits==1`); goldens regenerated on the canonical `$HOME/cutlass-4.5.0` build (the flashinfer cutlass copy tips the 27B tok6 razor to 234/235; cutlass-4.5.0 = 235/235), gate 16/16 both, 27B 235/235 + 35B 315/315 unchanged. RoPE-ON closes total tput 0.90×→0.97× (c1) / 0.62×→0.82× (c8) and c1 TTFT ratio 5.85×→2.27×, still `ACTIVE`. **SPEED RE-BOUND 2026-07-21 (same-session, matching-recipe) — TTFT residual RESOLVED, cutlass claim CORRECTED:** the "TTFT 2.27×/5.85×" + "c8 ITL 4.3×" were BAD-DENOMINATOR/num-prompts artifacts — a fresh same-session vLLM capture gives c1 TTFT ~152 ms & c8 ITL P99 ~130 ms, and OURS WINS TTFT at both concurrencies (c1 0.90×, c8 0.38×). **c1 = effective every-axis parity** (tput 0.98× / TPOT 1.01× / TTFT+ITL wins); **c8 residual = decode** (tput 0.93× / TPOT 1.10× / ITL P99 1.12×), which nsys shows is 93% GPU-busy/compute-bound (small-M=8 `cutlass_80_wmma` projections). The **qkv-merge** (new GQA `QkvSplit` op mirroring vLLM `QKVParallelLinear`) was implemented + MEASURED NEUTRAL (doesn't cut decode FLOPs) ⇒ default-OFF. **CUTLASS CLAIM CORRECTED: 27B `test_qwen27_paged_engine` = 235/235 on the FLASHINFER cutlass build** (the "flashinfer ⇒ 234/235" was a build artifact). Stays `ACTIVE`; named residual = c8 decode-GEMM efficiency (a decode-fusion sub-campaign). **SWEEP MODEL #1 — Qwen3-Coder-30B-A3B (`Qwen3MoeForCausalLM`) W0+W1 LANDED 2026-07-21** ([sweep-qwen3-coder-30b.md](specs/sweep-qwen3-coder-30b.md), `MODEL-TEXT-qwen3-moe-qwen3-moe-for-causal-lm` `ACTIVE`, `CLAIM-MODEL-QWEN3-CODER`): the first full-attention BF16 MoE, composed from the done dense attention + the done 35B MoE experts (ZERO runner change). W0 = registry stub (`qwen3_moe_registry.cpp` + `qwen3_moe.h`, one `REGISTER_VLLM_MODEL`, full-attn-only KV, `is_dense_model=false`, W2/W3 throwing stubs). W1 = three behaviour-preserving refactors making the two done pieces reusable: (#1) dense `AttnBlock` + glue extracted to `dense_attn_block.h` (Qwen3-dense byte-identical), (#2) bf16 `MoeBlock` exposed cross-TU via `RunMoeBlock` (`qwen3_5_moe_block.h`; 35B untouched), (#3) no-shared-expert guard (inert for the 35B). Gated: dgx CUDA `-Werror` 0-warn; Qwen3-dense 0.6B+4B 16/16 + 27B 235/235 + 35B 315/315 UNCHANGED; registry resolves `Qwen3MoeForCausalLM`; memcheck 0. W2 bf16 loader → W3 forward → W4 near-tie token-exact → W5 fast bf16 grouped-MoE GEMM remain. Then Llama dense (download), Mistral, MoE families **SWEEP MODEL — GLM-4-9B-0414 (`Glm4ForCausalLM`) G2 LANDED 2026-07-24 — CORRECTNESS COMPLETE** ([glm-dsa-latest-deepseek](specs/glm-dsa-latest-deepseek.md), `MODEL-TEXT-glm4-glm4-for-causal-lm` now `READY` per the [live-state audit](specs/live-state-audit-2026-08-06.md), `CLAIM-GLM-DSA-LATEST-DEEPSEEK` amended) **GLM-4.7-Flash (`Glm4MoeLiteForCausalLM`, 31.2B MLA+MoE) G1 LANDED 2026-07-24 — SACRED gate 8/8, CORRECTNESS COMPLETE** (`MODEL-TEXT-glm4-moe-lite-glm4-moe-lite-for-causal-lm` `ACTIVE`, speed pending): reuses the DeepSeek-V2 MLA stack; first e2e coverage of the q_lora branch + noaux_tc sigmoid router, closing the MLA campaign's two C2 coverage gaps: the FIRST GLM-family model. SACRED gate 16/16 vs vLLM 0.25.0 (STRICT 13/16 + near-tie 3/16, max gap 0 nats; vLLM K=5 self-deterministic ⇒ STRICT bar), speed PENDING. The two "new primitives" the spike named reduced to EXISTING infra: partial + interleaved `RopeFromCache` (`is_neox_style=false`, the DeepSeek-V2 decoupled-rope path) over `rotary_dim=64`; standalone `vt::RmsNorm` sandwich norms. Biased qkv, no QK-norm, GQA 32/2, untied lm_head. New files + one REGISTER, reusing the shared dense glue. **SWEEP MODEL — Llama-3.2 (`LlamaForCausalLM`) W0-W4 LANDED 2026-07-23 — CORRECTNESS COMPLETE** ([sweep-llama-3.2](specs/sweep-llama-3.2.md), `MODEL-TEXT-llama-llama-for-causal-lm` `ACTIVE`, `CLAIM-MODEL-LLAMA-3.2`): the roadmap's explicit "Llama-first" increment and the first mainstream non-Qwen/non-OPT dense arch. `LlamaForCausalLM` (Llama-3.2-1B) = the Qwen3-dense forward with exactly two ADDITIVE deltas — NO qk-norm (shared `AttnBlock` skips it when q_norm/k_norm empty) + llama3 rope-scaling (4 default-0 `RopeArgs` fields + a `Llama3ScaleFreq` kernel helper, no-op elsewhere; formula verified 2e-7 rel vs vLLM) — reusing the shared dense forward VERBATIM (`LlamaModel == Qwen3DenseModel`). 3 new Llama files, ZERO edit to runner/scheduler/platforms/attn-registry/`hf_config`/any qwen3-opt model. vLLM 0.25.0 greedy MEASURED DETERMINISTIC (K=6, 0 multi-valued cells) ⇒ STRICT token-exact bar, PASS **16/16 (12 strict + 4 near-tie band, max gap 0.0000 nats, 0 divergent)** — at all 13 divergent positions vLLM's own teacher-forced argmax given our prefix IS our token. A correctness-fatal tokenizer bug (Llama's `Sequence` post_processor wrapping `TemplateProcessing` ⇒ BOS 128000 never prepended, silently 1/16) was isolated by a CUDA prefill-argmax diagnostic (forward proven 4/4 correct given vLLM's exact tokens) and fixed byte-preservingly (Qwen/OPT/DeepSeek unaffected — all ByteLevel/top-level-TemplateProcessing). Regressions 27B 235/235 · 35B 315/315 · Coder 6/6 · Qwen3-dense 16/16 · OPT 6/6 · DeepSeek-V2 8/8 UNCHANGED; `-Werror` 0-warn; memcheck 0; DSR 67. SPEED pending (head_dim 64 → generic paged path; Llama-3.2-3B head_dim-128 is the FA2-toggle W-next). **MLA CAMPAIGN W0+W1 LANDED 2026-07-21** (`CLAIM-MLA-DEEPSEEK`; rows STAY `SPIKE` — W0/W1 make no model supported). **W0 grounded every fact the spike flagged as an unverified source read; ALL CONFIRMED, none contradicted:** DeepSeek-V2-Lite fetched to dgx (30 GB, 4 shards) and loading in the vLLM 0.25.0 oracle; the oracle's own DEBUG startup on sm_121 prints `Using TRITON_MLA attention backend out of potential backends: ['TRITON_MLA']` and `Using FLASH_ATTN MLA prefill backend` — so the dense-MLA decode + MLA-prefill targets are OBSERVED, not inferred, and the sm90/sm100-only MLA kernel class stays out of scope; the real `config.json` confirms every §5.1 number (`kv_lora_rank=512`, `qk_nope=128`, `qk_rope=64` -> the **576-wide latent**, `v_head_dim=128`, `q_lora_rank=null`, `n_group=topk_group=1`, softmax/greedy, 64+2 experts, 27 layers) plus `is_neox_style=False` and the mscale2 scale correction; and BOTH recorded coverage gaps (no `fused_qkv_a_proj` branch, no `e_score_correction_bias`) are confirmed real, so those pieces stay unit-gated only. **W1 = the behaviour-preserving spec-driven KV allocation, ZERO MLA math:** the attention cache is now sized `num_blocks * spec->page_size_bytes()` and viewed from the spec's own `block_size`/`num_kv_heads`/`head_size`/`dtype` instead of the hardcoded `2 * block * Hkv * Dh` with shape reconstructed from the HF config (`runner.cpp`), plus `MLAAttentionSpec` with upstream's factor-1 single-tensor page formula (`kv_cache_interface.py:397-398`) registered against the ORDINARY `FullAttentionManager` (`single_type_kv_cache_manager.py:1539`) — the spike's key finding, which is why block manager/prefix caching/eviction need no change. Gated: dgx clean CUDA `-Werror` 0 warnings/0 errors; **27B 235/235 + 35B 315/315 + Qwen3-Coder 6/6 + Qwen3-dense 16/16 ALL UNCHANGED** (behaviour-preserving proven, not assumed); `test_runner` 15/15, `test_kv_cache_interface` 21/21 (4 new MLA-spec cases), `test_llm_engine` 5/5; the new path is proven EXERCISED (not merely compiled) by `fa_page_size_bytes()` + a `page_size_padded` case no HF-config formula can produce. **W2 + W3 LANDED 2026-07-21** (base `a05437f`; rows STAY `SPIKE` — still no MLA attention math, no MLA model, no forward). **W2 = the MLA branch of `_get_backend_priorities` the pre-W2 comment deferred, ported as DATA:** the whole of `cuda.py:84-176` (BOTH branches — MLA sm_100 including the `:96-115` adaptive sparse tail, MLA sm_12x, MLA `else`, and the two pre-existing non-MLA arms) is now a TABLE in the new header `include/vllm/platforms/cuda_attn_priority.h`, one row per upstream arch arm keyed on `(use_mla, major)`, so a future arch is a ROW rather than a code path; putting it in a header (not the CUDA-only TU) also let the CPU test tier assert the REAL table and DELETED the hand-copied `FakeCudaPlatform` duplicate. On sm_121 a `use_mla=true` request now RESOLVES to `TRITON_MLA`, matching the W0 oracle observation. **The sparse/DSA seam is left OPEN and unit-proven:** GB10's row keeps both upstream entries and the sparse one loses to a real FILTER — `AttentionBackend::is_mla()`/`is_sparse()` checked against the request (`backend.py:307-360 validate_configuration`) — so a future DSA backend is selected purely by declaring `is_sparse() == true`, with ZERO edit to the table or the selector. `TritonMLABackend` lands the NAME plus upstream's 3-D `get_kv_cache_shape` (no K/V axis; `num_kv_heads != 1` REFUSED), `get_impl_cls()` deliberately still `nullptr`. MLA prefill priority ported too (GB10 -> `[FLASH_ATTN]` alone). **W3 = the two new `vt::` ops, both CPU-reference-gated.** `vt::ConcatAndCacheMla` mirrors `csrc/libtorch_stable/cache_kernels.cu:401-442` — and per the whole-chain rule this was VERIFIED, not assumed, to be vLLM's OWN csrc kernel (`_custom_ops.py:2532` -> `torch.ops._C_cache_ops`), with no flashinfer/cutlass variant in the dense-bf16 path; it concatenates the latent + rope part into ONE 576-wide entry, the write `ReshapeAndCache`'s K/V-pair signature cannot express, stride-driven so a per-layer cache slice and the two column halves of `kv_a_proj_with_mqa` both work copy-free. The **grouped-topk (`noaux_tc`) router** extension — flagged in `coordination.md` as SHARED with `CLAIM-GLM-DSA-LATEST-DEEPSEEK` and "must not be implemented twice" — is landed HERE and that claim now consumes it: additive `MoeRouterTopKArgs` fields + an optional `e_score_correction_bias`, with `num_expert_group == 0` still dispatching the ORIGINAL kernel so the 27B/35B/Coder/dense routers are byte-identical BY CONSTRUCTION. **Stated plainly: the `noaux_tc` correctness evidence is UNIT-ONLY.** V2-Lite has `n_group=topk_group=1` and no bias, so the e2e vehicle exercises none of it; the gate is `tests/vt/test_ops_moe_router_grouped.cpp` at DeepSeek-V3's REAL dimensions (256 experts, n_group=8, topk_group=4, sigmoid, scaling 2.5, WITH the bias) against an INDEPENDENT sort-based transcription of the upstream formula. **W4 LANDED 2026-07-22** (base `ed2c342`; rows STAY `SPIKE` — W4 adds a kernel and fills a `nullptr`, it makes no model supported). **`vt::MlaDecodeAttention` — the MQA decode over the compressed latent (QK 576 / V 512, `num_kv_heads=1`), a structure port of the two-stage split-KV pair W0 OBSERVED EXECUTING:** `MlaDecodeStage1` <- `_fwd_grouped_kernel_stage1` (`triton_decode_attention.py:278-458`, the `IS_MLA` branch whose `v = tl.trans(k)` at `:424-431` is the whole MLA trick — V is the leading 512 columns of the SAME latent row already loaded as K, so one shared-memory tile serves as both), `MlaDecodeStage2` <- `_fwd_kernel_stage2` (`:575-639`), `ComputeNumKvSplits` <- `_compute_num_kv_splits` (`triton_mla.py:40-47`), split workspace <- `_reserve_attn_logits_workspace` (`:57-78`) realized as the house grow-only per-stream scratch. **Honest reuse verdict:** our FA-2 split+combine machinery fit at the ALGORITHM level (the split schedule, the LSE merge algebra, the fixed-ascending no-atomicAdd determinism rule) and NOT at the code level — the vendored FA-2 launcher takes separate 4-D k/v caches and is instantiated for symmetric head_dim {128,256}, which cannot express a 3-D single-buffer cache with QK 576 / V 512; that is recorded in the TU header rather than forced. **Evidence is unit-level and deliberately strong** (there is no e2e model until W7): [`tests/vt/test_ops_mla_attn.cpp`](../tests/vt/test_ops_mla_attn.cpp), a port of `tests/kernels/attention/test_mla_decode_cpu.py` whose `ref_mla` becomes an INDEPENDENT TWO-PASS oracle (a different algorithm from the streaming online-softmax both impls use) plus its NaN-padding out-of-bounds detector, run at the REAL V2-Lite geometry (576/512/64, block 16, mscale^2 scale) over ragged / multi-block / single-block / every split boundary (`num_kv_splits` ∈ {1..512} incl. splits > seq_len) / 128-head V3 / non-BLOCK_H head counts / a 288-256 block-32 geometry / bf16 + f32, with run-to-run BIT-exactness. dgx sm_121: 11/11 cases, 2,303,193 assertions; `compute-sanitizer` memcheck **0 errors**, racecheck **0 hazards**, synccheck **0 errors**; clean CUDA build **0 warnings / 0 errors**; **27B 235/235 + 35B 315/315 + Coder 6/6 + Qwen3-dense 16/16 + OPT 6/6 ALL UNCHANGED**. `TritonMLABackend::get_impl_cls()` is no longer `nullptr` — it returns a real `TritonMLAImpl` whose `forward_mqa` is the 1:1 counterpart of `triton_mla.py:189-260`; PREFILL is W5 and `forward()` refuses a prefill-shaped batch BY NAME rather than producing wrong numbers. NO speed number — decode perf is W9. **W5 LANDED 2026-07-22** (base `5395203`; rows STAY `SPIKE`). **MLA PREFILL + the workspace-bounded CHUNKED-CONTEXT loop.** Three new ops — `vt::MlaPrefillAttention` (<- `mla/prefill/flash_attn.py:153-248`, the ONLY MLA prefill backend reachable on sm_121 and the one W0 OBSERVED the oracle logging), `vt::GatherMlaCache` (<- `cache_kernels.cu:992-1064`) and `vt::MergeAttnStates` (<- `merge_attn_states.cu:18-192`, both `-inf` edge cases verbatim) — plus the loop itself in the new `mla_chunked_context.h` (<- `mla_attention.py:1422-1451,1667-1745,2094-2199,2344-2425`), which is what keeps a long-context prefill inside a bounded workspace instead of materializing a 3 GB up-projected context. **The vendored FA-2 launcher WAS generalized, and W4's prediction that it would be tractable held for a reason worth recording: upstream does not ask FA-2 for asymmetric head dims either.** `requires_v_padding` is TRUE on GB10, so upstream ZERO-PADS V from 128 to 192 and slices the output back — the kernel stays a plain SYMMETRIC head_dim-192 instantiation. The whole change is two new explicit instantiations of an UNCHANGED generic template, one new launcher entry for the contiguous-varlen mode, and the pad/slice pair; the paged launcher every non-MLA prefill calls is textually untouched (211 insertions / **0 deletions**), and 27B 235/235 + 35B 315/315 + Coder 138/138 + Qwen3-dense 664/664 + OPT 36/36 are all UNCHANGED. Evidence is UNIT-ONLY (there is still no model): 4/4 cases / **2,377,052 assertions** and 5/5 / **306,037 assertions** on dgx sm_121 at the real QK 192 / V 128 geometry, against an INDEPENDENT double-precision two-pass oracle and — for the loop — a SINGLE-SHOT whole-sequence oracle that never chunks, over exact / +1 / -1 chunk boundaries, zero-context and zero-key-in-chunk requests, ADVERSARIAL reverse-interleaved block tables, NaN-poisoned outputs and run-to-run bit-exactness; memcheck **0**, racecheck **0 hazards**, synccheck **0**. A genuine upstream FA-2 quirk was found on the way and worked around rather than papered over: its EMPTY-K early exit ignores the unpadded-LSE flag, which a zero-key chunk request would turn into an out-of-bounds LSE write. **W6 LANDED 2026-07-22** (base `2846467`; rows STAY `SPIKE` — W6 adds an attention LAYER, not a model). **The MLA attention BLOCK + LOAD-TIME WEIGHT ABSORPTION — the piece that finally COMPOSES W3's cache write, W4's MQA decode and W5's MHA prefill into one layer:** the projections with BOTH `q_lora_rank` branches (`fused_qkv_a_proj` -> `q_a_layernorm` -> `q_b_proj`, or the direct `q_proj`), the two RMSNorms (the rope part deliberately NOT normed), the DECOUPLED RoPE (`is_neox_style=False`, only the trailing 64-dim slice rotates) with its YaRN cos/sin cache and the SEPARATE mscale^2 softmax-scale correction, the `kv_b_proj -> W_UK/W_UV` split, the prefill-MHA / decode-MQA dispatch with decode tokens packed FIRST, and the `kv_b_proj` up-projection callback W5 left open. **The spike's most useful prediction held: absorption needed NO new attention kernel** — it is a LOAD-TIME weight transform plus TWO batched GEMMs, so the entire new-kernel surface is two general primitives, `vt::BatchedMatmul` (<- `torch.bmm` at `mla_attention.py:789,1034`; on CUDA torch resolves that to cuBLAS `gemmStridedBatchedEx`, and ours is the cuBLASLt strided-batched form of the same GEMM) and `vt::ConcatMlaNopeRope` (<- `concat_mla_q`, generalized so one op also serves `_concat_k_nope_k_pe`). **The absorbed-vs-unabsorbed equivalence — the heart of W6 — is PROVEN NUMERICALLY, three independent ways, rather than argued:** an INDEPENDENT double-precision block oracle computes the attention BOTH ways and agrees to < 1e-11 (the identity itself); our absorbed decode reproduces the UNABSORBED oracle to < 2e-4 in f32; and — the strongest — the SAME batch is driven once through our ABSORBED MQA decode kernel (QK 576 / V 512, one KV head, K/V never materialized) and once through our UNABSORBED materialized-MHA prefill path (QK 192 / V 128 plus the chunked-context loop), agreeing to < 3e-4 (CPU f32) / < 4e-2 (CUDA bf16) with nothing but the weights shared between them. Evidence on dgx sm_121: `test_mla_attention_block.cpp` 10/10 cases / 2,372,644 assertions and `test_ops_mla_absorb.cpp` 9/9 / 1,644,807 (CUDA cases proven to EXECUTE; NaN-poisoned outputs; run-to-run BIT-exact), porting `tests/kernels/test_concat_mla_q.py` in both arms. memcheck 0, racecheck 0 hazards, synccheck 0 (the last needing `--num-cuda-barriers 65536` — the default table overflows on a binary driving this many kernel families and the tool then reports a bogus launch failure, a trap worth knowing). Clean CUDA build 0 warn / 0 err; **27B 235/235 + 35B 315/315 + Coder 138/138 + Qwen3-dense 664/664 + OPT 36/36 ALL UNCHANGED**. **Coverage stated plainly: the `q_lora` query branch has NO e2e coverage and cannot get any on GB10** — DeepSeek-V2-Lite has `q_lora_rank=null`, so it is unit-gated at DeepSeek-V3's real dimensions only; GLM-4.7-Flash (`q_lora_rank=768`, 58.2 GiB, fits) is what would close it. **W7 LANDED 2026-07-22** (base `ce43c51`; the row STILL stays `SPIKE`). **THE DEEPSEEK-V2 MODEL — registry + config parse + loader + forward: the first MLA model in the tree, and the first one that runs a real MLA checkpoint end to end.** Four new files plus ONE shared-code edit (a two-line additive condition in `runner.cpp` recognising a `kMlaAttention` KV group as the model's attention group — upstream maps MLA onto the ordinary `FullAttentionManager`, so block tables/prefix caching/eviction are untouched). **LOADER GATE PASSED on the real 4-shard DeepSeek-V2-Lite: 5291/5291 checkpoint tensors accounted for, ZERO unmapped and ZERO leftover** (4/4 cases / 37,331 assertions), every shape asserted including the LOAD-TIME `kv_b_proj -> W_UK_T [16,128,512]` / `W_UV [16,512,128]` absorption split — the same transform, at the same lifecycle point, as upstream's `process_weights_after_loading`. **V2-Lite takes the DIRECT `q_proj` query branch** (`q_lora_rank: null`), asserted with the fused branch EMPTY on every layer. **FORWARD GATE PASSED and obviously right, not merely finite: the real checkpoint prefill of `The capital of France is` -> argmax ` Paris`** (top-5 ` Paris`/` the`/` a`/` one`/` also`, run-to-run bit-exact) — the direct analogue of the Qwen3-Coder W3 sanity case. **BATCH-ORDERING GATE:** the ordering invariant W6 measured 0.86 relative error from is now VALIDATED, not assumed — `BuildMlaBatchSplit` throws (naming the request and citing the upstream line) if a decode follows a prefill or a with-context prefill follows a context-free one. **SHARED EXPERTS — new for this family and UNGATED unlike Qwen3.6's sigmoid-gated one — gated two ways:** a MoE layer with every routed expert zeroed is BIT-IDENTICAL to a dense layer holding the same MLP, and turning the shared expert off CHANGES the logits. **The CUDA path is EXERCISED, not merely compiled:** a case at the real MLA head geometry drives the CUDA MLA kernels and the CUDA-only grouped bf16 MoE GEMM, bit-exact on device and within 0.0061 worst relative logit error of the CPU reference path. 11/11 forward cases; memcheck/racecheck/synccheck all **0**; clean CUDA build **0 warn / 0 err**; **regression set UNCHANGED**. **Only `DeepseekV2ForCausalLM` is REGISTERED** — `DeepseekForCausalLM` (plain MHA), V3 (fp8/671B) and V3.2 (DSA indexer) are REFUSED BY NAME in the config parse rather than falsely claimed. A pre-existing tree-wide hazard was found on the way and recorded: the shared `DevicePool` is a process-wide singleton keyed only on a byte size class, so a single process driving BOTH a CPU and a CUDA forward hands the second backend the first's recycled pointers. **NEXT: W8 — the SACRED token-exact gate on DeepSeek-V2-Lite** (wire the paged engine to produce the MLA batch order the model already validates, capture oracle goldens, run the STRICT form W0 determined). A loading, forwarding model is NOT a supported model, so no model row moves until that gate passes. **W8 LANDED 2026-07-22 — THE SACRED CORRECTNESS GATE PASSES 8/8, and `MODEL-TEXT-deepseek-v2-deepseek-v2-for-causal-lm` moves `SPIKE` -> `ACTIVE` (correctness COMPLETE, speed PENDING). NOT `DONE` — that additionally requires vLLM-speed parity on every axis, which is W9 and has NO number yet; the other four campaign rows stay `SPIKE`.** An 8-prompt battery is driven through the FULL paged `LLMEngine` and compared to the pinned vLLM 0.25.0 oracle: **8/8 PASS — STRICT token-exact 5/8, near-tie band 3/8, 92/128 tokens strictly exact, max teacher-forced gap 0.25 nats, 0 forward-divergent** (223 assertions). **The bar was ARRIVED AT by measurement, not chosen:** vLLM is DETERMINISTIC on this model at batch=1 (W0's K=5 8/8, re-confirmed by W8's own capture at T=16 with 0 multi-valued cells), so the STRICT form ran FIRST and came out 5/8; the ratified TEACHER-FORCING diagnostic then showed **36 divergent positions with 35 at gap EXACTLY 0.0000 nats** — vLLM's own argmax GIVEN OUR PREFIX is our token, so they are the downstream tail of one earlier flip — **exactly ONE root flip with any gap at all (prompt[3] tok 9, 0.2500 nats, inside the ratified 0.5-nat band and equal to the landed Qwen3-dense 4B gate's worst)**, and **ZERO tokens outside vLLM's top-20**, with the per-position nats COMMITTED as goldens and anything beyond the band still FAILING. **W8's first job — the scheduler/runner wiring — turned out to need NO new code, for a non-accidental reason:** `runner.cpp:671` already reorders with `decode_threshold = 1`, exactly MLA's `reorder_batch_threshold` (`mla_attention.py:1420`), and its `decode -> short_extend -> long_extend -> pure_prefill` ordering satisfies BOTH MLA invariants (decodes form a batch prefix; with-context prefills lead the prefill tail). W8 PROVES that end to end rather than duplicating it, with new DIAGNOSTIC `MlaBatchSplitStats` counters and a non-vacuity bar: the battery is admitted CONCURRENTLY with staggered arrival, producing **7 genuinely MIXED decode+prefill steps at up to 8 concurrent requests** with `BuildMlaBatchSplit` (which throws naming the request) never firing, plus a prefix-cache-driven **with-context prefill**, and a phase-0 check that the engine really allocated the compact MLA cache (`fa_page_size_bytes = 36864`, no factor 2). **THE REAL BLOCKER WAS THE TOKENIZER, NOT THE MODEL:** the first run REFUSED to load (`unsupported normalizer "Sequence"`), and behind it sat a whole NEW pre-tokenizer family — DeepSeek's is a HF `Sequence` PIPELINE of SEVEN stages (five `Split(Isolated)` over ENUMERATED codepoint ranges, then `Digits(individual_digits=true)`, then `ByteLevel(use_regex=false)`), whose stage ORDER is load-bearing because stage 2's punctuation class spans 0x3A-0x7E and CONTAINS A-Z/a-z. Landed as `SplitPattern::kDeepSeek` with the five patterns compared VERBATIM at load, and MEASURED token-for-token against the REAL HF `tokenizers` library over a stage-stress corpus (**6/6 cases / 2461 assertions**). **The TOKENIZATION goldens earned their keep by REFUTING a fix that was already written:** `tokenizer_config.json` declares `add_bos_token: true`, which reads as exactly the OPT missing-BOS bug — but vLLM's resolved tokenizer (`TokenizersBackend`) adds NO BOS, our loader already matched bit-for-bit, and the "fix" would have BROKEN a passing gate; it was reverted and the measured behaviour PINNED by a guard case ([[ground-premises-before-dispatching]]). Regression set UNCHANGED (27B 235/235, 35B 315/315, Coder 6/6, Qwen3-dense 16/16, OPT 6/6, plus every tokenizer test — W8 touches SHARED tokenizer code, so that was proved, not assumed); clean CUDA rebuild 0 warn/0 err; local CPU suite 151/151; memcheck/racecheck/synccheck 0. Batch invariance is REPORTED (6/8) and deliberately NOT a bar, because the ORACLE itself changed on 3/8 of this battery under batched generation (W0). One W9 input recorded: the oracle must run `moe_backend='triton'` — vLLM's auto-selected FlashInfer CUTLASS unquantized MoE REBOOTED dgx three times on GB10's unified memory. **W9 SPEED CLOSE LANDED 2026-07-22 — the track has its FIRST binding speed number, and it is an ATTRIBUTED MISS: `MODEL-TEXT-deepseek-v2-deepseek-v2-for-causal-lm` STAYS `ACTIVE` (correctness COMPLETE, speed SHORT), NOT `DONE`** ([grid](../docs/BENCHMARKS.md), [spike §W9](specs/mla-deepseek-campaign.md)). Denominator SETTLED with evidence — CUTLASS MoE has now rebooted dgx **five times** (two more at W9, the second on a pristine box with a 0 GiB page cache and every mitigation applied, both deaths at the identical post-`torch.compile` phase), so `--moe-backend triton` IS vLLM's best STABLE GRAPHED configuration here and is the bar; the substitution does not flatter us, we lose to it. `nsys` (both sides, `--cuda-graph-trace=node`) overrode the plan: the lever was not the planned MLA fusion recipes but `MlaDecodeStage1` sitting at **44.7% of all GPU time and ~180x off its own memory-bound floor** on a **2-CTA grid at batch 1**; applying upstream's own occupancy target made it **18.3x faster** (837 -> 45.8 us) for **+69.5%/+53.3%/+32.0%/+19.5%** end-to-end at c1/c2/c4/c8, while the planned decode-graph sibling is worth only ~+2% (this decode is GPU-bound). Grid vs vLLM: output throughput **0.87/0.95/0.86/0.88** (was 0.50 at c1), TTFT **1.06/1.14/0.96/0.88** (we WIN at c4/c8), TPOT **1.11/0.97/1.16/1.17**. SACRED gate **8/8 UNCHANGED** with both levers default-ON; a real latent CUDA-graph use-after-free in the MLA metadata upload was found and fixed (its whole class now guarded); regression set UNCHANGED; clean rebuild 0 warn/0 err; sanitizers 0. **NEXT LEVER, NAMED: route the batch-1 dense projections off cuBLAS `gemvx` (31.8% of our GPU time) onto a tensor-core GEMM — vLLM splits the same work `gemvx` 12.7% + `nvjet_sm121_tst_mma_*` 6.6%.** **W10 BLOCKED-ROW HONESTY PASS LANDED 2026-07-22 — the campaign's W-plan is COMPLETE; records only (no code, no build, no GPU work, nothing downloaded, no number claimed).** Rows set to their final honest state: `MODEL-TEXT-deepseek-v2-deepseek-v3-for-causal-lm` (V3 + V3.2, and Kimi-K2/K2.5's text backbone by config composition) and `MODEL-TEXT-minimax-m2-mini-max-m2-for-causal-lm` move `SPIKE` -> `BLOCKED`, joined cross-claim by `MODEL-TEXT-deepseek-v2-glm-moe-dsa-for-causal-lm` (GLM-5) under `CLAIM-GLM-DSA-LATEST-DEEPSEEK`; each is HW-BLOCKED on 119 GiB (~642 GiB fp8 / ~428 GiB / 1404 GiB) and the two DSA models are additionally DEP-BLOCKED — for a SPARSE model the XOR filter eliminates `TRITON_MLA`, leaving `FLASHINFER_MLA_SPARSE_SM120` alone, whose sm12x dispatch goes to flashinfer's DENSE-ONLY XQA backend that discards `sparse_mla_top_k` (upstream's own test monkeypatches the probe and asserts nothing numerical). `MODEL-TEXT-deepseek-v2-deepseek-for-causal-lm` stays `SPIKE` with the record repaired to say it is plain MHA and needs NO MLA; Kimi-Linear stays `SPIKE` (MLA half unlocked, KDA a separate kernel campaign, HW-MARGINAL). Each blocked row states what CAN still be gated (config resolution, weight-map on a slice, unit parity at real dimensions) versus what CANNOT (anything e2e). **Two PERMANENT coverage gaps now stated in the rows:** the `noaux_tc` grouped router and the `q_lora` query branch have NO e2e coverage and are unit-gated only, because V2-Lite is `n_group=topk_group=1`/softmax with no `e_score_correction_bias` and `q_lora_rank=null`. **NAMED NEXT VEHICLE: GLM-4.7-Flash** (`MODEL-TEXT-glm4-moe-lite-glm4-moe-lite-for-causal-lm`, 31.2B / 58.2 GiB, FITS GB10) — the only reachable checkpoint that closes BOTH gaps. **BLOCK NOT CLOSEABLE, nothing archived:** the DeepSeek-V2 row is `ACTIVE`, not `DONE`, so the plan/spec stay LIVE; the single open item is the named `gemvx` -> tensor-core dispatch lever. **MISTRAL FIFTH FAMILY W0-W3 LANDED 2026-07-23** ([sweep-mistral](specs/sweep-mistral.md), `MODEL-TEXT-mistral-mistral-for-causal-lm` `ACTIVE`, `CLAIM-MODEL-MISTRAL`): the closest-to-Llama dense arch (vLLM `mistral.py` = "Mistral adaptation of the LLaMA architecture") — plain rope θ1e6 (no rope_scaling) + qk-norm-optional + untied lm_head + null sliding_window, all PRE-EXISTING ⇒ NO new primitive, 3 new files + additive CMake/registry-test rows only, ZERO shared-code edit. **MODEL forward gate 30/30 greedy tokens vs vLLM 0.25.0** (tokenizer-free: fed vLLM's exact prompt tokens through our CUDA prefill; 29 STRICT token-exact + 1 near-tie, 0 forward-divergent; vLLM greedy det 4/5 K=3). W2 loader real-weights 1541 assertions. **REAL FINDING:** Mistral's SentencePiece/Metaspace tokenizer is unsupported by our ByteLevel-BPE tokenizer → the FULL paged-engine SACRED gate is BLOCKED, the pre-inventoried `LOAD-SENTENCEPIECE` row (SentencePiece tokenizer family). `-Werror` 0-warn, DSR 32, regressions UNCHANGED (Llama paged 16/16, Qwen3-dense forward 1031, registry 299; MoE/GDN gates unaffected by construction). SPEED + full paged gate both PENDING (row `ACTIVE`, not `DONE`). **OLMo-2 SPIKED 2026-07-24** ([sweep-olmo2](specs/sweep-olmo2.md), `CLAIM-SWEEP-OLMO2`): one row `INVENTORIED` → `SPIKE` (`MODEL-TEXT-olmo2-olmo2-for-causal-lm`, covering `Olmo2ForCausalLM` + its `Olmo3ForCausalLM` alias). Answers the breadth-sweep §B.3 Tier-2 rank-8 "GLM4 / Olmo2-3" item (GLM-4 + Gemma landed; OLMo-2 next). **HEADLINE: OLMo-2 is the cleanest dense bring-up yet — ZERO new compute kernels.** The two distinctive facts both reduce to WIRING over landed ops: (1) the **pure post-norm (`norm_after`) placement** is a strict SUBSET of the GLM-4/Gemma sandwich (keeps ONLY the standalone-output-norm op `glm4.cpp:174-178` — the exact primitive flagged — DROPS the pre-norms, plain residual add); (2) the **QK-norm is FULL-WIDTH not per-head** → reuses `vt::RmsNorm` at a `[T,q_size]`/`[T,kv_size]` shape but CANNOT use the fused per-head `kAttnQkNormRopeGate`. Everything else REUSES (plain RMSNorm, SiLU SwiGLU, NeoX rope, GQA paged glue, Gemma-3 sliding-window for Olmo-3, tied embeddings, packed loader, ByteLevel BPE). Gate vehicle `allenai/OLMo-2-0425-1B` (1.485B, ~2.77 GiB, fits GB10 tight ~30 GiB free); Olmo-3 rides the same row (0.25.0 oracle-support UNVERIFIED). OLMo-1 (non-parametric LayerNorm), OLMoE/FlexOlmo (MoE), OlmoHybrid (SSM) stay `INVENTORIED`. Nothing implemented. | -| 2a | `ROAD-V1-C2-LOCAL-BF16` | Close local discrete-Blackwell Qwen3.5 plain-BF16 production parity | [`MODEL-MM-qwen3-5-qwen3-5-for-conditional-generation`](model-matrix.md), [`LOAD-SAFETENSORS-DIRECT-DENSE`](engine-matrix.md), [`SERVE-CLI-BENCH`](engine-matrix.md), [`KERNEL-SSM-MAMBA`](kernel-matrix.md) | Exact `(sequence, 8-token chunk)` GDN conv dispatch is byte-identical and default ON. Rebased-main same-binary reprofile: **718.704→233.955 ms (3.072x)** and whole run **+2.272%**; sealed vLLM remains 145.421 ms (**1.609x residual**). The reviewed combined default-off K4 causal-conv + 16-token post-conv local A/B is exact and improves total throughput **0.7175%**, TTFT **1.340%**, TPOT **0.556%** and E2E **0.716%**. The #206 frontend now defaults to pre-timed tokenization with exact timed-string rollback and is CPU-GREEN with identical prompt/output IDs. Cross-engine throughput and TTFT remain **VOID/PENDING** until the reviewed GPU rerun; TPOT and VRAM remain OPEN | `GATING` | Fresh mutation review, operator CPU gate, then counterbalanced `STRING→TOKENS→TOKENS→STRING` GPU A/B and pinned-vLLM rerun. Latency, VRAM and 27B/35B correctness remain gates; no 4B-to-gate-model extrapolation. Evidence: [campaign spec](specs/sm120-qwen35-pareto-2026-08-09.md), [production baseline](../docs/bench-evidence/qwen35-4b-sm120-main-20260807.md) | +| 2a | `ROAD-V1-C2-LOCAL-BF16` | Close local discrete-Blackwell Qwen3.5 plain-BF16 production parity | [`MODEL-MM-qwen3-5-qwen3-5-for-conditional-generation`](model-matrix.md), [`LOAD-SAFETENSORS-DIRECT-DENSE`](engine-matrix.md), [`SERVE-CLI-BENCH`](engine-matrix.md), [`KERNEL-SSM-MAMBA`](kernel-matrix.md) | Exact `(sequence, 8-token chunk)` GDN conv dispatch is byte-identical and default ON. Rebased-main same-binary reprofile: **718.704→233.955 ms (3.072x)** and whole run **+2.272%**; sealed vLLM remains 145.421 ms (**1.609x residual**). The reviewed combined default-off K4 causal-conv + 16-token post-conv local A/B is exact and improves total throughput **0.7175%**, TTFT **1.340%**, TPOT **0.556%** and E2E **0.716%**. The #206 pretoken frontend is CPU-GREEN, but its first real A/B failed output identity (98/128 requests; 15,507/16,384 positions) because one-at-a-time AsyncLLM publication changed the execution wave; all apparent timings are VOID. Atomic wave admission matching vLLM's add-all-before-step behavior is now specified. Cross-engine throughput and TTFT remain **VOID/PENDING**; TPOT and VRAM remain OPEN | `GATING` | Fresh implementer: RED-first atomic batch queue/AsyncLLM wave admission and benchmark refill-wave dispatch; fresh mutation review; operator CPU gate; then restart counterbalanced `STRING→TOKENS→TOKENS→STRING` GPU A/B and pinned-vLLM rerun. Latency, VRAM and 27B/35B correctness remain gates; no 4B-to-gate-model extrapolation. Evidence: [campaign spec](specs/sm120-qwen35-pareto-2026-08-09.md), [production baseline](../docs/bench-evidence/qwen35-4b-sm120-main-20260807.md) | | 3 | `ROAD-V1-C3` | MTP k=1 + GDN speculative path, then DFlash, DSpark and heterogeneous-vocabulary TLI | [engine matrix](engine-matrix.md), [coverage view §8](feature-matrix.md#8-speculative-decoding) | MTP and DFlash specs exist. **M-mtp-0 CLOSED 2026-07-24** - the standalone MTP draft head is oracle-parity-proven on BOTH gate checkpoints (op-level parity vs a dumped k=1 vLLM oracle, not a token-generation SACRED gate). **I2 SCHEDULER-HALF LANDED 2026-07-24** ([mtp-spec-decode §2.7](specs/mtp-spec-decode.md)): host-side spec-decode scheduler/engine plumbing + the FROZEN spec-metadata ABI that I3 (rejection sampler) and I5 (verify/propose runner) build against - `SpeculativeConfig`, `DraftTokenIds`, `Request::spec_token_ids`/`NumTokensWithSpec`, the first population of `scheduled_spec_decode_tokens`, `Scheduler::update_draft_token_ids`, the `take_draft_token_ids` seam, `EngineCore::post_step`, `InputBatch::num_accepted_tokens`/`update_req_spec_token_ids`; DEFAULT-OFF and INERT (no `SpeculativeConfig` => `num_lookahead_tokens == 0` => byte-identical engine). `SPEC-MTP` **STAYS `GATING`** because M-mtp-1..4 (greedy rejection, GDN spec slots, k>1, CUDA graphs) are still open, so spec decode remains user-invisible. DSpark is user-promoted scope with DeepSeek-V4/Qwen3 draft models, reduced-vocabulary handling and full-CUDA-graph behavior inventoried under `SPEC-DSPARK`; tokenizer-agnostic target<->draft mapping is separately inventoried as `SPEC-TLI`. Their dedicated spikes are not written **I3 GREEDY REJECTION SAMPLER LANDED 2026-07-24** (`SPEC-REJECTION` `READY` -> `ACTIVE`): per-request logits expansion to `1 + k_i` rows plus the greedy accept rule (accept a draft iff it equals the target argmax; on the first mismatch emit the target argmax and stop; bonus token when all k accept), CUDA==CPU bit-exact at vocab 248320. **I4 GDN-HALF LANDED 2026-07-24** (`SPEC-GDN-SEGMENTS` -> `ACTIVE`): the GDN spec metadata split + decode->prefill reclassification, the `T>1`/`IS_SPEC` recurrence with per-timestep snapshots, the conv sliding window honouring `num_accepted`, k+1 slot allocation - bit-exact rollback. **I5a GDN LAYER ROUTING + runner spec-metadata upload LANDED 2026-07-24** (`CLAIM-SPEC-MTP-I5A`): `GdnBlockPaged` now routes a pure-spec batch through `vt::GdnSpecDecode`/`vt::CausalConv1dSpecUpdate` and the runner uploads I4's six spec device tensors - first of the scoped M-mtp-1 sub-increments (I5a GDN wiring -> I5b prepare_prefill -> I5c MTP paged propose -> I5d config+runner-loop+27B token gate, spec §5), DEFAULT-OFF INERT, bit-exact vs the I4 ops, no e2e loop yet. **I5b `prepare_prefill_inputs` LANDED 2026-07-24** (`CLAIM-SPEC-MTP-I5B`): the drafter prefill input-prep - shift each request's `input_ids` left one within its query span, splice the just-sampled next token, `query_len -= num_rejected`, emit last-token index / query_start_loc / seq_lens into the `SpecPrefillInputs` struct; a HOST routine (no new CUDA kernel; mirrors our DEVICE-NEUTRAL `prepare_inputs`/`combine_sampled_and_draft_tokens` family), unit-gated 7 cases / 27 assertions RED-first, DEFAULT-OFF INERT, additive by construction. **I5c MTP PAGED PROPOSE + DRAFT KV LAYER LANDED 2026-07-24** (`CLAIM-SPEC-MTP-I5C`): `Qwen3_5MTPModel::ForwardPaged` runs the head + one full_attention decoder layer over the head's OWN paged draft KV layer (ReshapeAndCache + PagedAttention over the target's block table / slot mapping); `MakeQwen3_5KVCacheSpec(num_spec>0)` adds that draft KV layer (`fa_draft` FullAttentionSpec group, index num_hidden_layers); `ForwardDeviceTap` exposes the `[T,H]` post-final-norm hidden tap (INERT); and `MtpProposePrefill` is the callable k=1 propose (I5b shift-splice -> one paged forward -> argmax at last_token_indices, early-exit). CORE PROOF: the paged forward reproduces I1's standalone head logits/argmax on BOTH gate checkpoints; a two-step drive proves the draft-KV write/read (RED control diverges). DEFAULT-OFF INERT (no spec config -> draft KV layer not allocated, tap nullptr, target forward byte-identical); NOT wired into the runner step loop. **I5d-pre REGISTRY/FORWARD-SEAM ENABLING REFACTOR LANDED 2026-07-25** (`CLAIM-SPEC-MTP-I5D-PRE`): a scoping pass found the model seam is fully TYPE-ERASED, so the runner cannot reach the concrete target weights / hidden-state tap / loaded MTP weights the I5d loop needs. Four ADDITIVE, inert-when-spec-off access paths + one latent-bug fix - the `hidden_tap` out-field on the type-erased `ModelForwardInput` (routes to the existing `ForwardDeviceTap`), a `LoadedModel::BuildMtpDraft` virtual (typed path to the draft, null for non-MTP), MTP weight loading + shard retention in `FromModelDir` behind `EngineParams::speculative_config`, and the `GPUModelRunner` ctor widened with optional draft/draft-KV/`SpeculativeConfig`; PLUS the latent `initialize_kv_cache` fix (select the FIRST non-eagle full-attn group as the target so a third `fa_draft` group can't displace it; byte-identical at num_spec==0). DEFAULT-OFF INERT, unit-gated RED-first, spec-off SACRED gates byte-identical. **I5d CONFIG RUNTIME + VERIFY/PROPOSE RUNNER LOOP LANDED as a spec-off-byte-identical PARTIAL 2026-07-25** (`CLAIM-SPEC-MTP-I5D`): `--speculative-config` JSON parse -> `EngineParams` -> `LoadedEngine` resolution (widened KV `MakeQwen3_5KVCacheSpec(num_spec>0)`, `BuildMtpDraft`, forced sync scheduling, `MakeScheduler(spec)`, `EngineCore(check_for_draft=true)`) + the full runner loop (draft splice, hidden-tap capture, GDN builder spec-overload feed, k+1 GDN state-slot remap + widened conv cache + draft-KV alloc, `MtpProposePrefill`, `take_draft_token_ids`, acceptance telemetry). CUDA `-Werror` 0 warnings, cutlass-ON; spec-OFF SACRED byte-identical (27B 235/235, 35B 315/315, Coder 138/138 + spec unit tests ALL PASS). **The three-way 27B token gate is NOT yet passing**: the spec-ON engine RUNS the loop end to end and MEASURES the blocker (`test_qwen27_spec_decode`) - it throws on the FIRST prefill step at `gdn_state_gather: working/cache row shapes must match` (`src/vt/ops.cpp:1773`) because I4's spec conv rollback needs the conv row widened to `(K-1)+num_spec` while the non-spec GDN conv ops assume `(K-1)`. `SPEC-MTP` STAYS `GATING`. **I5e LANDED 2026-07-25 (`CLAIM-SPEC-MTP-I5E`): the non-spec GDN conv ops made widened-cache-aware (mirror vLLM `state_len=KERNEL_WIDTH-1` + physical `stride_conv_state_tok`, leading `(K-1)` sub-window, byte-identical at `num_spec==0`) AND the async input-combine forced off under spec (it overwrote the verify batch's draft position with the committed token -> 0 acceptance, RCA'd on the real 27B). **THE THREE-WAY 27B SINGLE-REQUEST GREEDY GATE PASSES**: our-spec-ON == vLLM `--speculative-config mtp` greedy == our-spec-OFF token-for-token, **acceptance 16/16 drafts accepted** (~16 target steps saved); spec-OFF SACRED byte-identical (27B 235/235, 35B 315/315, Coder 138/138), compute-sanitizer 0 on the spec step. `SPEC-MTP` LEAVES `GATING` (single-request greedy correctness PROVEN); NOT `DONE` - the MIXED `GdnBlockPaged` concurrency split/merge + the throughput A/B vs vLLM same-config are I6. **I6 LANDED 2026-07-25 (`CLAIM-SPEC-MTP-I6`, `benchmark_binding=true`): the §5 c1 THROUGHPUT GATE — ours spec-ON AT/ABOVE vLLM spec-ON on EVERY measured axis at c1** (TPOT 66.2/62.95 vs 69.1/65.3 ms prose/code, ours ~1.04x faster; output tput +4.6%/+3.9%; ITL/TTFT lower; acceptance ours 0.85/0.92 vs vLLM 0.838, within noise; spec helps both ~1.5-1.6x TPOT; ours ~4% faster spec-OFF too), via an additive example-only `--speculative-config` bench flag (NO engine code touched). STAYS `ACTIVE`: the c>1 mixed-batch path is still refused + owes a c>1 A/B, and no server-facing spec flag yet. | `ACTIVE` | M-mtp-0, I2 scheduler-half, I3 rejection sampler, I4 GDN spec slots, I5a GDN layer routing, I5b prepare_prefill, I5c MTP paged propose + draft KV, I5d-pre the registry/forward enabling seam, and I5d config runtime + verify/propose runner loop (spec-off byte-identical) are landed; next (before `SPEC-MTP` leaves `GATING`) is closing the measured I5d gate blocker - make the non-spec GDN conv ops widened-cache-aware (mirror vLLM `causal_conv1d` `state_len=width-1+(seqlen-1)`) + the MIXED `GdnBlockPaged` split/merge - then the passing M-mtp-1 27B k=1 greedy three-way token gate + acceptance, then M-mtp-2 35B, then DFlash, the DSpark spike/gates and TLI. **DFlash D0-redo + D1 LANDED 2026-07-26 (`CLAIM-DFLASH-D0D1`, [dflash-spec-decode §0](specs/dflash-spec-decode.md)): `SPEC-DFLASH` UNBLOCKED + `ACTIVE` on the advanced pin `555967922`/vLLM 0.26.0.dev0.** The prior 0.25.0 ORACLE-BLOCKED verdict is SUPERSEDED — under `VLLM_USE_V2_MODEL_RUNNER=1` (vllm#40898 resolved) the mixed-SWA/full z-lab 27B draft CONSTRUCTS and the drafter is ALIVE (acceptance 2.21/8.80/4.75/4.57 > 1, `num_spec=16`, flashinfer-native fp8-KV; goldens committed). Gate FORM measured STRICT MODE-MATCHED (vLLM-ON run-deterministic K>=3 but != vLLM-OFF at k=16 near-ties — NOT the MTP three-way identity). **D1 `DF-AUX-TAPS` DONE:** the single hidden tap is generalized to the multi-tap `[T,H×taps]` (`ForwardDeviceMultiTap` capturing `(hidden+res)` at `target_layer_ids`), config-gated byte-identical off; unit gate 598 assertions (RED-first), CUDA 697/697 + sanitizer 0, 27B MTP e2e 9/9 + 27B SACRED 235/235 byte-identical (inertness). **D2-D5 LANDED 2026-07-26 (`CLAIM-DFLASH-D2`/`D3`/`D4D5`/`D5`):** the drafter model + the project's first non-causal in-block attention (D2, GPU parity vs the real vLLM draft), context-KV precompute + `prepare_dflash_inputs` (D3, GPU numeric-parity 61/61), the non-autoregressive whole-block propose brick + `dflash` config-select (D4), and the RUNNER-LOOP INTEGRATION + 27B e2e (D5): the full verify/propose loop is wired (separate z-lab draft load + target-shared bf16 embed/lm_head, aux-tap capture, per-request combined-feature context accumulation honoring num_rejected, `propose_drafts_dflash`) and RUNS end to end - `test_qwen27_dflash_spec_decode` 2/4 STRICT token-exact vs the vLLM-DFlash-ON golden + acceptance ~ vLLM on ALL 4 (19/39/29/25 vs 17/39/30/25). The 2 divergences are SINGLE bf16 near-tie flips (ratified near-tie ROOT = the D3 inline context-KV recompute envelope), NOT a wiring bug; inertness SACRED 235/235 + MTP 9/9 byte-identical; CUDA `-Werror` clean, no new kernel. NOT a clean strict-4/4 pass - STRICT 4/4 token-identity + the speed A/B = D6 (persistent paged draft-KV bit-matching vLLM's fused projections + the uniform-1+k FULL CG). Capture tool + goldens: `scripts/spec/d{0,2,3}_dflash_*.py`, `tests/parity/goldens/dflash_27b{,_draft,_kvprep}/`. **D6-D9 SPEED CAMPAIGN 2026-07-27:** D6 c1 A/B + STRICT-4/4 bf16-irreducibility RCA; D7 device-resident within-step forward (bit-identical); D8 acceptance RCA + FINAL golden A/B (ours 0.69× vLLM). **D9 (`CLAIM-DFLASH-D9`) PERSISTENT PAGED DRAFT-KV LANDED (bit-identical, +22.7%): `AppendContextKVHost` + `ForwardBlockLogitsWithPrecomputedKV` replace the O(context²) per-step recompute with an append-only per-request store; ours-ON 20.99→25.75 tok/s = 0.917× vLLM-ON (28.09, was 0.69×); e2e 27/27 SAME tokens, SACRED 235/235 + MTP 9/9 byte-identical, CUDA `-Werror` clean, no new kernel. D8's "bf16 acceptance ceiling" REFUTED — same-trajectory per-step acceptance == vLLM (ratio 1.00) and ours realized acceptance (3.68/step) > vLLM (3.31); the SOLE residual (~8%) is the FULL uniform-(1+k) CUDA graph (eager-vs-graphed), a closeable increment. SPEC-DFLASH stays `ACTIVE` (speed not yet ≥ vLLM).** **SPEC-MTP → `DONE` 2026-07-26 (`CLAIM-SPEC-MTP-DONE`, records-only, closing commit I7 `72f9fb1`):** the user ratified the c>1 near-tie+SPEED criterion, closing both I6-owed items (mixed-batch concurrency + server/CLI/C-ABI `--speculative-config`); MTP k=1 is COMPLETE + gated. **M-mtp-2 CLOSED 2026-07-26 (`CLAIM-SPEC-MTP-M-MTP-2`): the 35B `Qwen3_5MoeMTP` full e2e three-way token gate PASSES** — our spec-ON == our spec-OFF == vLLM 0.25.0 `--speculative-config mtp` greedy == vLLM spec-OFF, 16/16 vs the `greedy_ids` anchor (STRICT, c1), acceptance 16/16 both sides; c1 spec-ON 1.19x TPOT / +16.3% output-tput vs spec-OFF (acceptance 0.908) — MoE speedup transfers; spec-OFF byte-identical (test+docs-only). MTP is now `DONE` on BOTH gate models (`MODEL-SPEC-qwen3-5-mtp-qwen3-5-moe-mtp` `GATING`→`DONE`). **DFlash D11+D12 2026-07-27 — the FULL uniform-(1+k) CUDA graph is being built in three parts:** D11 (`CLAIM-DFLASH-D11`) landed Part A (the device-store primitive, CPU-gated); **D12 (`CLAIM-DFLASH-D12`) landed A-wire (the D11 device store is now the PRODUCTION path; GPU-gated e2e 27/27 all-exact acceptance 19/39/29/25 + SACRED 235/235 + MTP 9/9 byte-identical) + Part B (`vt::DFlashPagedBlockAttention`, the capture-safe paged kernel; `test_ops_dflash_paged_block_attn` 795648/795648 CPU==CUDA + cross-check vs materialized `DFlashBlockAttention` + compute-sanitizer 0; NO function-local host `cu_seqlens` upload = capture-UAF fixed).** Speed UNCHANGED 0.917× (A-wire eager + Part B not yet wired). The SOLE remaining piece is Part C (static-shape capture + device mask-scatter + `BeginCapture`/replay + the ≥vLLM c1 A/B); if ours-ON-graphed ≥ vLLM-ON → SPEC-DFLASH DONE → C3 complete. C3 stays `ACTIVE` (DFlash Part C + DSpark/TLI remain) | | 4 | `ROAD-V1-C4` | Quantization: llama.cpp breadth/speed, NVFP4/FP8/MX, MLX native | [quantization matrix](quantization-matrix.md) | coverage spike merged; `QUANT-GGUF-CPU-THREADPOOL` W1-W3 implemented and correctness-gated, still `GATING` (its reproduction now exists — same-binary 1-vs-20-thread A/B is prefill 12.47x / decode 8.05x / RSS 1.000x, so **decode misses the >=10x bar**). **GGUF COMPUTE-IN-QUANT IS NOW LIVE AND DEFAULT-ON (2026-07-22, `CLAIM-QUANT-GGUF-CIQ-G4-1`):** [compute-in-quant GEMM](specs/gguf-compute-in-quant-gemm.md) **G1-G4** — block dtypes + traits, the Q8_0/Q8_K activation quantizers, the six generic `vec_dot`, `kMatmulBTQuant`, and now the ROUTING (`vt::MatmulBT` dispatches a block-dtype weight to it) — plus [keep-quant loader](specs/gguf-keep-quant-loader.md) **L1-L4**, whose master switch defaults ON wherever that op is registered for the running device (CPU today; a CUDA build still expands). Six encodings (Q4_0, Q8_0, Q3_K, Q4_K, Q5_K, Q6_K) now carry `C` = `Y`. **Correctness held exactly**: the 35B GGUF gate is 16/16 token-exact vs the same-file llama.cpp oracle with the quant path on, and the bench model's output tokens are byte-identical across the pre-G4, post-G4 and `VT_CPU_REF=1` arms — no golden regenerated. **Binding CPU A/B** (idle dgx aarch64, one flock, same binary, 3 reps): decode **3.45x**, prefill **4.16x**, peak RSS **1.16x less**; vs llama.cpp we went from 11.7x / 34.1x / 2.66x behind to **3.38x / 8.20x / 2.29x**. The projected 9-17x did **not** hold, for a measured reason: 60 % of that file's weight bytes are `f16`, which no block encoding covers. Keep-quant loader **L4** is therefore MEASURED-and-NOT-MET on RSS; other leaf specs open. **THAT #1 LEVER IS NOW LANDED (2026-07-22, `CLAIM-KERNEL-CPU-ELEM-GEMM-1`, new row [`KERNEL-GEMM-CPU-ELEM`](kernel-matrix.md)):** [the elementwise CPU GEMM](specs/cpu-elementwise-gemm.md) **E1-E4** — per-dtype specialization out of the K loop, 16 independent accumulators instead of one, AArch64 NEON + x86-64 SSE2/F16C tiers behind a runtime probe, and M-blocking — all **BYTE-IDENTICAL** to the historical kernel (`memcmp` gate, exhaustive 65,536-pattern widening check, same token md5), so nothing was regenerated. Binding same-binary A/B: prefill **3.41x**, decode **3.11x**; op-level bf16 18-24 -> 69-351 GFLOP/s. **vs llama.cpp: decode AT PARITY (1.03x), prefill 2.34x behind, RSS 2.29x worse. THEN loader L5 LANDED (2026-07-23, `CLAIM-QUANT-GGUF-KEEPQ-L5-1`):** [keep-quant loader](specs/gguf-keep-quant-loader.md) **L5** — mmap in-place residency (borrow kept blocks out of the read-only mapping, refcounted), tied-head sharing (one bf16 vocab matrix for embed+lm_head), and a read-once page release (port of llama.cpp `unmap_fragment`) — took **peak RSS 6.401 -> 3.884 GiB (2.29x -> 1.39x llama.cpp)** with decode UNCHANGED and output tokens byte-identical (md5 `d235db12f2cd304007530286a1755c95`). The remaining ~1.09 GiB over llama.cpp is the f16 expansion (no block encoding covers f16). | `PARTIAL` | **THE OWED FRESH PROFILE IS DONE (2026-07-23) and it re-ranks the plan.** A `vt::GetOp` hook (100% of wall time, reverted before binding) on the CURRENT binary: prefill is **no longer GEMM-bound** — kMatmulBTQuant 37%, **kGdnPrefill 25%**, kMatmul 12%, kMatmulBT 10%, **kPagedAttention 10%**; the two non-GEMM kernels (GDN linear-attention recurrence + paged attention) run **SINGLE-THREADED** on the CPU and are now the top prefill levers. Decode is memory-bound matmul at parity, no kernel work owed. **RE-RANK: G5/G6/G7 all only speed the already-fast quant GEMM and rank BELOW the two serial non-GEMM kernels; the new #1 CPU lever is threading kGdnPrefill + kPagedAttention.** **THAT #1 LEVER IS NOW LANDED (2026-07-23, `CLAIM-CPU-THREAD-GDN-PAGED-1`, [two-kernel threading](specs/cpu-thread-gdn-paged-2026-07-23.md)):** kGdnPrefill chunks over the (sequence, value-head) axis and kPagedAttention over query-token rows, both via the existing `ParallelForRows`, both **byte-identical** (qwen35 output-token md5 `d235db12f2cd304007530286a1755c95` unchanged at threads 1/4/20 + `VT_CPU_REF=1`, determinism battery extended, CPU ctest 158/158). **Binding dgx aarch64 (idle): prefill 1.382× same-binary (73.0→100.9 t/s), 2.43×→1.76× behind llama.cpp pp128; decode at parity; op-scaling 1→20 GdnPrefill 7.08× / PagedAttention 8.96×; fresh profile shows the two kernels 35%→8.6% of prefill and re-ranks the NEW bottleneck to the GEMMs (kMatmulBTQuant 50% + kMatmul 16% + kMatmulBT 14% = 80%) ⇒ next CPU lever is the SIMD/repack GEMM tiers (G5/G6/G7).** **THE FIRST SUCH TIER IS NOW LANDED (2026-07-23, `CLAIM-QUANT-GGUF-CIQ-G6-1`, [compute-in-quant GEMM](specs/gguf-compute-in-quant-gemm.md) G6):** the Arm **i8mm mmla `nrc==2`** `vec_dot` 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 (decode M=1 → portable, unchanged); runtime `HWCAP2_I8MM` probe + `VT_CPU_QUANT_MMLA` defeat + per-file `+i8mm`. **BYTE-IDENTICAL** where the math allows (q8_0/q4_0 bit-exact to the scalar tier, q4_K/q6_K within NMSE ≤ 5e-4), bit-identical across threads 1/2/4/20, e2e token md5 `d235db12f2cd304007530286a1755c95` byte-identical (mmla on/off/`VT_CPU_REF=1`), 35B GGUF gate 16/16 vs llama.cpp on both files. **Op-level portable→i8mm: q4_K 7–8.4×, q6_K 3.8–4.5×, q8_0 ~1.2×**; e2e prefill same-binary **1.084×** on the q8_0-dominant bench file (1.56×→**1.44× behind** llama.cpp pp128, Amdahl-bounded — the big k-quant win lands on the APEX 35B files). Fresh bottleneck: the elementwise f16/f32 GEMM (~30%, unchanged) is now co-dominant on this mixed file. CUDA `-Werror` 0-warn, regression set UNCHANGED. `docs/BENCHMARKS.md` ACCEPTED. RSS deficit closed to 1.39x by L5; the last RSS lever is an f16 keep-as-is compute path, not this loader. **THEN the GDN split-projection orientation LANDED (2026-07-23, `CLAIM-CPU-GDN-ORIENT-1`, [GDN projection orientation](specs/cpu-gdn-proj-orientation-2026-07-23.md)):** a fresh op-dispatch profile of the current binary (warm prefill, `vt::GetOp` hook + per-GEMM shape histogram, reverted before binding) found the four GDN input projections (`in_proj_qkv/z/b/a`, 72 GEMMs, **17.9%** of prefill: `kMatmulBTQuant 50.7% / kMatmul 17.9% / kMatmulBT 14.9%`) were the LAST weight family `LoadGdnGguf` still transposed into [K,N] (nk=false → the N-striding `kMatmul`, no M-blocking) after G4's `expand_nk` gave every other expanded weight the file's own [N,K] order. New `GgufLoadPolicy::gdn_expand_nk` + `MakeGdnProj` keep them [N,K] nk=true → the M-blocked `kMatmulBT`; **BYTE-IDENTICAL** (same sequential f32 K-reduction, only the weight offset differs — token md5 `d235db12f2cd304007530286a1755c95` unchanged across on/`VT_GGUF_GDN_NK=0`/`VT_CPU_REF=1` and threads 1/4/20), `test_qwen36_gguf_engine` 2/2·28/28·16/16 on APEX. **Binding same-binary prefill 1.090× / decode 1.09× (44.1→40.4 ms TPOT = 1.01× llama tg32, at parity), 1.44×→1.32× behind llama.cpp pp128, RSS unchanged.** Fresh post-change profile: `kMatmul` **17.9%→0% (72→0 calls, ELIMINATED)**, absorbed into `kMatmulBT` (14.9%→27.7%); **next CPU prefill lever = the quant GEMM (kMatmulBTQuant, now 55%): G7 repack-at-load.** **G7 LANDED 2026-07-23 (`CLAIM-QUANT-GGUF-CIQ-G7-1`, [compute-in-quant GEMM](specs/gguf-compute-in-quant-gemm.md) G7):** q8_0 repacked once at load into the i8mm `block_q8_0x4` interleave (ported llama.cpp `repack.cpp` `q8_0_4x8`), `kMatmulBTQuant` dispatches a pre-shuffled gemm/gemv with no per-block register shuffles. **BIT-IDENTICAL** (byte-permute weight + non-fused `vmlaq_f32`, 305-assertion memcmp across decode/prefill/bf16-out/threads, token md5 `d235db12f2cd304007530286a1755c95` unchanged on/`VT_CPU_QUANT_REPACK=0`/`VT_CPU_REF=1`; a `ResidentWeight`/`MakeTensor` flag-drop that produced all-zero tokens was caught by the E2E gate and fixed). Op-level q8_0 **3.7–5.9×** (518→2401 / 583→3456 / 514→1902 GFLOP/s); **E2E prefill 1.92× same-binary (1096→572 ms), 223.8 t/s vs llama.cpp pp128 177.3 = 1.26× — AT/BEYOND PARITY** (was ~1.5× behind), decode at parity, RSS 3.884 GiB unchanged. Fresh profile: q8_0 GEMM 55%→~21%; **the CPU prefill-lever search is CLOSED — the sole remaining gap to llama.cpp is peak RSS (1.39×), not prefill.** CUDA-inert (gated off any non-CPU-quant device), CUDA `-Werror` 0-warn, regression set UNCHANGED. **THE RSS GAP IS NOW CORRECTLY ATTRIBUTED (2026-07-23, `CLAIM-QUANT-GGUF-KEEPF16-L6-1`, [keep-quant loader](specs/gguf-keep-quant-loader.md) L6): it is NOT the f16 expansion.** L6 implemented keep-f16 residency (keep the file's 56 F16 weights + tied head resident as F16 and compute on them, mirroring llama.cpp `ggml_vec_dot_f16`) and MEASURED it **RSS-NEUTRAL** (3.884 → 3.832 GiB, −52 MB) and prefill-regressive (TTFT 577 → ~1000 ms, from 1.25× ahead of llama.cpp to 0.72× behind) — because L5's page-release had ALREADY dropped the f16 file pages, so keep-f16 only swaps an anonymous bf16 buffer for equal-size file-backed f16 pages. smaps attribution proves our weight residency is at llama.cpp parity (file-backed 2.63 ≈ 2.68 GiB); **the residual ~1.08 GiB is the engine's ANONYMOUS activation/KV workspace, not weights — the real, separate CPU RSS lever.** keep-f16 ships DEFAULT OFF (`VT_GGUF_KEEP_F16=1` opt-in), tokens byte-identical (md5 `d235db1…`), `test_gguf_keep_quant` 35/35 (x86+aarch64), regressions UNCHANGED (27B 235/235, 35B 315/315, Coder 138, dense 184, OPT 63, DeepSeek 223, Llama 92, GGUF engine 28/28). **NEXT CPU RSS lever: profile + shrink the engine's activation/KV working set, NOT the weight loader** | | 5 | `ROAD-V1-C5` | Sliding window, local attention, YaRN/long context | [engine matrix](engine-matrix.md), [coverage view §§2,11](feature-matrix.md#2-kv-cache--memory), [joint spike](specs/sliding-local-yarn-long-context.md) | **CUDA GPU CLOSURE 2026-07-27 (`CLAIM-ROADMAP-C5`, dgx GB10 sm_121a, clean build of `489f7771`, oracle vLLM 0.26.0.dev0):** the shared scaled-RoPE + local-mask CUDA path COMPILES `-Werror`-clean and RUNS on GB10; the C5 feature-positive correctness gates that were the stated `GATING` blocker now PASS — SWA (Gemma-2/Gemma-3 48/48), LongRoPE (Phi-4-mini 16/16, RED-first), llama3-rope (Llama-3.2-1B 16/16), dynamic-NTK (InternLM2 16/16); both RoPE 0.26-oracle recaptures BIT-IDENTICAL to goldens (zero drift). Leaves `ATTN-SLIDING-WINDOW`/`ATTN-ROPE-{LLAMA3,LONGROPE,DYNAMIC-NTK}`/`ATTN-YARN` → `ACTIVE` | `PARTIAL` | (RI) **Honest residual (vehicle-blocked, not skipped):** YaRN model e2e (no cached Nomic/gpt-oss consumer) + chunked-local model e2e (no Llama4 row) are REACHABLE-BLOCKED — operator/formula stay GPU/G3-gated; long-context positive-mask (prompt > W) SWA model e2e + the KV memory-optimization G8; and the roadmap-wide every-axis SPEED tail (all C5 leaves correctness-complete, speed-pending, mirroring their model consumers). Not row-DONE until speed + the blocked vehicles close | diff --git a/.agents/specs/sm120-qwen35-pareto-2026-08-09.md b/.agents/specs/sm120-qwen35-pareto-2026-08-09.md index d21a7824a..572b58ed9 100644 --- a/.agents/specs/sm120-qwen35-pareto-2026-08-09.md +++ b/.agents/specs/sm120-qwen35-pareto-2026-08-09.md @@ -6,8 +6,10 @@ VRAM gaps vs vLLM](https://github.com/mudler/vllm.cpp/issues/206) **Rows:** `KERNEL-SSM-MAMBA`, `SERVE-CLI-BENCH`, feeding `ROAD-V1-C2-LOCAL-BF16` -**Lifecycle:** combined prefill opt-in accepted locally; frontend timing parity -correction implemented and CPU-gated, with fresh review and GPU A/B pending +**Lifecycle:** combined prefill opt-in accepted locally; frontend pretokenization +is CPU-gated but its first real GPU A/B failed token identity because sequential +enqueue changed the first scheduling wave; atomic wave admission is specified +below and pending implementation/review **Target:** local NVIDIA GeForce RTX 5070 Ti (`sm_120`), Qwen3.5-4B plain BF16. This vehicle does not substitute for the hardware-unavailable 27B/35B @@ -142,9 +144,80 @@ rollback produce identical prompt and output token IDs in submission order. A real `TemplateProcessing` BOS fixture proves raw `Encode` is not an equivalent substitute, a pure preparation/clock seam proves all default tokenization ends before timing begins, and report capture proves the resolved mode remains -auditable. This is not a performance acceptance: fresh mutation re-review, -operator verification and the counterbalanced production GPU A/B remain -pending. +auditable. This was not a performance acceptance. Fresh mutation review and +operator verification later passed; the first production GPU pair then failed +the separate execution-wave identity gate below. + +### Real-GPU failure and atomic wave-admission amendment (2026-08-09) + +The clean `a33993a7cd8fdcf33b2b91112ef7a172b8f63fe1` binary passed the +frontend CPU contract and review, but the first production A/B pair failed the +hard token-identity gate. Both arms reported 131,784 input tokens and produced +128 requests x 128 output tokens. STRING-r1 and TOKENS-r1 matched only 98/128 +requests and 15,507/16,384 output positions; differing request IDs were +`1,7,10,23,30,41,46,50,51,60,61,68,70,74,77,80,82,83,85,86,88,92,95,110,113,116,122,124,126,127`. +Their output JSON SHA-256 values are respectively +`83fcdc45f79ddb06a634c7d7d95eba3384543b3cd781a45a8db1fc4e2a453545` +and `be20ffbceb61f0264ca21d972bfc5fc51e855e64f2b945de71669cae666aa702`. +The series stopped immediately; its apparent timing movement receives no +credit. + +This is an execution-batch mismatch, not evidence that pretokenized prompt IDs +are semantically different. Each local `AsyncLLM::add_request` registers one +collector and immediately calls `InprocClient::add_request_async` +(`src/vllm/v1/engine/async_llm.cpp:36-123`), which appends one ADD and notifies +the background core (`src/vllm/v1/engine/core_client.cpp:56-72` and +`include/vllm/v1/engine/core_proc.h:90-101`). The core blocks for the first item +and then drains only the items already available +(`src/vllm/v1/engine/core_proc.cpp:142-178`). Sequential string admission does +tokenization between publishes, while the pretokenized path publishes rapidly; +therefore the core may schedule different partial first waves. The pinned +closed-loop wrapper instead calls the synchronous engine's `add_request` for +the complete initial concurrency before its first explicit `engine.step`, and +does the same for every refill after a step +(`tools/bench/vllm_closed_loop_metrics.py:57-102`). Matching only the tokenizer +boundary is insufficient: the selected request wave must become visible as one +unit. + +The smallest additive repair is a batch/wave admission path with this binding +contract: + +1. `AsyncLLM` accepts one ordered wave of request ID, exactly one prompt form + (string or token IDs), sampling parameters and priority, and returns the + corresponding `AsyncRequest`s in that order. Existing one-request overloads + remain unchanged and a one-item wave has identical behavior. +2. Before publishing any request, the batch path processes every selected + input, builds every core `Request`, and builds every collector. The benchmark + records all wave arrivals first and then makes exactly one batch call. In + rollback mode, string tokenization therefore remains after `t0` and after + arrival recording; in the default path, token IDs remain prepared before + `t0`. Only the frontend work differs; execution-wave composition does not. +3. Under the existing output-processor admission/shutdown lock, recheck engine + life and register all collectors before core publication. A duplicate, + input/factory/allocation failure, enqueue failure or concurrent shutdown + publishes no partial core wave and removes every frontend state newly + registered by that call without manufacturing terminal output. Existing + requests and request-ID reusability are preserved. +4. `InprocClient` publishes the prepared core requests through a + `BlockingQueue` batch operation that holds the queue mutex while appending + every ADD, rolls back any partial append before releasing the lock, and + performs exactly one notification after success. The consumer can observe + either zero or all wave items, never a prefix. +5. `RunBench` uses this path for the full initial concurrency and every refill + wave. Its pure dispatch seam must still select exactly one string or token + callback, and the resolved mode remains present in `BenchResult` and the + printed report. + +RED-first CPU tests must prove zero partial queue visibility while a batch is +being assembled, ordered all-at-once visibility after publication, exact +single-callback dispatch, string/token input and output ID equality on the +asynchronous synthetic BOS fixture, duplicate/failure/shutdown rollback, and +one-item compatibility. Mutation review must independently break the queue +lock span/notify count, publish before full preparation, callback selection, +rollback and benchmark wave call and observe focused failures. The real gate +then restarts the complete counterbalanced +`STRINGa -> TOKENSa -> TOKENSb -> STRINGb` series; all prompt/output IDs must +match before any latency or throughput value is interpreted. ## Correctness, tests, review and acceptance diff --git a/docs/BENCHMARKS.md b/docs/BENCHMARKS.md index caab28b2b..d3eb43938 100644 --- a/docs/BENCHMARKS.md +++ b/docs/BENCHMARKS.md @@ -41,7 +41,7 @@ The binding comparison. vLLM runs its **production graphed config**, never | Qwen3.6-27B | NVFP4 (`nvidia` @`0893e160`, ModelOpt `modelopt_mixed`) | 0.25.0 | 0/4 | **BEHIND, uniformly 0.85x** on decode throughput (was 0.72x before the FP8 tower fix); greedy continuation IDENTICAL to vLLM. A different model from the `unsloth` row (NVFP4 MLP + FP8 W8A8 GDN/attn tower) | | Qwen3.6-35B-A3B | NVFP4 `modelopt_mixed` | 0.25.0 | 2/18 | 3-rep grid 2026-08-05 @`1ea26427`: 0.93-1.03x (c4 wins), c16 0.93x. Both c16 levers A/B'd NEG: drain event -1.9%, mirror 0.999x. ★ probe found a prod async batch-1 greedy DEGENERATION bug the mirror fixes | | DeepSeek-V2-Lite | bf16 MLA | 0.25.0 | 4/25 | Attributed miss, row stays `ACTIVE` | -| Qwen3.5-4B | bf16 direct-load | 0.26.0.dev0 | host PSS; frontend GPU measurement pending | #206 CPU 10/10 (special IDs/pre-clock/report); GPU pending. Combined +0.718% tput/-1.340% TTFT/-0.556% TPOT; cross-engine tput/TTFT VOID; TPOT/VRAM OPEN ([data](bench-evidence/qwen35-4b-sm120-main-20260807.md)) | +| Qwen3.5-4B | bf16 direct-load | 0.26.0.dev0 | host PSS; frontend exactness failed | #206 pretoken A/B FAILED (98/128); timings VOID, atomic-wave retry pending. Opt-ins: +0.718% tput/-0.556% TPOT; cross-engine tput/TTFT VOID, TPOT/VRAM OPEN ([data](bench-evidence/qwen35-4b-sm120-main-20260807.md)) | ### GDN prefill causal-convolution by GPU diff --git a/docs/STATUS.md b/docs/STATUS.md index e9dfc9f44..e5068a5c8 100644 --- a/docs/STATUS.md +++ b/docs/STATUS.md @@ -78,7 +78,7 @@ token-for-token correctness against the pinned oracle. | Qwen3.6-27B (NVFP4) text generation | Correctness-complete; speed is CHECKPOINT-dependent | Token-exact GB10 on both. `unsloth` @`890bdef7` beats vLLM 0.25.0 every c (1.007-1.045x), 115/124; `nvidia` @`0893e160` (ModelOpt FP8 tower) is **0.85x BEHIND**, decode ~100% GPU-busy | | Qwen3.6-35B-A3B (NVFP4, GDN MoE) | Correctness-complete; binding grid @`a0fa12c7` FLAT 0.935x-0.979x over c1-c32 (CoV <0.81%); the prior 0.87x c2 / 0.92x c8 "weak cells" were harness mismatch, not code; memory PSS 3.81x, GPU 1.40x | Token-exact SYNC+ASYNC; `VT_ASYNC_DEVICE_MIRROR` ON fixes async batch-1 token-0 degeneration; `VT_ASYNC_EXECUTOR` Option A NEUTRAL → OFF | | Qwen3 / Qwen2 dense (BF16) | Correctness-complete, speed-pending. Async-serving P0 FIXED (`ROW-SERVE-ASYNC-DENSE-MIRROR`): classic-dense `Qwen3ForCausalLM` now honors the async device token-ids mirror; CPU-only -Werror test-guard fixes x2 | Near-tie-robust token-exact vs vLLM (Qwen3-0.6B, Qwen3-4B); c1 effective parity, c8 decode residual. **Async device-mirror (`ROW-SERVE-ASYNC-DENSE-MIRROR`, `f9c969ae`): the #31 fix ported to the classic dense family, dgx-VERIFIED.** The shared dense `EmbedInto` (qwen3.cpp) raced the async combine's device input-ids write against a stale host upload → token-0 degeneration on the depth-2 AsyncLLM serving path (quant-independent). `EmbedInto` now consumes the device override published by `ForwardQwen3ForCausalLM`'s `DeviceTokenIdsScope` (27B-dense template); gate `test_qwen3_dense_async_serving` RED on `VT_ASYNC_DEVICE_MIRROR=0`, GREEN default, byte-identical mirror-off. dgx GB10: async gate RED→GREEN 0.6B+4B, SACRED 0.6B+4B 184/184 unchanged (byte-neutral sync path), memcheck 0 errors; Yi30/Qwen3-8B-MXFP4 default-config e2e coherent + 3/4 token-exact (p2 = oracle-ratified near-tie, gap 0.0000), closing the QUANT-CT-MXFP4 async-default residual. RESIDUAL: sibling InternLM2/Mistral/Llama scope one-liner; W4 bench RAN; FA2 GQA-swap default-ON, c2-c8 <1.0x. `FLASH-PTXAS` #82: codegen at PARITY (no ptxas lever); gap=engine context. **D1 (2026-07-31, `CLAIM-D1-BF16-MERGED-QKV`): the bf16 merged-QKV path (`Qwen3QkvMergeEnabled`/`VT_QWEN3_QKV_MERGE`) is now default-ON** — one `vt::MatmulBT` over the merged `[qdim+2kdim,H]` owner + a contiguous `vt::QkvSplit` (OLMo-2 exemplar), replacing three per-shard GEMMs. Bit-exact GEMM math (A/B unit `test_ops_qkv_merge` byte-identical, RED-first); the wider-N cuBLASLt K-reduction flips the 0.6B genuine bf16 near-tie so the SACRED 0.6B golden was regenerated (all tokens within the near-tie band, max 0.125 nats), while Qwen3-4B is byte-neutral (0 diffs, stays STRICT). Re-gated 0.6B 16/16 + 4B 16/16; consistency/launch-count fold (measured NEUTRAL on 4B decode), no new throughput owed | -| Qwen3.5-4B plain BF16 direct loading on discrete CUDA | Correctness-complete; frontend CPU-green/GPU-pending; TPOT/VRAM open | #206 frontend CPU 10/10: special IDs, pre-clock timing, mode report. Cross-engine tput/TTFT VOID pending GPU; TPOT 1.025x, VRAM +233 MiB OPEN ([data](bench-evidence/qwen35-4b-sm120-main-20260807.md)) | +| Qwen3.5-4B plain BF16 direct loading on discrete CUDA | Correctness-complete; atomic frontend wave pending; TPOT/VRAM open | #206 real pretoken A/B FAILED exactness (98/128); timings VOID. Atomic wave admission pending; TPOT 1.025x, VRAM +233 MiB OPEN ([data](bench-evidence/qwen35-4b-sm120-main-20260807.md)) | | Qwen3-Coder-30B-A3B MoE (BF16) | Correctness-complete, speed-pending | Near-tie-robust token-exact 6/6; 11 of 16 binding grid cells at or above vLLM. **D1 (2026-07-31): inherits the default-ON bf16 merged-QKV via the shared dense `AttnBlock` — byte-neutral (0 token diffs, golden UNCHANGED); re-gated 6/6** | | Llama-3.x dense (BF16) | Correctness-complete, speed-pending | Near-tie-robust token-exact 16/16 (Llama-3.2-1B); llama3 RoPE scaling | | Mistral dense (BF16) | Correctness-complete, speed-pending | Paged-engine token-exact 16/16 (Mistral-7B-v0.3) | From b3f4b48e4f04c95f3ffdada7e890a29590f23786 Mon Sep 17 00:00:00 2001 From: Richard Palethorpe Date: Sun, 9 Aug 2026 20:58:27 +0100 Subject: [PATCH 08/17] perf(bench): admit request waves atomically (#206) Prepare and register complete AsyncLLM string/token waves before one all-or-zero core queue publish, and make the benchmark use that boundary for initial and refill admissions. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: Codex:gpt-5.6-sol [Codex] --- .agents/NOW.md | 4 +- .agents/engine-matrix.md | 2 +- .agents/roadmap_v1.md | 2 +- .../specs/sm120-qwen35-pareto-2026-08-09.md | 36 ++- docs/BENCHMARKS.md | 2 +- docs/STATUS.md | 2 +- docs/USAGE.md | 5 +- examples/bench/bench_core.h | 91 ++++-- include/vllm/v1/engine/async_llm.h | 52 ++++ include/vllm/v1/engine/core_client.h | 16 ++ include/vllm/v1/engine/core_proc.h | 24 +- include/vllm/v1/engine/output_processor.h | 8 + src/vllm/v1/engine/async_llm.cpp | 118 ++++++++ src/vllm/v1/engine/core_client.cpp | 17 ++ src/vllm/v1/engine/output_processor.cpp | 22 ++ tests/examples/test_bench.cpp | 84 ++++++ tests/vllm/v1/test_async_llm.cpp | 260 +++++++++++++++++- tests/vllm/v1/test_engine_core_proc.cpp | 136 +++++++++ 18 files changed, 849 insertions(+), 32 deletions(-) diff --git a/.agents/NOW.md b/.agents/NOW.md index f32f071aa..2e6bc0e55 100644 --- a/.agents/NOW.md +++ b/.agents/NOW.md @@ -21,7 +21,7 @@ Work: exact-chunks on main `1ce0d662b`; sm_120 measured at `3d2581551`. | MiniMax-H3 | **PRUNED ckpts RUN (#241): Q8_0 renders, seam 0.9941** | same-binary A/B | | Kimi-Linear-48B | 122/128 held; grouped router parallelised, e2e NOT ESTABLISHED | ckpt is tiktoken-only: no warm server | | 35B binding grid | @`a0fa12c7`: **flat 0.935-0.979x, NO c2/c8 weak cell** (CoV <0.81%); mem PSS 3.81x | Attribute the flat ~5% mid-band; TTFT c2 0.872x | -| Qwen3.5-4B sm_120 | #206 pretoken A/B FAIL: 98/128 exact; timings VOID. Sequential publish changes wave | Atomic admission; rerun | +| Qwen3.5-4B sm_120 | #206 pretoken A/B FAIL: 98/128 exact; timings VOID. Atomic wave CPU-green | Review; GPU A/B | | RPi5 A76 CPU | **R5 asm GREEN; llama NOT MET**: 0.461x pf, 0.653x dec | W6: BF16 GEMM | | MXFP4 parity | c1 1.020, c2-c8 0.962-0.969. **#82 CLOSED: ptxas-lineage REFUTED** | TERMINAL: at parity | | SERVE-ASYNC-DENSE-MIRROR | **LANDED+VERIFIED** (`f9c969ae`): async mirror, dense Qwen3; SACRED 184/184 | Sibling scope one-liner | @@ -53,7 +53,7 @@ latency/memory on every axis, both gate models, reproduced 2–3x idle. See of roof. Dense-marlin +0.5%; Triton-AOT GDN a WASH. 2. **Spike the Parakeet encoder row** (vLLM: `nano_nemotron_vl.py`; the transducer half is NOT in vLLM: separate call). -3. **Qwen3.5-4B #206:** atomic admission, review, then GPU/vLLM rerun. +3. **Qwen3.5-4B #206:** atomic admission CPU-green; review, GPU/vLLM rerun. 2. **Merge the invocation-parity prevention** (CI guard + AGENTS.md checklist); CUDA build-verify the byte-exact `kGemvHeuristicAlgos` refactor on dgx. 4. **Restore `local-ai-worker`** on dgx at campaign end (`--restart=always`). diff --git a/.agents/engine-matrix.md b/.agents/engine-matrix.md index fb8b32f87..6542bf70a 100644 --- a/.agents/engine-matrix.md +++ b/.agents/engine-matrix.md @@ -204,7 +204,7 @@ claims it. | `SERVE-HTTP-TRANSPORT` | Serving-socket transport parity: mirror vLLM's uvicorn/asyncio default `TCP_NODELAY` on every accepted SSE socket so per-token stream frames are not held by Nagle against the peer's delayed ACK. Implemented + CPU-tested; the non-binding localhost A/B sizing is COMPLETE and NEUTRAL within noise on c1/c2 ITL/TPOT/throughput (loopback ACKs are instant, so Nagle never coalesces ~100 ms-cadence token frames) — no gate-axis credit expected; the mirror stays for real-network parity. Future keep-alive / read-write-timeout / listening-socket option parity noted, not done | T0 | vLLM serves via uvicorn over asyncio `vllm/entrypoints/launcher.py:71,76`, `vllm/entrypoints/openai/api_server.py:591,630`; asyncio disables Nagle per accepted TCP stream socket `asyncio/base_events.py:192-197` (`_set_nodelay`) called from `asyncio/selector_events.py:950`; cpp-httplib default-off `third_party/httplib/httplib.h:142`, applied on accept only when set `third_party/httplib/httplib.h:12083` | `src/vllm/entrypoints/openai/api_server.cpp:69` (`set_tcp_nodelay(true)` in the ApiServer setup) | behavioral accepted-socket `getsockopt(TCP_NODELAY)` case `tests/vllm/entrypoints/openai/test_api_server.cpp:1076` (helper `:380`); RED accepted `TCP_NODELAY` 0 → GREEN 1, full `test_openai_api_server` **22/22 cases / 242 assertions**; non-binding sizing root `~/work/vllm.cpp-tcpnodelay-sizing/ff915e8…` (raw-set SHA `f5b52900…2128`) neutral within noise; closure [ledger](parity-ledger.md#L451) | [serve-tcp-nodelay.md](specs/serve-tcp-nodelay.md) | `DONE` | `ff915e8` | | `SERVE-C-ABI` | Stable LocalAI-style C FFI (**19** exported `VLLM_API` symbols at `VLLM_ABI_VERSION 10`; blocking and nonblocking request handles. Count corrected 2026-07-24 from a stale `17`, which predated ABI v4/v5 adding `tool_parser`/`reasoning_parser` and the chat entry points; `include/vllm.h` is the source of truth and README:231 already said 19). **ABI v9 2026-07-28 (`CLAIM-CAPI-ENGINE-CONFIG-V9`): the ABI carried strictly LESS engine config than `EngineParams` does** - `max_num_batched_tokens`, the scheduler `scheduling_policy` (`fcfs` / `priority` / `lpm`), and `kv_transfer_config` (the external KV connector / LMCache JSON) were reachable from the bundled server's flags and from NO embedder. All three added, inert at their defaults (zero-filled v8 growth == byte-identical pre-v9 engine); the connector NAME is validated against `KVConnectorFactory` at load, mirroring the server's startup check. `tokenizer_config_path` stopped being a declared-since-v1 no-op and now selects the chat template's source file. Malformed `speculative_config`/`kv_transfer_config` documents now report `VLLM_ERR_INVALID_ARGUMENT` (the contract vllm.h documented since v6) instead of `VLLM_ERR_MODEL_LOAD`, via a catch scoped to the parse block so a real `FromModelDir` failure still reports MODEL_LOAD. Driver: the LocalAI vllm-cpp backend could not expose LMCache or the prefill budget in a model config) | T0 | Original project ABI; pinned vLLM has no C ABI | `include/vllm.h:143,181,207`; `src/capi/vllm_c.cpp:229,264,327,391` | `tests/capi/test_capi.cpp:320,428,505,574,606,640`; `tests/capi/test_dlopen.cpp:77,86`; `tests/capi/c_header_compile.c:1` | [c-api-library.md](specs/c-api-library.md) | `ANCHOR-BACKFILL` | `CLAIM-SERVE-C-ABI-SPIKE` | | `SERVE-CPP-API` | Rich `LLM` and `AsyncLLM` C++ API | T1 | `vllm/entrypoints/llm.py:66,422`; `vllm/v1/engine/async_llm.py:70` | - | - | `planned: specs/cpp-api.md` | `INVENTORIED` | - | -| `SERVE-CLI-BENCH` | Serve and latency/throughput/serve benchmark modes | T0 | `vllm/entrypoints/cli/serve.py:44`; `vllm/entrypoints/cli/benchmark/main.py:29`; production queue `vllm/v1/engine/core.py:200-231,622-669`; pinned comparison pretokenizes before timing and synchronously admits a complete concurrency wave before each explicit step `tools/bench/vllm_closed_loop_metrics.py:57-102,137-167` | separate binaries + explicit scheduler-capacity flags `examples/server/main.cpp:63,96,116,170`; production `AsyncLLM` benchmark frontend + auditable scheduler depth; #206 default pre-encodes every prompt before `t0` and admits token IDs, exact `VT_BENCH_PRETOKENIZE=0` retains timed-string admission, and `BenchResult` reports the resolved path (`examples/bench/bench_core.h:190,212,582-586,698,780`). The current one-at-a-time publish changes EngineCore's first partial wave between string/token modes; additive atomic wave admission is specified but absent | server-help/production-frontend/metric assertions plus #206 parser/callback identity, preparation-before-clock, special-token/InputProcessor parity, report-mode capture and synthetic default/rollback exact prompt/output-ID A/B: `tests/examples/test_bench.cpp:123-212,316` 10/10, 70 assertions CPU-GREEN. Real `a33993a7` A/B **FAILED** token identity (98/128 requests, 15,507/16,384 positions); no timing credit. Atomic queue visibility, rollback/shutdown/duplicate, one-item, async exact-ID and real counterbalanced gates are pending | [CLI/serve/benchmark spike](specs/cli-serve-bench.md); [#206 campaign contract](specs/sm120-qwen35-pareto-2026-08-09.md) | `PARTIAL` | - | +| `SERVE-CLI-BENCH` | Serve and latency/throughput/serve benchmark modes | T0 | `vllm/entrypoints/cli/serve.py:44`; `vllm/entrypoints/cli/benchmark/main.py:29`; production queue `vllm/v1/engine/core.py:200-231,622-669`; pinned comparison pretokenizes before timing and synchronously admits a complete concurrency wave before each explicit step `tools/bench/vllm_closed_loop_metrics.py:57-102,137-167` | separate binaries + explicit scheduler-capacity flags `examples/server/main.cpp:63,96,116,170`; production `AsyncLLM` benchmark frontend + auditable scheduler depth; #206 default pre-encodes every prompt before `t0` and admits token IDs, exact `VT_BENCH_PRETOKENIZE=0` retains timed-string admission, and `BenchResult` reports the resolved path (`examples/bench/bench_core.h:190-233,616-665`). Atomic queue/wave admission prepares every request and collector before one all-or-zero ordered core publish (`include/vllm/v1/engine/core_proc.h:90-126`; `src/vllm/v1/engine/async_llm.cpp:125-237`) | server-help/production-frontend/metric assertions plus #206 parser/callback identity, preparation-before-clock, special-token/InputProcessor parity, report-mode capture and synthetic exact-ID A/B (`tests/examples/test_bench.cpp:123-186,347-376`). Atomic queue 3/3·13 (`tests/vllm/v1/test_engine_core_proc.cpp:233-299`); ordered/rollback/shutdown wave gates and complete async suite 12/12·433 (`tests/vllm/v1/test_async_llm.cpp:288-484`), all CPU-GREEN. Real `a33993a7` A/B **FAILED** token identity (98/128 requests, 15,507/16,384 positions); no timing credit. Fresh mutation review, operator gate and real counterbalanced retry remain pending | [CLI/serve/benchmark spike](specs/cli-serve-bench.md); [#206 campaign contract](specs/sm120-qwen35-pareto-2026-08-09.md) | `PARTIAL` | - | | `SERVE-GATE-ONLINE` | Same-corpus online correctness, TTFT/TPOT/ITL, throughput and peak-memory gate vs vLLM v0.25.0 | T0 | `vllm/benchmarks/serve.py:1,581-615`; [v0.25 audit](sync/2026-07-12-702f481.md); `tests/benchmarks/test_serve_cli.py:1` | Schema-v5 harness plus [trace controller](../include/vt/cuda/cuda_profiler_control.h#L13), [production component driver](../scripts/dgx-gdn-packed-component.sh), and fail-closed [component finalizer](../tools/bench/gdn_packed_component.py) | **BINDING `9ecd9d0`: 114/124** (async default ON; mem 4/4, c1 20/20, c2 20/20, c16 19/20, c4 & c32 18/20, c8 15/20; `benchmark_binding` refers here, superseding `3f256ab` 55/124 and `246a23c` 49/124, both retained immutable). Two-grid totality with `f0fb727` (111/124) is 115/124 effective parity vs vLLM 0.25.0 (27B). Async CLOSED the c16/c32 ITL tails (ours now BEATS vLLM: c16 p99 1.055, c32 p90 1.034/p99 1.078) and leaves a stable c8 `p99_itl` ~0.86 residual, ROOT-CAUSED (2026-07-18, `CLAIM-C8-P99-TAIL-1`, [spec](specs/c8-p99-itl-tail-2026-07-18.md)) as IRREDUCIBLE-AS-MIRRORED: our deterministic synchronous forward keeps co-admitted c8 requests in byte-identical lockstep where vLLM's async-future jitter de-phases them; the c16/c32 INVERSION proves this is the trailing edge of the per-step determinism that wins c16/c32 + throughput, not a capability gap (scheduler + async placeholder byte-identical, `tests/vllm/v1/test_scheduler_wave.cpp:265`, [tail spec](specs/tail-stall-analysis-2026-07-16.md)). Full grid + per-binding forensics: roadmap_v1.md + parity ledger; no packed speed credit | [online serving gate](specs/cuda-online-serving-gate.md); [merged GDN projections](specs/gdn-merged-input-projections.md); [packed decode](specs/gdn-packed-decode.md) | `ANCHOR-BACKFILL` | CLAIM-SERVE-GATE-1 | | `SERVE-E2E-NIGHTLY` | Server conformance and real-model nightly suites for all release gates | T0 | `tests/entrypoints/openai/`; `tests/v1/e2e/`; `.buildkite/test-pipeline.yaml` | current unit/conformance tests only; no scheduled DGX suite | `tests/vllm/entrypoints/openai/test_conformance.cpp:1`; `tests/parity/test_qwen36_paged_engine.cpp:78`; `tests/parity/test_qwen27_paged_engine.cpp:110` | `planned: specs/server-e2e-nightly.md` | `INVENTORIED` | - | | `ENG-RELEASE-BINARIES` | Downloadable host-ABI-specific `vllm-server` bundles: adaptive CPU and fat CUDA primary artifacts, optional per-SM diagnostics, and literal-static feasibility boundary | T0 | vLLM release lanes `.buildkite/release-pipeline.yaml:1-18,34-170` @ `555967922`; release-image dependency boundary `docker/Dockerfile.cpu:262-290` | Required W1-W11/W13 implementation is complete in draft PR #196: ten-SM gencode/AOT, adaptive CPU tiers, extracted-archive validation and supply chain, least-privilege immutable handoff, eight primary bundles, byte-derived indexes, attestation, and exact-file publication; W12 remains optional/non-primary | Local manifest/archive/workflow/mutation gates, full adaptive x86 tier execution, clean CPU archive, and Vulkan 35/35 backend + 11/11 cross-device archive run are green; hosted ten-SM completion, full eight-tuple dry run, matching-hardware gates, and tagged publication remain pending, so no published-binary claim exists | [release-binary-matrix.md](specs/release-binary-matrix.md) | `ACTIVE` | `CLAIM-ENG-RELEASE-BINARIES-W1-W13` | diff --git a/.agents/roadmap_v1.md b/.agents/roadmap_v1.md index 01fb2bed7..55a2bf0c6 100644 --- a/.agents/roadmap_v1.md +++ b/.agents/roadmap_v1.md @@ -101,7 +101,7 @@ models we already ship + benchmark. Full seam map + M0–M5 W-plan: | MEM | `ROAD-V1-MEM` | **Memory budgeting: auto-size to the declared workload by default, optional total-footprint cap, pre-flight error instead of an OOM (user-directed 2026-08-06, [#83](https://github.com/mudler/vllm.cpp/issues/83))** — the user-facing wart that every operator hits before they hit any perf question: vLLM makes you compute your own VRAM budget, express it as a PERCENT, and nail it or OOM | [coverage view §2](feature-matrix.md#2-kv-cache--memory), [porting inventory](porting-inventory.md) | — **M1+M2 LANDED 2026-08-08** ([`specs/kv-sizing.md`](specs/kv-sizing.md)): the absolute `--kv-cache-memory` knob sizes the pool via a group-aware `KVBytesPerBlock` divisor, `--num-blocks` is the override, both mirrored on the C ABI at v16; `ResolveNumBlocks` precedence `num_blocks > bytes > 256`, CPU-gated (`KVBytesPerBlock` 5/5 + `test_capi` v16). M3 (the `gpu_memory_utilization` profile run) stays dgx-gated. (M0 design grounded in vLLM `config/cache.py` + `gpu_worker.py:497-599`; GB10 unified-pool caveat carried) | `M1+M2 DONE` | **Source-verified 2026-08-06 (records-only, NO code).** WE ARE CURRENTLY BEHIND vLLM ON THIS AXIS, not ahead: there is NO memory profiling at all and the KV pool is a RAW BLOCK COUNT the user types by hand — `EngineParams::num_blocks = 256` (`include/vllm/entrypoints/model_loader.h:58`, beside `block_size = 32` `:57` / `max_model_len` `:59` / `max_num_seqs = 8` `:60`), exposed verbatim as `--num-blocks N` (`examples/server/main.cpp:100,203-204,370`), carried on the C ABI as `vllm_model_params.num_blocks` at the same 256 default (`src/capi/vllm_c.cpp:429,485`), landing as `BlockPool(num_gpu_blocks, ...)` which asserts `> 0` and otherwise TRUSTS it (`include/vllm/v1/core/block_pool.h:96,223`; `src/vllm/v1/core/block_pool.cpp:51`). So a user must convert "40 GB free, 32k context, concurrency 8" into a block count themselves — strictly worse ergonomics than a percentage. UPSTREAM HAS THREE KNOBS, all `config/cache.py`, all T0, all rowed at [porting-inventory.md](porting-inventory.md) §T0: `gpu_memory_utilization` (fraction of TOTAL, default 0.9), `kv_cache_memory_bytes` (absolute KV pool) and `num_gpu_blocks_override` (exact block pin), sized off a profile run as `total x utilization - non-torch - peak activation`. **Mirroring that is necessary but does NOT solve the complaint**, for three source-grounded reasons: (1) the fraction is of TOTAL not FREE, so on any shared card the right fraction is a function of what someone else already holds — exactly the arithmetic the user is being asked to do; (2) WEIGHTS LOAD BEFORE THE KNOB ENGAGES (utilization sizes the KV pool AFTER the model is resident), so an oversized model OOMs during load and never reaches the check — which is the failure operators actually hit; (3) 0.9 is taken whether or not it is needed (a 4B model on an 80 GB card reserves 72 GB it will never touch and blocks everything else on the device). THE DESIGN (user-ratified 2026-08-06) is three modes: **Mode 1 default = SIZE TO THE DECLARED WORKLOAD** — per-allocation-class accounting BEFORE allocating anything (weights from checkpoint metadata / safetensors header / GGUF manifest, known before reading a byte of tensor data; CUDA context measured at creation; peak activation from a profile run at `max_num_batched_tokens`; KV for `max_model_len x max_num_seqs` at the resolved `block_size`/KV dtype; CUDA-graph capture-set footprint) → allocate exactly that and LEAVE THE REMAINDER OF THE DEVICE FREE. This is the SURPASS over vLLM, which takes its 90% regardless of whether the workload needs 8 GiB or 80. **Mode 2 = a cap on the TOTAL ENGINE FOOTPRINT** (weights + activations + KV + graph pools + context), NOT on the KV pool alone — load-bearing, because a KV-only cap cannot prevent the weight-load OOM; three spellings of the same cap: `--memory-limit 40GiB` (absolute, the primary form), `--gpu-memory-utilization 0.85` (vLLM's exact flag name and fraction semantics so existing vLLM launch lines port unchanged, per [[mirror-vllm-always-no-asking]]) and `--num-gpu-blocks-override N` (upstream's reproducibility escape hatch — this is where today's `--num-blocks` GOES, demoted from primary knob to explicit override), with precedence spelled out and TESTED, not left to argument order. **Mode 3 = REFUSE BEFORE ALLOCATING** with the full per-class breakdown and remedies COMPUTED FROM THE ACTUAL BUDGET (`--max-model-len N` / `--max-num-seqs M` / `--kv-cache-dtype fp8` / smaller quant, each with the GiB it recovers) — "you are 43.9 GiB over and here are the three levers that close it" is the difference between a usable error and a stack trace. **UNIFIED-MEMORY HAZARD (not hypothetical):** on GB10 the ~119 GiB pool is UNIFIED, a fraction-of-total setting reserves HOST RAM too, and `gpu_memory_utilization=0.85` has HARD-REBOOTED our DGX three separate times ([[gb10-unified-memory-oom-reboots-box]]) — hence absolute bytes is the PRIMARY form with the percentage kept only for vLLM compatibility, and hence the accounting must know whether the pool is unified, which makes free/total + an is-unified predicate a PLATFORM-SEAM question belonging behind `ROAD-V1-C1`'s abstraction (note `Platform::needs_weight_staging()` was deliberately NOT `is_unified_memory()` because the latter FLIPS GB10 — the distinction matters here) rather than a CUDA-specific branch, since discrete and unified devices need different safety margins. CORRECTNESS: pool size changes preemption/scheduling TIMING but not emitted tokens, so the SACRED token-exact gates are unaffected — and M2's gate makes that explicit by re-running them with no block flag at all. **Next gate = M0 the `specs/kv-sizing.md` spike (accounting model + precedence rules + upstream `file:line`); then M1 a `MemoryBudget` computing required bytes per class WITHOUT allocating plus the Platform free/total + is-unified seam, unit-gated predicted-vs-actual weight bytes; M2 auto-sizing as the default with `--num-blocks` demoted to `--num-gpu-blocks-override`, gated by every existing model gate running with NO block flag and staying token-exact; M3 the three caps + precedence through the server flags and the C ABI, gated by our KV pool matching vLLM's own at a matched `--gpu-memory-utilization`; M4 the pre-flight refusal, gated by a deliberately over-subscribed config exiting cleanly on GB10 (non-zero exit, no OOM, NO BOX REBOOT) covering the unified-pool path specifically; M5 (optional) a runtime guard failing the REQUEST rather than the engine.** Docs (README, [STATUS](../docs/STATUS.md)) update in the SAME change as whichever milestone shifts externally-visible behaviour, per [[keep-readme-current]] | | 1 | `ROAD-V1-C1` | **Extensibility-first (USER PRIORITY 2026-07-18):** drop-in kernel ABI + the MISSING Platform seam + model self-registration — make new GPUs/archs/models ADDITIVE (plan: [extensibility-platform-seam-2026-07-18.md](specs/extensibility-platform-seam-2026-07-18.md)) | [`BACKEND-ABI-VT`](backend-matrix.md), [kernel matrix](kernel-matrix.md) | exhaustive kernel/dependency inventory and [raw-pointer adapter ABI](specs/dropin-kernel-abi.md) accepted; additive W0 implemented and CPU 94/94. `CLAIM-BACKEND-ABI-W0-GPU-1` repaired the GCC13/doctest blocker without runtime changes; exact sm_121a all-target build, focused CUDA/ABI sanitizer, and both gate-model tests pass at `1141b79`. Cross-arch/trace/A-B and scalar-forwarder/backend-shim debts remain explicit | `PARTIAL` | **★ NEW ORDER-1 HEAD (user-directed 2026-07-19): the PORTABLE AUTOMATIC OP-FUSION FRAMEWORK (`KERNEL-FUSION-FRAMEWORK`, spike [portable-fusion-framework.md](specs/portable-fusion-framework.md), `SPIKE`).** The extensibility cornerstone: fusions DECLARED ONCE (backend-agnostic `constexpr FusedRecipe` catalog above `vt::`, transcribing vLLM's finite pattern-pass set `passes/fusion/*` @ `pass_manager.py:138-200`, mirroring the `CustomOp` `forward_native`/`forward_cuda` seam `custom_op.py:103`) and REALIZED PER-BACKEND through the existing `vt::` op table (Tier-0 composite = the CPU oracle inherited free by any backend; Tier-1 interpreter = one kernel port per backend lights up every recipe). Makes a new vLLM fusion PR a ONE-declaration port, a new GPU an additive catalog realization, a new model an additive pattern declaration — the PR-#4 remedy, composed with the Platform/attn-registry/model-registry seams below. The TDR Phase-0 skeleton is already LANDED (`fused_recipe.h`/`recipes.h` one recipe + `OpId::kFusedChain` Tier-0/1 on CPU+CUDA + byte-exact `test_ops_fused_chain.cpp`). **W0 ADOPTED 2026-07-19 (`CLAIM-FUSION-FRAMEWORK-W0`):** the seam is now used in production at ONE real site — the 35B `RunLayerPaged` post-attention layernorm routes its plain add+residual+gemma-RMSNorm through `vt::FusedChain(kFusedAddRmsNorm)` (`VT_FUSED_CHAIN_ADOPT` default-ON / `=0` rollback), behaviour-preserving + byte-identical to the prior hand-call (DGX: clean CUDA `-Werror` 0 warn, byte-exact composite==interp==golden incl. H=2048, 35B 315/315 + 27B 235/235 token-exact BOTH arms, memcheck 0 errors) — proving the declare-once/realize-per-backend seam end-to-end; the current 3-opcode/4-role POD sufficed byte-identically for the plain add+rmsnorm class, so W0 needed NO generalization. **W1 GENERALIZED the POD 2026-07-20 (`CLAIM-FUSION-FRAMEWORK-W1`, `1115648`):** full activation/norm/quant/rope opcode set + indexed operand table; all 5 quant-fused chains declared byte-exact; Tier-0 composite = ONE device-agnostic walker (kills CPU/CUDA oracle drift); infrastructure only, no call site changed (DGX: `-Werror` 0-warn, byte-exact CPU 196 + CUDA 361, memcheck 0, 27B 235/235 + 35B 315/315 both arms). **W2 MIGRATED the hand-fusions 2026-07-20 (`CLAIM-FUSION-FRAMEWORK-W2`):** the framework now OWNS the fusion dispatch — each recipe binds (new backend-agnostic `FusedRecipe.fast_op`) to its EXISTING single-launch fused kernel, so `FusedChain(recipe)` dispatches to the SAME fast kernel the model called directly pre-migration (byte-identical + perf-neutral by construction; composite is the graceful fallback + oracle). SIX call sites migrated behind `VT_FUSED_CHAIN_ADOPT` (`kSiluMulFp4Quant`/`kSigmoidGateFp4Quant`/`kRmsNormGatedQuantFp8`×2/`kRmsNormQuantFp8`/`kAttnQkNormRopeGate`×2). DGX: `-Werror` 0-warn, byte-exact CPU 228 + CUDA 420, memcheck 0, 27B 235/235 + 35B 315/315 BOTH arms. **W3 MECHANICAL-SYNC PROOF LANDED 2026-07-20 (`CLAIM-FUSION-FRAMEWORK-W3`):** ported a NEW, previously-unported vLLM fusion pass — `SiluMulFp8StaticQuantPattern` (`act_quant_fusion.py:81` → `_C.silu_and_mul_quant`, the static-per-tensor-FP8 sibling of `kSiluMulFp4Quant`) — as ONE `constexpr FusedRecipe kSiluMulQuantFp8` + its byte-exact test, touching EXACTLY 2 shared files (`recipes.h` + `test_ops_fused_chain.cpp`), NO kernel/dispatch/composite-walker/model-site edit and NO new primitive (composite = existing `vt::MoeSiluMul` + `vt::QuantFp8Static`; `fast_op=kNoFastOp`). The PR-#4 additivity test made concrete: a whole new fusion pattern = one declaration. DGX: `-Werror` 0-warn, byte-exact CUDA 432, memcheck 0, no token regression (recipe declared-only) 27B 235/235 + 35B 315/315. **W4 BACKEND-ADDITIVITY PROOF LANDED 2026-07-20 (`CLAIM-FUSION-FRAMEWORK-W4`) — the W-series proof milestone is DONE:** made the additivity claim EXECUTABLE — new test `test_fused_chain_additivity.cpp` treats the EXISTING CPU backend AS the 'second backend' relative to CUDA (no mock `DeviceType` — that would edit the core enum + every switch, ironically non-additive) and, in ONE generic loop over the WHOLE catalog (all 7 recipes), asserts each runs byte-exact on the CPU backend via the Tier-0 composite — 4 CPU-full end-to-end + 3 fp8-prefix (byte-exact prefix + the FULL composite asserted to THROW on CPU, documenting the CUDA-only static-fp8 backend-negotiated tail, §3b/§6). Additivity evidence: catalog `recipes.h` grew 1→6→7 while the composite walker stayed ONE per-OPCODE function (12 `FOp::` cases) + the CPU/CUDA `kFusedChain` registration ONE line each + `cpu_ops.cpp` never `#include`s `recipes.h` — W3's whole new recipe `kSiluMulQuantFp8` is in ZERO backend TUs, inherited free. CPU `-Werror` 0-warn, `test_fused_chain_additivity` 17/17 + `test_ops_fused_chain` 228/228; engine byte-identical (no `src/`/`include/` change) ⇒ 27B 235/235 + 35B 315/315 structurally unchanged; memcheck N/A. Honest deferred (named, non-blocking the ORDER-1 milestone): Tier-1 perf interpreter for the quant chains (composite-only today), a REAL Metal/Vulkan catalog realization (M4 HW-blocked), and per-recipe fast single-launch kernels. Honest payoff: perf ceiling ~3.5%/step compute-bound on 35B (NOT a perf lever — tasks #61/#62; W0 is perf-neutral by construction); primary value = extensibility + mechanical upstream-sync + CPU/CUDA oracle-drift elimination. Incremental W0 adopt-one **(DONE)** → W1 generalize POD **(DONE)** → W2 migrate hand-fusions **(DONE)** → W3 mechanical-sync proof **(DONE)** → W4 backend-additivity proof **(DONE)** → Wn honest re-measure (optional, off the extensibility critical path). **W-SERIES ORDER-1 PROOF MILESTONE DONE 2026-07-20 (`CLAIM-FUSION-FRAMEWORK-W4`).** **PRIOR extensibility items (all LANDED, the seams this composes with):** **#1 extensibility item — extract the Platform seam — LANDED 2026-07-18** (`BACKEND-PLATFORM` `ACTIVE`, `CLAIM-BACKEND-PLATFORM-1`): `include/vllm/platforms/interface.h` + `src/vllm/platforms/{platform,cpu,cuda}.cpp` mirror `vllm/platforms/interface.py:134-229` 1:1; `CurrentPlatform()` self-registered per `DeviceType`; the 7 memory-model/residency `device.type == kCUDA` sites (of PR #4's ~37) now route through it → new-GPU memory model is ONE additive `platforms/.cpp`. Behavior-preserving (clean CPU build + `test_platform` + full CPU CTest green; DGX 235/235 + 315/315 pending). **Item 2 residency-as-Platform-capability LANDED 2026-07-19** (`CLAIM-BACKEND-PLATFORM-2`): the host-free / load-stream / DevicePool-cap decisions in `qwen3_5.cpp` now READ `GetPlatform(.device.type).residency_policy()` (per-device) via the pure `ShouldReleaseHostWeights`/`ShouldInterleaveLoadStream` helpers + `device_pool_cap_bytes`, not an inline `device.type`/env gate; `CudaPlatform.release_host_weights_after_upload` flipped false→true (now CONSUMED ⇒ reproduces today's GB10 host-free-after-Marlin + ~4 GiB load peak EXACTLY); `MarlinMoeEnabled()` stays the orthogonal kernel-path gate. **A new (discrete) GPU sets `residency_policy()` values ⇒ ZERO model edit.** Behavior-preserving (clean CPU build + `test_platform` consumption cases 7/43 + full CPU CTest + tools 164/164 green; **DGX-CONFIRMED @ `62fc0e0`: clean CUDA `-Werror`, 27B 235/235 + 35B 315/315 token-exact, 35B VmHWM ≈ 4.0 GiB load-stream win preserved, memcheck 0 errors**). Then item 3 drop-in ABI family migration. **Item 4 attn-backend registry LANDED 2026-07-19** (`CLAIM-ATTN-REGISTRY-1`, `BACKEND-ATTN-REGISTRY`): NEW `include/vllm/v1/attention/registry.{h,cpp}` `(DeviceType,name)` registry + `SelectAttentionBackendName` selector (mirror `registry.py` self-registration + `cuda.py:361-470` `get_attn_backend_cls`/`_get_backend_priorities`); `Platform::get_attn_backend_priority()` filled (was the item-1 STUB) → capability-ordered name lists on `CudaPlatform` (major-10 vs else) + `CpuPlatform`; FLASH_ATTN/GDN self-register. **Adding a backend's attention = 1 self-registering TU + 1 priority slot, ZERO selector/model/runner edit.** Behavior-preserving — the walk returns FLASH_ATTN on CUDA+CPU (the same FA2 attention runs); clean CPU build + `test_attn_backend_registry` (8/25) + full CPU CTest, tools 164/164, checkers green; **DGX-CONFIRMED @ `2c732e7`: 27B 235/235 + 35B 315/315 token-exact (FA2 sm_121a), memcheck 0/315**. **Item 5 model self-registration LANDED 2026-07-19** (`CLAIM-MODEL-SELFREG-1`, `MODEL-FACTORY-registry`): the fixed `kRegistrations` array → `REGISTER_VLLM_MODEL(...)` static-`Registrar` idiom (`model_registry.h:167-189`) + Qwen dense/MoE arch entry points split into per-variant TUs (`qwen3_5_dense.cpp`/`qwen3_5_moe.cpp`) over shared `qwen3_5_common.{h,cpp}`, so **adding a model = 1 new TU + 1 REGISTER line, ZERO shared-array edit**; behavior-preserving (clean CPU build + `test_model_registry` extension + full CPU CTest, tools 164/164, checkers green; DGX 27B/35B token-exact pending). Deep `qwen3_5.cpp` machinery factoring deferred. Metal/MLX bring-up proves the seams (needs M4). **★ THE ARCH HALF OF THIS ITEM IS NOW PROVEN BY MEASUREMENT, NOT ARGUED (2026-07-22, `CLAIM-CUDA-SM120-BRINGUP`, [spec §W8](specs/cuda-arch-additivity.md)):** consumer-Blackwell `sm_120a` was brought up as a BUILD-supported target and required **ZERO kernel, model, runner, sampler or feature-table edits** — the additive seams (per-arch FEATURE TABLE, capability-keyed tactic registry keyed on `major == 12`, Platform auto-probe, `pageable && integrated` residency classification) already covered it, so the entire diff is build configuration, a configure-tier test and records. That is the PR-#4 additivity test passed on a real second architecture. It is deliberately NOT a runtime-support claim: no sm_120 board exists here. **★ THE MODEL/QUANT HALF NOW ADVANCES TOO (2026-07-23, `CLAIM-BACKEND-SEAM-S4-1`): the `model_executor/layers/` `LinearMethod`/`QuantizationConfig` seam the [accelerator-seam audit](specs/accelerator-seam-audit.md) §9 named ABSENT now EXISTS in part.** `S4` landed byte-identical: the dense model's projections route through a `method.Apply()` chosen ONCE from the checkpoint (retiring the per-call `IsNvfp4()` tensor-name probe), and 18 shared-layer `device==kCUDA` availability gates became `vt::OpRegistered` op-table queries — the policy(scheme)/implementation(kernel) split the audit's binding rule requires. **DSR 86 → 67**; all six SACRED gates byte-identical on dgx (27B/35B/Coder/dense/OPT/DeepSeek-V2); the fragile 27B-W4A4/fp8-recipe gates are correctly deferred to `S6` behind `S5`'s reference tier. **★ `S6` ASSESSED 2026-07-23 (`CLAIM-BACKEND-SEAM-S6-1`) → NO-OP / BLOCKED, DSR stays 67 (§11):** the deferred fp4/fp8 gates convert ZERO sites byte-identically — every one bottoms out at a **dual-registered** (CPU+CUDA) bespoke op (none CUDA-only, unlike S4's convertible gates), so `OpRegistered(op,dev)` is TRUE on `kCPU` ⇒ the class-A swap is bit-changing on the CPU reference/emulation path (two numerics per device); S5's reference tier does not change this (those CPU kernels are present natively, never a miss). No `src/`/`include/`/test byte changed, no baseline moved. The genuine byte-identical unlock is re-scoped to **`S3`** (Platform capability fields mirroring `supports_fp8`/`cutlass_fp4_supported` — the audit's own class-D fix) and **`S7`** (layer extraction); the plan's `~37` S6 target assumed the class-A `OpRegistered` swap was byte-identical, which holds only for CUDA-only ops (all taken by S4). **★ `S3` LANDED 2026-07-23 (`CLAIM-BACKEND-SEAM-S3-1`) — the byte-identical unlock S6 re-scoped:** mirrors vLLM's `Platform` capability surface (`supports_fp8`/`cutlass_fp4_supported`/`opaque_attention_op`/`is_integrated_gpu`/`support_static_graph_mode`/`is_device_capability_family`, base false in `interface.h`; CudaPlatform answers GB10 values in `cuda.cpp`/`platform.cpp`) and converts **12** deferred `qwen3_5.cpp` gates onto it (7 fp4-act `cutlass_fp4_supported`, 3 fp8-fused `supports_fp8`, 2 decode-graph `support_static_graph_mode`) — byte-identical because a capability answers the base false off CUDA, exactly what `device==kCUDA` did (where S6's `OpRegistered` was TRUE on `kCPU`), and it DECOUPLES (a future accelerator answers for itself). **DSR 67 → 55** (`kcuda` 25→13), baseline lowered same commit, ratchet + 24-case suite green; all seven SACRED gates byte-identical on dgx (27B 235/235 · 35B 315/315 · Coder · dense-32B · OPT · DeepSeek-V2 · Llama), `test_platform` CUDA-leg proves each predicate == former `device==kCUDA` on GB10, memcheck 0 errors, clean CUDA+CPU `-Werror`. Residency/stream/FA2-dtype/merged-layout sites LEFT for `S7`. **★ `S7` LANDED 2026-07-23 (`CLAIM-BACKEND-SEAM-S7-1`) — the seam campaign's TERMINAL runtime-decoupling state, the closest this extensibility work comes to a finish line:** ALL 23 remaining runtime `kCUDA`/`is_cuda()` sites in the shared model layer hoisted onto capabilities — new `Platform::needs_weight_staging()` (the CUDA device-resident staging policy, NOT `is_unified_memory()` which would FLIP GB10; covers residency/merged-GDN/packed-decode/direct-load), `Platform::supports_fa2_attention()` (FA2 dtype), `Backend::SupportsAuxStream()` (MoE aux-stream), reusing S3's `supports_fp8`/`cutlass_fp4_supported`/`support_static_graph_mode`/`is_integrated_gpu` (runner combine/scatter) and `vt::OpRegistered(kMoeGroupedGemmBf16)` (a CUDA-only op). Each returns the former `device==kCUDA` value on GB10 → byte-identical. **DSR 55 → 32 — the IRREDUCIBLE build-gate floor:** the shared model layer holds ZERO runtime device tests; the 32 residual are all `#ifdef VT_*` compile-time gates for kernels that only build on one GPU family (a kernel that only compiles on one arch is legitimately irreducible), so the audit's `<10` is NOT reachable and this is the honest answer to "how additive can the shared layer get" — every runtime device coupling is gone. baseline lowered same commit, ratchet + 24-case suite green; all seven SACRED gates byte-identical on dgx (27B 235/235 · 35B 315/315 · Coder 6/6 · dense-32B 16/16 · OPT 6/6 · DeepSeek-V2 8/8 · Llama 16/16), new `test_platform`/`test_backend`/`test_cuda_backend` cases green, memcheck 0, clean CUDA+CPU `-Werror`. The `layers/`-library physical relocation (shrinking `qwen3_5.cpp` toward `qwen3_next.py`'s 802-line shape) is a follow-on refactor; the device coupling it was to remove is already gone | | 2 | `ROAD-V1-C2` | Model families: Llama/Qwen3/Mistral, MoE, Qwen3-Next | [model matrix](model-matrix.md) | current pin has 353 static IDs; v0.25.0 adds three sync-target rows (MOSS-Transcribe-Diarize, Laguna DFlash, Bailing hybrid MTP), yielding 356 after pin advance. **FIRST ADDITIVE-MODEL BRING-UP W0-W4 LANDED 2026-07-20 — CORRECTNESS COMPLETE (0.6B + 4B gates PASS 16/16; SPEED pending)** ([first-additive-model-qwen3-dense.md](specs/first-additive-model-qwen3-dense.md), `MODEL-TEXT-qwen3-qwen3-for-causal-lm` `ACTIVE`(correctness COMPLETE, speed pending), runner generalization `ENG-RUNNER-MODELSHAPE`, `CLAIM-MODEL-QWEN3-DENSE`) **MLA CAMPAIGN SPIKED 2026-07-21** ([mla-deepseek-campaign](specs/mla-deepseek-campaign.md), `CLAIM-MLA-DEEPSEEK`): five rows `INVENTORIED` -> `SPIKE` (DeepSeek V2 / V3+V3.2 / v1-MHA, Kimi-Linear, MiniMax-M2). **KIMI-LINEAR-48B W0 DEDICATED SPIKE 2026-08-05** ([kimi-linear.md](specs/kimi-linear.md), `CLAIM-KIMI-LINEAR-W0`): full dedicated W0 spike for `MODEL-TEXT-kimi-linear-*` (stays `SPIKE` — actively claimed) — the ONE Kimi text model that FITS one GB10 (91.5 GiB, 0.77x pool) with a real e2e SACRED gate; HEAVY reuse (MLA + sigmoid/`noaux_tc` MoE + GDN family + KDA host refs landed), NET-NEW = the KDA device kernel + NoPE-MLA branch + hybrid schedule/loader; W1 implementation can start. Answers the Tier-3 "MLA = new attention, new campaign" item in [breadth-sweep-plan](specs/breadth-sweep-plan.md) §B.3. Key determinations: GB10/sm_121 selects **`TRITON_MLA`** for dense MLA decode and **`FLASH_ATTN`** for MLA prefill (`vllm/platforms/cuda.py:129-133`, `mla/prefill/selector.py:74-77`), so the sm90/sm100-only MLA kernels are out of reach and out of scope; the cross-cutting cost is the **compressed-latent KV cache** (one 576-wide latent per token, `num_kv_heads=1`, no separate V), which our allocator and `vt::ReshapeAndCache`/`vt::PagedAttention` cannot express; and **only DeepSeek-V2-Lite (~29.3 GiB bf16) fits GB10** — V3/V3.2, Kimi-K2.5, MiniMax-M2/M3 are HW-BLOCKED e2e, Kimi-Linear-48B is HW-MARGINAL. W0-W10 plan recorded; nothing implemented. **GLM + DSA + LATEST-DEEPSEEK SPIKED 2026-07-21** ([glm-dsa-latest-deepseek](specs/glm-dsa-latest-deepseek.md), `CLAIM-GLM-DSA-LATEST-DEEPSEEK`): seven rows `INVENTORIED` -> `SPIKE` (ChatGLM, Glm, Glm4, Glm4Moe, Glm4MoeLite, GlmMoeDsa, DeepSeek-V4). Answers the user's "also glm, and deepseek latest versions" priority. Headline: **`Glm4MoeLiteForCausalLM` / `zai-org/GLM-4.7-Flash` (31.2B, 58.2 GiB bf16) FITS GB10 and is a SECOND MLA gate vehicle that closes BOTH coverage gaps the MLA campaign named as unit-gated-only** (it has `q_lora_rank=768` and `noaux_tc`/`e_score_correction_bias`, which DeepSeek-V2-Lite lacks). **DSA is DOUBLY BLOCKED on GB10:** the sparse XOR filter eliminates `TRITON_MLA` for sparse models leaving `FLASHINFER_MLA_SPARSE_SM120` as the sole candidate, and that path is non-functional on flashinfer 0.6.12 (XQA backend is dense-only, discards `sparse_mla_top_k`); GLM-5 is 1404 GiB and V3.2 642 GiB regardless. DeepSeek-V4 is a NEW architecture (Sinkhorn-normalized Manifold Hyper-Connections, CSA/HCA compressor with recurrent state, hash-routed MoE) and HW-BLOCKED at 148.7 GiB — but its TOKENIZER risk is LOW (standard HF fast BPE; only the chat template needs porting, with upstream golden fixtures). Glm4/Glm need two primitives we have NONE of: partial rotary factor and sandwich norms. Nothing implemented. **NVFP4A16 (W4A16)** on the already-done dense `Qwen3ForCausalLM` (`RedHatAI/Qwen3-32B-NVFP4A16`, 64L) — the QUANT-SCHEME additivity experiment, serving user priorities #2 (models) and #4 (quants) at once. KERNEL LAYER FULLY ADDITIVE (ZERO new kernel code: vLLM FORCES Marlin for `use_a16`, OBSERVED `Using MarlinNvFp4LinearKernel`, and that is the GEMM we already vendored for the 35B). **CORRECTNESS CLOSED 2026-07-21 (W4b):** the strict gate's 4/6 was diagnosed by the ratified TEACHER-FORCING isolation — all 29 divergent positions gap <= 0.0625 nats with 28/29 EXACTLY 0.0, one root flip an EXACT bf16 tie at which vLLM's teacher-forced argmax is OURS and vLLM contradicts its own greedy. **NOT a W4A16 defect: it is the pre-existing dense-forward bf16 near-tie drift, recorded against `MODEL-TEXT-qwen3-qwen3-for-causal-lm`.** Gate closes **6/6** under the ratified near-tie-robust bar with the nats evidence committed. SPEED still pending ⇒ `ACTIVE`, not `DONE`. **GEMMA FAMILY SPIKED 2026-07-24** ([sweep-gemma](specs/sweep-gemma.md), `CLAIM-SWEEP-GEMMA`): four rows `INVENTORIED` → `SPIKE` (Gemma 1/2/3/4). Answers the user's "and then we do gemma" ("gemma 4") next-target. **The newest registered Gemma is Gemma 4** (real, public, but ALL checkpoints multimodal-wrapped `Gemma4*ForConditionalGeneration`, ≥12B, 0.25.0 oracle-support unverified, needs a PLE/YOCO/MoE/k_eq_v stack) — it leads the characterization but is gate-BLOCKED as a first vehicle. **The recent-first gate vehicle that FITS + is oracle-certain is Gemma 3** (`Gemma3ForCausalLM` on `google/gemma-3-1b-it`). Headline: Gemma reduces MOSTLY to landed infra — gemma-RMSNorm, sandwich norms (glm4 `b568d20`), SentencePiece (names "Gemma"), sliding-window (FA-2 + SlidingWindow/ChunkedLocalAttention specs), the `kAttnQkNormRopeGate` QK-norm+rope recipe, tied embeddings are ALL REUSE; the one genuinely-new compute kernel is GeGLU (`gelu_pytorch_tanh`+mul; we have only SiLU), plus the final logit soft-cap + qpas/embed-scale scalars + dual-rope routing. Per-version delta: Gemma-2 has an attn logit soft-cap, Gemma-3 removed it for QK-norm. **GEMMA-3 W0-W2 LANDED 2026-07-24 — CORRECTNESS COMPLETE, the FIRST Gemma family** (`MODEL-TEXT-gemma3-gemma3-for-causal-lm` `ACTIVE`, speed pending): `Gemma3ForCausalLM` on `google/gemma-3-1b-it`. W1 = two additive default-inert vt ops `kGeluAndMul` (GeGLU `gelu_pytorch_tanh`) + `kMulScalar` (bf16 embed-scale), CUDA+CPU, unit 12/12. W2 = `gemma3.{h,cpp}`/`gemma3_weights.cpp`/`gemma3_registry.cpp` reusing the GLM-4 sandwich-norm layout + `dense_attn_block.h` + FA-only KV: GemmaRMSNorm `(1+w)`, per-head Gemma q/k norm, dual per-layer RoPE theta, qpas scale, per-layer sliding window, GeGLU MLP, `sqrt(hidden)` embed-scale, tied lm_head. **SACRED gate STRICT token-exact 48/48** greedy vs vLLM 0.25.0 (K=5 ALL-DETERMINISTIC → STRICT; BOS-verified; tokenizer-free like Mistral's `LOAD-SENTENCEPIECE` path). Loader 340 tensors, registry 23/23, clean `-Werror` 0 warn. **GEMMA-2 + GEMMA-1 W3-W6 LANDED 2026-07-24 — CORRECTNESS COMPLETE** (`MODEL-TEXT-gemma2-gemma2-for-causal-lm` + `MODEL-TEXT-gemma-gemma-for-causal-lm` `ACTIVE`, speed pending): W3 = the logit soft-cap primitives (`vt::SoftCap` final cap + `PagedAttentionArgs.logits_soft_cap` attention cap threaded into the native/CPU/FA-2 attention, default-inert). W4 `Gemma2ForCausalLM` (gemma-2-2b-it) = the inverse of Gemma-3 (BOTH soft-caps, NO QK-norm, single rope) — **near-tie-band SACRED 48/48** (44/48 strict + 4/48 at 0.0-nat vLLM-own ties, 0 forward-divergent; soft-cap PROVEN applied by a cap-on≠cap-off A/B). W5 `GemmaForCausalLM` (gemma-2b) = the original Gemma (two fused norms, head_dim scale) — **STRICT 48/48**. W6 = Gemma-4 honesty pass (HW/DEP-BLOCKED, not registered). Regressions byte-identical (Gemma-3 48/48, Qwen3-dense 184/184, OPT 63/63, Llama 92/92, Mistral 92/92) + DeepSeek-V2 asserts-on 223/223; compute-sanitizer 0; clean `-Werror` 0 warn. Gemma-4 stays `BLOCKED`. | `PARTIAL` | **ACTIVE: the first additive-model bring-up = Qwen3 dense (`Qwen3ForCausalLM`) on `Qwen3-0.6B` BF16 — W0+W1 landed 2026-07-20.** W0 (config+registry stub: new TU `qwen3_dense.cpp`+`qwen3.h`, one `REGISTER_VLLM_MODEL`, full-attention-only KV spec, forward stub) + W1 (the RUNNER GENERALIZATION `ENG-RUNNER-MODELSHAPE`) are DONE and gated: dgx CUDA `-Werror` 0-warn, **27B 235/235 + 35B 315/315 token-exact UNCHANGED** (behaviour-preserving), new CPU runner tests RED(SIGSEGV)→GREEN, registry resolves `Qwen3ForCausalLM`, ASan/UBSan clean. The runner is now MODEL-SHAPE-AGNOSTIC (a full-attention-only KV config allocates+steps without the hybrid GDN path) → every future dense/non-hybrid arch adds new-files-only. Qwen3-0.6B is the only standard-dense arch with a checkpoint + runnable vLLM 0.25.0 oracle on dgx TODAY (no Llama/Mistral checkpoint present → Llama-first needs a download, sequenced as W-next for genuine cross-family additivity). **W2 loader + W3 forward LANDED 2026-07-20:** dense forward `qwen3.cpp` (`Qwen3DenseModel::Forward/ForwardDevice`) composed from vt:: ops + the fusion catalog (2 new byte-exact recipes: `kFusedAddRmsNormStd`, `kAttnQkNormRope`); bf16 attention numerics mirror vLLM. The first pure-dense bf16 model forced out + FIXED 2 genuine latent bugs: tokenizer `kQwen2Classic` (classic Qwen2/Qwen3 pre-tokenizer was hard-rejected) and `cuda_paged_attn.cu` WMMA prefill mistokenizing at head_dim≠256 (now gated to the validated d=256). **W4 CORRECTNESS COMPLETE 2026-07-20 — near-tie-robust gate PASSES on 0.6B AND a bigger 4B.** The 2026-07-20 razor's "vLLM greedy non-deterministic" premise was a BATCHING artifact: per-prompt (batch=1, the gate regime) vLLM 0.25.0 greedy is DETERMINISTIC (0.6B 0-multi/K=10, 4B 0-multi/K=5). Forward correctness is PROVEN by teacher-forcing vLLM on OUR exact prefix (`scripts/qwen3-neartie-gap.py`): at all-but-2 positions vLLM's own argmax given our prefix IS our token (gap 0.0000, bit-identical logprobs — our forward matches vLLM's prefill logits); residual flips are bf16 near-ties (0.6B ≤0.125 nats, 4B ≤0.25) where vLLM's own prefill argmax disagrees with its decode. Gate = our token within 0.5 nats of vLLM's teacher-forced argmax (strict where equal): **Qwen3-0.6B 16/16** (strict 12 + near-tie 4) and the **bigger-model complete-correctness proof Qwen3-4B (36L, GQA 32/8, hidden 2560, same forward code) 16/16** (strict 10 + near-tie 6). Regression 27B 235/235 + 35B 315/315 UNCHANGED, CUDA `-Werror` 0-warn, memcheck 0. Correctness-complete. **SPEED — d128 FA2 PREFILL + DECODE DEFAULT-ON 2026-07-20 (`Qwen3-4B` vs vLLM 0.25.0 production/graphed, in1024/out128) — big gap-close, still below vLLM, `MODEL-TEXT-qwen3-qwen3-for-causal-lm` stays `ACTIVE` NOT `DONE`:** implemented the dominant prefill lever (a d128 FlashAttention-2 varlen prefill — generalized the vendored FA2 launcher to head_dim 128, `VT_FA2_PREFILL_QWEN3` default-ON) and flipped the FA2 varlen d128 decode default ON (near-tie gate re-passes 16/16 on 0.6B + 4B). Total tput now 0.90× (c1)/0.62× (c8) (up from 0.80×/0.48×), c1 decode at parity (TPOT 1.04×, ITL P99 0.98× win); prefill A/B = +7%/+41% total, −55%/−48% TTFT. STILL failing TTFT median 5.85×/10.2× + total <1×: the full prefill STEP (not the attention kernel, now vLLM's FA2 family) is ~6× vLLM = non-attention glue (GEMM/MLP fusion) + host-side launch overhead (un-graphed prefill); plus c8 decode batch efficiency (TPOT 1.38×). Dominant residual lever = portable prefill-step fusion + graphed prefill (secondary = c8 split-KV decode occupancy). **RoPE cos/sin cache flipped DEFAULT-ON 2026-07-20** (`VT_QWEN3_ROPE_CACHE`): the opt-in blocker (an alleged FA2-split-KV-combine run-to-run nondeterminism) was GROUNDED + DISPROVEN — the paged engine is byte-deterministic run-to-run and the combine never launches on the gate (`num_splits==1`); goldens regenerated on the canonical `$HOME/cutlass-4.5.0` build (the flashinfer cutlass copy tips the 27B tok6 razor to 234/235; cutlass-4.5.0 = 235/235), gate 16/16 both, 27B 235/235 + 35B 315/315 unchanged. RoPE-ON closes total tput 0.90×→0.97× (c1) / 0.62×→0.82× (c8) and c1 TTFT ratio 5.85×→2.27×, still `ACTIVE`. **SPEED RE-BOUND 2026-07-21 (same-session, matching-recipe) — TTFT residual RESOLVED, cutlass claim CORRECTED:** the "TTFT 2.27×/5.85×" + "c8 ITL 4.3×" were BAD-DENOMINATOR/num-prompts artifacts — a fresh same-session vLLM capture gives c1 TTFT ~152 ms & c8 ITL P99 ~130 ms, and OURS WINS TTFT at both concurrencies (c1 0.90×, c8 0.38×). **c1 = effective every-axis parity** (tput 0.98× / TPOT 1.01× / TTFT+ITL wins); **c8 residual = decode** (tput 0.93× / TPOT 1.10× / ITL P99 1.12×), which nsys shows is 93% GPU-busy/compute-bound (small-M=8 `cutlass_80_wmma` projections). The **qkv-merge** (new GQA `QkvSplit` op mirroring vLLM `QKVParallelLinear`) was implemented + MEASURED NEUTRAL (doesn't cut decode FLOPs) ⇒ default-OFF. **CUTLASS CLAIM CORRECTED: 27B `test_qwen27_paged_engine` = 235/235 on the FLASHINFER cutlass build** (the "flashinfer ⇒ 234/235" was a build artifact). Stays `ACTIVE`; named residual = c8 decode-GEMM efficiency (a decode-fusion sub-campaign). **SWEEP MODEL #1 — Qwen3-Coder-30B-A3B (`Qwen3MoeForCausalLM`) W0+W1 LANDED 2026-07-21** ([sweep-qwen3-coder-30b.md](specs/sweep-qwen3-coder-30b.md), `MODEL-TEXT-qwen3-moe-qwen3-moe-for-causal-lm` `ACTIVE`, `CLAIM-MODEL-QWEN3-CODER`): the first full-attention BF16 MoE, composed from the done dense attention + the done 35B MoE experts (ZERO runner change). W0 = registry stub (`qwen3_moe_registry.cpp` + `qwen3_moe.h`, one `REGISTER_VLLM_MODEL`, full-attn-only KV, `is_dense_model=false`, W2/W3 throwing stubs). W1 = three behaviour-preserving refactors making the two done pieces reusable: (#1) dense `AttnBlock` + glue extracted to `dense_attn_block.h` (Qwen3-dense byte-identical), (#2) bf16 `MoeBlock` exposed cross-TU via `RunMoeBlock` (`qwen3_5_moe_block.h`; 35B untouched), (#3) no-shared-expert guard (inert for the 35B). Gated: dgx CUDA `-Werror` 0-warn; Qwen3-dense 0.6B+4B 16/16 + 27B 235/235 + 35B 315/315 UNCHANGED; registry resolves `Qwen3MoeForCausalLM`; memcheck 0. W2 bf16 loader → W3 forward → W4 near-tie token-exact → W5 fast bf16 grouped-MoE GEMM remain. Then Llama dense (download), Mistral, MoE families **SWEEP MODEL — GLM-4-9B-0414 (`Glm4ForCausalLM`) G2 LANDED 2026-07-24 — CORRECTNESS COMPLETE** ([glm-dsa-latest-deepseek](specs/glm-dsa-latest-deepseek.md), `MODEL-TEXT-glm4-glm4-for-causal-lm` now `READY` per the [live-state audit](specs/live-state-audit-2026-08-06.md), `CLAIM-GLM-DSA-LATEST-DEEPSEEK` amended) **GLM-4.7-Flash (`Glm4MoeLiteForCausalLM`, 31.2B MLA+MoE) G1 LANDED 2026-07-24 — SACRED gate 8/8, CORRECTNESS COMPLETE** (`MODEL-TEXT-glm4-moe-lite-glm4-moe-lite-for-causal-lm` `ACTIVE`, speed pending): reuses the DeepSeek-V2 MLA stack; first e2e coverage of the q_lora branch + noaux_tc sigmoid router, closing the MLA campaign's two C2 coverage gaps: the FIRST GLM-family model. SACRED gate 16/16 vs vLLM 0.25.0 (STRICT 13/16 + near-tie 3/16, max gap 0 nats; vLLM K=5 self-deterministic ⇒ STRICT bar), speed PENDING. The two "new primitives" the spike named reduced to EXISTING infra: partial + interleaved `RopeFromCache` (`is_neox_style=false`, the DeepSeek-V2 decoupled-rope path) over `rotary_dim=64`; standalone `vt::RmsNorm` sandwich norms. Biased qkv, no QK-norm, GQA 32/2, untied lm_head. New files + one REGISTER, reusing the shared dense glue. **SWEEP MODEL — Llama-3.2 (`LlamaForCausalLM`) W0-W4 LANDED 2026-07-23 — CORRECTNESS COMPLETE** ([sweep-llama-3.2](specs/sweep-llama-3.2.md), `MODEL-TEXT-llama-llama-for-causal-lm` `ACTIVE`, `CLAIM-MODEL-LLAMA-3.2`): the roadmap's explicit "Llama-first" increment and the first mainstream non-Qwen/non-OPT dense arch. `LlamaForCausalLM` (Llama-3.2-1B) = the Qwen3-dense forward with exactly two ADDITIVE deltas — NO qk-norm (shared `AttnBlock` skips it when q_norm/k_norm empty) + llama3 rope-scaling (4 default-0 `RopeArgs` fields + a `Llama3ScaleFreq` kernel helper, no-op elsewhere; formula verified 2e-7 rel vs vLLM) — reusing the shared dense forward VERBATIM (`LlamaModel == Qwen3DenseModel`). 3 new Llama files, ZERO edit to runner/scheduler/platforms/attn-registry/`hf_config`/any qwen3-opt model. vLLM 0.25.0 greedy MEASURED DETERMINISTIC (K=6, 0 multi-valued cells) ⇒ STRICT token-exact bar, PASS **16/16 (12 strict + 4 near-tie band, max gap 0.0000 nats, 0 divergent)** — at all 13 divergent positions vLLM's own teacher-forced argmax given our prefix IS our token. A correctness-fatal tokenizer bug (Llama's `Sequence` post_processor wrapping `TemplateProcessing` ⇒ BOS 128000 never prepended, silently 1/16) was isolated by a CUDA prefill-argmax diagnostic (forward proven 4/4 correct given vLLM's exact tokens) and fixed byte-preservingly (Qwen/OPT/DeepSeek unaffected — all ByteLevel/top-level-TemplateProcessing). Regressions 27B 235/235 · 35B 315/315 · Coder 6/6 · Qwen3-dense 16/16 · OPT 6/6 · DeepSeek-V2 8/8 UNCHANGED; `-Werror` 0-warn; memcheck 0; DSR 67. SPEED pending (head_dim 64 → generic paged path; Llama-3.2-3B head_dim-128 is the FA2-toggle W-next). **MLA CAMPAIGN W0+W1 LANDED 2026-07-21** (`CLAIM-MLA-DEEPSEEK`; rows STAY `SPIKE` — W0/W1 make no model supported). **W0 grounded every fact the spike flagged as an unverified source read; ALL CONFIRMED, none contradicted:** DeepSeek-V2-Lite fetched to dgx (30 GB, 4 shards) and loading in the vLLM 0.25.0 oracle; the oracle's own DEBUG startup on sm_121 prints `Using TRITON_MLA attention backend out of potential backends: ['TRITON_MLA']` and `Using FLASH_ATTN MLA prefill backend` — so the dense-MLA decode + MLA-prefill targets are OBSERVED, not inferred, and the sm90/sm100-only MLA kernel class stays out of scope; the real `config.json` confirms every §5.1 number (`kv_lora_rank=512`, `qk_nope=128`, `qk_rope=64` -> the **576-wide latent**, `v_head_dim=128`, `q_lora_rank=null`, `n_group=topk_group=1`, softmax/greedy, 64+2 experts, 27 layers) plus `is_neox_style=False` and the mscale2 scale correction; and BOTH recorded coverage gaps (no `fused_qkv_a_proj` branch, no `e_score_correction_bias`) are confirmed real, so those pieces stay unit-gated only. **W1 = the behaviour-preserving spec-driven KV allocation, ZERO MLA math:** the attention cache is now sized `num_blocks * spec->page_size_bytes()` and viewed from the spec's own `block_size`/`num_kv_heads`/`head_size`/`dtype` instead of the hardcoded `2 * block * Hkv * Dh` with shape reconstructed from the HF config (`runner.cpp`), plus `MLAAttentionSpec` with upstream's factor-1 single-tensor page formula (`kv_cache_interface.py:397-398`) registered against the ORDINARY `FullAttentionManager` (`single_type_kv_cache_manager.py:1539`) — the spike's key finding, which is why block manager/prefix caching/eviction need no change. Gated: dgx clean CUDA `-Werror` 0 warnings/0 errors; **27B 235/235 + 35B 315/315 + Qwen3-Coder 6/6 + Qwen3-dense 16/16 ALL UNCHANGED** (behaviour-preserving proven, not assumed); `test_runner` 15/15, `test_kv_cache_interface` 21/21 (4 new MLA-spec cases), `test_llm_engine` 5/5; the new path is proven EXERCISED (not merely compiled) by `fa_page_size_bytes()` + a `page_size_padded` case no HF-config formula can produce. **W2 + W3 LANDED 2026-07-21** (base `a05437f`; rows STAY `SPIKE` — still no MLA attention math, no MLA model, no forward). **W2 = the MLA branch of `_get_backend_priorities` the pre-W2 comment deferred, ported as DATA:** the whole of `cuda.py:84-176` (BOTH branches — MLA sm_100 including the `:96-115` adaptive sparse tail, MLA sm_12x, MLA `else`, and the two pre-existing non-MLA arms) is now a TABLE in the new header `include/vllm/platforms/cuda_attn_priority.h`, one row per upstream arch arm keyed on `(use_mla, major)`, so a future arch is a ROW rather than a code path; putting it in a header (not the CUDA-only TU) also let the CPU test tier assert the REAL table and DELETED the hand-copied `FakeCudaPlatform` duplicate. On sm_121 a `use_mla=true` request now RESOLVES to `TRITON_MLA`, matching the W0 oracle observation. **The sparse/DSA seam is left OPEN and unit-proven:** GB10's row keeps both upstream entries and the sparse one loses to a real FILTER — `AttentionBackend::is_mla()`/`is_sparse()` checked against the request (`backend.py:307-360 validate_configuration`) — so a future DSA backend is selected purely by declaring `is_sparse() == true`, with ZERO edit to the table or the selector. `TritonMLABackend` lands the NAME plus upstream's 3-D `get_kv_cache_shape` (no K/V axis; `num_kv_heads != 1` REFUSED), `get_impl_cls()` deliberately still `nullptr`. MLA prefill priority ported too (GB10 -> `[FLASH_ATTN]` alone). **W3 = the two new `vt::` ops, both CPU-reference-gated.** `vt::ConcatAndCacheMla` mirrors `csrc/libtorch_stable/cache_kernels.cu:401-442` — and per the whole-chain rule this was VERIFIED, not assumed, to be vLLM's OWN csrc kernel (`_custom_ops.py:2532` -> `torch.ops._C_cache_ops`), with no flashinfer/cutlass variant in the dense-bf16 path; it concatenates the latent + rope part into ONE 576-wide entry, the write `ReshapeAndCache`'s K/V-pair signature cannot express, stride-driven so a per-layer cache slice and the two column halves of `kv_a_proj_with_mqa` both work copy-free. The **grouped-topk (`noaux_tc`) router** extension — flagged in `coordination.md` as SHARED with `CLAIM-GLM-DSA-LATEST-DEEPSEEK` and "must not be implemented twice" — is landed HERE and that claim now consumes it: additive `MoeRouterTopKArgs` fields + an optional `e_score_correction_bias`, with `num_expert_group == 0` still dispatching the ORIGINAL kernel so the 27B/35B/Coder/dense routers are byte-identical BY CONSTRUCTION. **Stated plainly: the `noaux_tc` correctness evidence is UNIT-ONLY.** V2-Lite has `n_group=topk_group=1` and no bias, so the e2e vehicle exercises none of it; the gate is `tests/vt/test_ops_moe_router_grouped.cpp` at DeepSeek-V3's REAL dimensions (256 experts, n_group=8, topk_group=4, sigmoid, scaling 2.5, WITH the bias) against an INDEPENDENT sort-based transcription of the upstream formula. **W4 LANDED 2026-07-22** (base `ed2c342`; rows STAY `SPIKE` — W4 adds a kernel and fills a `nullptr`, it makes no model supported). **`vt::MlaDecodeAttention` — the MQA decode over the compressed latent (QK 576 / V 512, `num_kv_heads=1`), a structure port of the two-stage split-KV pair W0 OBSERVED EXECUTING:** `MlaDecodeStage1` <- `_fwd_grouped_kernel_stage1` (`triton_decode_attention.py:278-458`, the `IS_MLA` branch whose `v = tl.trans(k)` at `:424-431` is the whole MLA trick — V is the leading 512 columns of the SAME latent row already loaded as K, so one shared-memory tile serves as both), `MlaDecodeStage2` <- `_fwd_kernel_stage2` (`:575-639`), `ComputeNumKvSplits` <- `_compute_num_kv_splits` (`triton_mla.py:40-47`), split workspace <- `_reserve_attn_logits_workspace` (`:57-78`) realized as the house grow-only per-stream scratch. **Honest reuse verdict:** our FA-2 split+combine machinery fit at the ALGORITHM level (the split schedule, the LSE merge algebra, the fixed-ascending no-atomicAdd determinism rule) and NOT at the code level — the vendored FA-2 launcher takes separate 4-D k/v caches and is instantiated for symmetric head_dim {128,256}, which cannot express a 3-D single-buffer cache with QK 576 / V 512; that is recorded in the TU header rather than forced. **Evidence is unit-level and deliberately strong** (there is no e2e model until W7): [`tests/vt/test_ops_mla_attn.cpp`](../tests/vt/test_ops_mla_attn.cpp), a port of `tests/kernels/attention/test_mla_decode_cpu.py` whose `ref_mla` becomes an INDEPENDENT TWO-PASS oracle (a different algorithm from the streaming online-softmax both impls use) plus its NaN-padding out-of-bounds detector, run at the REAL V2-Lite geometry (576/512/64, block 16, mscale^2 scale) over ragged / multi-block / single-block / every split boundary (`num_kv_splits` ∈ {1..512} incl. splits > seq_len) / 128-head V3 / non-BLOCK_H head counts / a 288-256 block-32 geometry / bf16 + f32, with run-to-run BIT-exactness. dgx sm_121: 11/11 cases, 2,303,193 assertions; `compute-sanitizer` memcheck **0 errors**, racecheck **0 hazards**, synccheck **0 errors**; clean CUDA build **0 warnings / 0 errors**; **27B 235/235 + 35B 315/315 + Coder 6/6 + Qwen3-dense 16/16 + OPT 6/6 ALL UNCHANGED**. `TritonMLABackend::get_impl_cls()` is no longer `nullptr` — it returns a real `TritonMLAImpl` whose `forward_mqa` is the 1:1 counterpart of `triton_mla.py:189-260`; PREFILL is W5 and `forward()` refuses a prefill-shaped batch BY NAME rather than producing wrong numbers. NO speed number — decode perf is W9. **W5 LANDED 2026-07-22** (base `5395203`; rows STAY `SPIKE`). **MLA PREFILL + the workspace-bounded CHUNKED-CONTEXT loop.** Three new ops — `vt::MlaPrefillAttention` (<- `mla/prefill/flash_attn.py:153-248`, the ONLY MLA prefill backend reachable on sm_121 and the one W0 OBSERVED the oracle logging), `vt::GatherMlaCache` (<- `cache_kernels.cu:992-1064`) and `vt::MergeAttnStates` (<- `merge_attn_states.cu:18-192`, both `-inf` edge cases verbatim) — plus the loop itself in the new `mla_chunked_context.h` (<- `mla_attention.py:1422-1451,1667-1745,2094-2199,2344-2425`), which is what keeps a long-context prefill inside a bounded workspace instead of materializing a 3 GB up-projected context. **The vendored FA-2 launcher WAS generalized, and W4's prediction that it would be tractable held for a reason worth recording: upstream does not ask FA-2 for asymmetric head dims either.** `requires_v_padding` is TRUE on GB10, so upstream ZERO-PADS V from 128 to 192 and slices the output back — the kernel stays a plain SYMMETRIC head_dim-192 instantiation. The whole change is two new explicit instantiations of an UNCHANGED generic template, one new launcher entry for the contiguous-varlen mode, and the pad/slice pair; the paged launcher every non-MLA prefill calls is textually untouched (211 insertions / **0 deletions**), and 27B 235/235 + 35B 315/315 + Coder 138/138 + Qwen3-dense 664/664 + OPT 36/36 are all UNCHANGED. Evidence is UNIT-ONLY (there is still no model): 4/4 cases / **2,377,052 assertions** and 5/5 / **306,037 assertions** on dgx sm_121 at the real QK 192 / V 128 geometry, against an INDEPENDENT double-precision two-pass oracle and — for the loop — a SINGLE-SHOT whole-sequence oracle that never chunks, over exact / +1 / -1 chunk boundaries, zero-context and zero-key-in-chunk requests, ADVERSARIAL reverse-interleaved block tables, NaN-poisoned outputs and run-to-run bit-exactness; memcheck **0**, racecheck **0 hazards**, synccheck **0**. A genuine upstream FA-2 quirk was found on the way and worked around rather than papered over: its EMPTY-K early exit ignores the unpadded-LSE flag, which a zero-key chunk request would turn into an out-of-bounds LSE write. **W6 LANDED 2026-07-22** (base `2846467`; rows STAY `SPIKE` — W6 adds an attention LAYER, not a model). **The MLA attention BLOCK + LOAD-TIME WEIGHT ABSORPTION — the piece that finally COMPOSES W3's cache write, W4's MQA decode and W5's MHA prefill into one layer:** the projections with BOTH `q_lora_rank` branches (`fused_qkv_a_proj` -> `q_a_layernorm` -> `q_b_proj`, or the direct `q_proj`), the two RMSNorms (the rope part deliberately NOT normed), the DECOUPLED RoPE (`is_neox_style=False`, only the trailing 64-dim slice rotates) with its YaRN cos/sin cache and the SEPARATE mscale^2 softmax-scale correction, the `kv_b_proj -> W_UK/W_UV` split, the prefill-MHA / decode-MQA dispatch with decode tokens packed FIRST, and the `kv_b_proj` up-projection callback W5 left open. **The spike's most useful prediction held: absorption needed NO new attention kernel** — it is a LOAD-TIME weight transform plus TWO batched GEMMs, so the entire new-kernel surface is two general primitives, `vt::BatchedMatmul` (<- `torch.bmm` at `mla_attention.py:789,1034`; on CUDA torch resolves that to cuBLAS `gemmStridedBatchedEx`, and ours is the cuBLASLt strided-batched form of the same GEMM) and `vt::ConcatMlaNopeRope` (<- `concat_mla_q`, generalized so one op also serves `_concat_k_nope_k_pe`). **The absorbed-vs-unabsorbed equivalence — the heart of W6 — is PROVEN NUMERICALLY, three independent ways, rather than argued:** an INDEPENDENT double-precision block oracle computes the attention BOTH ways and agrees to < 1e-11 (the identity itself); our absorbed decode reproduces the UNABSORBED oracle to < 2e-4 in f32; and — the strongest — the SAME batch is driven once through our ABSORBED MQA decode kernel (QK 576 / V 512, one KV head, K/V never materialized) and once through our UNABSORBED materialized-MHA prefill path (QK 192 / V 128 plus the chunked-context loop), agreeing to < 3e-4 (CPU f32) / < 4e-2 (CUDA bf16) with nothing but the weights shared between them. Evidence on dgx sm_121: `test_mla_attention_block.cpp` 10/10 cases / 2,372,644 assertions and `test_ops_mla_absorb.cpp` 9/9 / 1,644,807 (CUDA cases proven to EXECUTE; NaN-poisoned outputs; run-to-run BIT-exact), porting `tests/kernels/test_concat_mla_q.py` in both arms. memcheck 0, racecheck 0 hazards, synccheck 0 (the last needing `--num-cuda-barriers 65536` — the default table overflows on a binary driving this many kernel families and the tool then reports a bogus launch failure, a trap worth knowing). Clean CUDA build 0 warn / 0 err; **27B 235/235 + 35B 315/315 + Coder 138/138 + Qwen3-dense 664/664 + OPT 36/36 ALL UNCHANGED**. **Coverage stated plainly: the `q_lora` query branch has NO e2e coverage and cannot get any on GB10** — DeepSeek-V2-Lite has `q_lora_rank=null`, so it is unit-gated at DeepSeek-V3's real dimensions only; GLM-4.7-Flash (`q_lora_rank=768`, 58.2 GiB, fits) is what would close it. **W7 LANDED 2026-07-22** (base `ce43c51`; the row STILL stays `SPIKE`). **THE DEEPSEEK-V2 MODEL — registry + config parse + loader + forward: the first MLA model in the tree, and the first one that runs a real MLA checkpoint end to end.** Four new files plus ONE shared-code edit (a two-line additive condition in `runner.cpp` recognising a `kMlaAttention` KV group as the model's attention group — upstream maps MLA onto the ordinary `FullAttentionManager`, so block tables/prefix caching/eviction are untouched). **LOADER GATE PASSED on the real 4-shard DeepSeek-V2-Lite: 5291/5291 checkpoint tensors accounted for, ZERO unmapped and ZERO leftover** (4/4 cases / 37,331 assertions), every shape asserted including the LOAD-TIME `kv_b_proj -> W_UK_T [16,128,512]` / `W_UV [16,512,128]` absorption split — the same transform, at the same lifecycle point, as upstream's `process_weights_after_loading`. **V2-Lite takes the DIRECT `q_proj` query branch** (`q_lora_rank: null`), asserted with the fused branch EMPTY on every layer. **FORWARD GATE PASSED and obviously right, not merely finite: the real checkpoint prefill of `The capital of France is` -> argmax ` Paris`** (top-5 ` Paris`/` the`/` a`/` one`/` also`, run-to-run bit-exact) — the direct analogue of the Qwen3-Coder W3 sanity case. **BATCH-ORDERING GATE:** the ordering invariant W6 measured 0.86 relative error from is now VALIDATED, not assumed — `BuildMlaBatchSplit` throws (naming the request and citing the upstream line) if a decode follows a prefill or a with-context prefill follows a context-free one. **SHARED EXPERTS — new for this family and UNGATED unlike Qwen3.6's sigmoid-gated one — gated two ways:** a MoE layer with every routed expert zeroed is BIT-IDENTICAL to a dense layer holding the same MLP, and turning the shared expert off CHANGES the logits. **The CUDA path is EXERCISED, not merely compiled:** a case at the real MLA head geometry drives the CUDA MLA kernels and the CUDA-only grouped bf16 MoE GEMM, bit-exact on device and within 0.0061 worst relative logit error of the CPU reference path. 11/11 forward cases; memcheck/racecheck/synccheck all **0**; clean CUDA build **0 warn / 0 err**; **regression set UNCHANGED**. **Only `DeepseekV2ForCausalLM` is REGISTERED** — `DeepseekForCausalLM` (plain MHA), V3 (fp8/671B) and V3.2 (DSA indexer) are REFUSED BY NAME in the config parse rather than falsely claimed. A pre-existing tree-wide hazard was found on the way and recorded: the shared `DevicePool` is a process-wide singleton keyed only on a byte size class, so a single process driving BOTH a CPU and a CUDA forward hands the second backend the first's recycled pointers. **NEXT: W8 — the SACRED token-exact gate on DeepSeek-V2-Lite** (wire the paged engine to produce the MLA batch order the model already validates, capture oracle goldens, run the STRICT form W0 determined). A loading, forwarding model is NOT a supported model, so no model row moves until that gate passes. **W8 LANDED 2026-07-22 — THE SACRED CORRECTNESS GATE PASSES 8/8, and `MODEL-TEXT-deepseek-v2-deepseek-v2-for-causal-lm` moves `SPIKE` -> `ACTIVE` (correctness COMPLETE, speed PENDING). NOT `DONE` — that additionally requires vLLM-speed parity on every axis, which is W9 and has NO number yet; the other four campaign rows stay `SPIKE`.** An 8-prompt battery is driven through the FULL paged `LLMEngine` and compared to the pinned vLLM 0.25.0 oracle: **8/8 PASS — STRICT token-exact 5/8, near-tie band 3/8, 92/128 tokens strictly exact, max teacher-forced gap 0.25 nats, 0 forward-divergent** (223 assertions). **The bar was ARRIVED AT by measurement, not chosen:** vLLM is DETERMINISTIC on this model at batch=1 (W0's K=5 8/8, re-confirmed by W8's own capture at T=16 with 0 multi-valued cells), so the STRICT form ran FIRST and came out 5/8; the ratified TEACHER-FORCING diagnostic then showed **36 divergent positions with 35 at gap EXACTLY 0.0000 nats** — vLLM's own argmax GIVEN OUR PREFIX is our token, so they are the downstream tail of one earlier flip — **exactly ONE root flip with any gap at all (prompt[3] tok 9, 0.2500 nats, inside the ratified 0.5-nat band and equal to the landed Qwen3-dense 4B gate's worst)**, and **ZERO tokens outside vLLM's top-20**, with the per-position nats COMMITTED as goldens and anything beyond the band still FAILING. **W8's first job — the scheduler/runner wiring — turned out to need NO new code, for a non-accidental reason:** `runner.cpp:671` already reorders with `decode_threshold = 1`, exactly MLA's `reorder_batch_threshold` (`mla_attention.py:1420`), and its `decode -> short_extend -> long_extend -> pure_prefill` ordering satisfies BOTH MLA invariants (decodes form a batch prefix; with-context prefills lead the prefill tail). W8 PROVES that end to end rather than duplicating it, with new DIAGNOSTIC `MlaBatchSplitStats` counters and a non-vacuity bar: the battery is admitted CONCURRENTLY with staggered arrival, producing **7 genuinely MIXED decode+prefill steps at up to 8 concurrent requests** with `BuildMlaBatchSplit` (which throws naming the request) never firing, plus a prefix-cache-driven **with-context prefill**, and a phase-0 check that the engine really allocated the compact MLA cache (`fa_page_size_bytes = 36864`, no factor 2). **THE REAL BLOCKER WAS THE TOKENIZER, NOT THE MODEL:** the first run REFUSED to load (`unsupported normalizer "Sequence"`), and behind it sat a whole NEW pre-tokenizer family — DeepSeek's is a HF `Sequence` PIPELINE of SEVEN stages (five `Split(Isolated)` over ENUMERATED codepoint ranges, then `Digits(individual_digits=true)`, then `ByteLevel(use_regex=false)`), whose stage ORDER is load-bearing because stage 2's punctuation class spans 0x3A-0x7E and CONTAINS A-Z/a-z. Landed as `SplitPattern::kDeepSeek` with the five patterns compared VERBATIM at load, and MEASURED token-for-token against the REAL HF `tokenizers` library over a stage-stress corpus (**6/6 cases / 2461 assertions**). **The TOKENIZATION goldens earned their keep by REFUTING a fix that was already written:** `tokenizer_config.json` declares `add_bos_token: true`, which reads as exactly the OPT missing-BOS bug — but vLLM's resolved tokenizer (`TokenizersBackend`) adds NO BOS, our loader already matched bit-for-bit, and the "fix" would have BROKEN a passing gate; it was reverted and the measured behaviour PINNED by a guard case ([[ground-premises-before-dispatching]]). Regression set UNCHANGED (27B 235/235, 35B 315/315, Coder 6/6, Qwen3-dense 16/16, OPT 6/6, plus every tokenizer test — W8 touches SHARED tokenizer code, so that was proved, not assumed); clean CUDA rebuild 0 warn/0 err; local CPU suite 151/151; memcheck/racecheck/synccheck 0. Batch invariance is REPORTED (6/8) and deliberately NOT a bar, because the ORACLE itself changed on 3/8 of this battery under batched generation (W0). One W9 input recorded: the oracle must run `moe_backend='triton'` — vLLM's auto-selected FlashInfer CUTLASS unquantized MoE REBOOTED dgx three times on GB10's unified memory. **W9 SPEED CLOSE LANDED 2026-07-22 — the track has its FIRST binding speed number, and it is an ATTRIBUTED MISS: `MODEL-TEXT-deepseek-v2-deepseek-v2-for-causal-lm` STAYS `ACTIVE` (correctness COMPLETE, speed SHORT), NOT `DONE`** ([grid](../docs/BENCHMARKS.md), [spike §W9](specs/mla-deepseek-campaign.md)). Denominator SETTLED with evidence — CUTLASS MoE has now rebooted dgx **five times** (two more at W9, the second on a pristine box with a 0 GiB page cache and every mitigation applied, both deaths at the identical post-`torch.compile` phase), so `--moe-backend triton` IS vLLM's best STABLE GRAPHED configuration here and is the bar; the substitution does not flatter us, we lose to it. `nsys` (both sides, `--cuda-graph-trace=node`) overrode the plan: the lever was not the planned MLA fusion recipes but `MlaDecodeStage1` sitting at **44.7% of all GPU time and ~180x off its own memory-bound floor** on a **2-CTA grid at batch 1**; applying upstream's own occupancy target made it **18.3x faster** (837 -> 45.8 us) for **+69.5%/+53.3%/+32.0%/+19.5%** end-to-end at c1/c2/c4/c8, while the planned decode-graph sibling is worth only ~+2% (this decode is GPU-bound). Grid vs vLLM: output throughput **0.87/0.95/0.86/0.88** (was 0.50 at c1), TTFT **1.06/1.14/0.96/0.88** (we WIN at c4/c8), TPOT **1.11/0.97/1.16/1.17**. SACRED gate **8/8 UNCHANGED** with both levers default-ON; a real latent CUDA-graph use-after-free in the MLA metadata upload was found and fixed (its whole class now guarded); regression set UNCHANGED; clean rebuild 0 warn/0 err; sanitizers 0. **NEXT LEVER, NAMED: route the batch-1 dense projections off cuBLAS `gemvx` (31.8% of our GPU time) onto a tensor-core GEMM — vLLM splits the same work `gemvx` 12.7% + `nvjet_sm121_tst_mma_*` 6.6%.** **W10 BLOCKED-ROW HONESTY PASS LANDED 2026-07-22 — the campaign's W-plan is COMPLETE; records only (no code, no build, no GPU work, nothing downloaded, no number claimed).** Rows set to their final honest state: `MODEL-TEXT-deepseek-v2-deepseek-v3-for-causal-lm` (V3 + V3.2, and Kimi-K2/K2.5's text backbone by config composition) and `MODEL-TEXT-minimax-m2-mini-max-m2-for-causal-lm` move `SPIKE` -> `BLOCKED`, joined cross-claim by `MODEL-TEXT-deepseek-v2-glm-moe-dsa-for-causal-lm` (GLM-5) under `CLAIM-GLM-DSA-LATEST-DEEPSEEK`; each is HW-BLOCKED on 119 GiB (~642 GiB fp8 / ~428 GiB / 1404 GiB) and the two DSA models are additionally DEP-BLOCKED — for a SPARSE model the XOR filter eliminates `TRITON_MLA`, leaving `FLASHINFER_MLA_SPARSE_SM120` alone, whose sm12x dispatch goes to flashinfer's DENSE-ONLY XQA backend that discards `sparse_mla_top_k` (upstream's own test monkeypatches the probe and asserts nothing numerical). `MODEL-TEXT-deepseek-v2-deepseek-for-causal-lm` stays `SPIKE` with the record repaired to say it is plain MHA and needs NO MLA; Kimi-Linear stays `SPIKE` (MLA half unlocked, KDA a separate kernel campaign, HW-MARGINAL). Each blocked row states what CAN still be gated (config resolution, weight-map on a slice, unit parity at real dimensions) versus what CANNOT (anything e2e). **Two PERMANENT coverage gaps now stated in the rows:** the `noaux_tc` grouped router and the `q_lora` query branch have NO e2e coverage and are unit-gated only, because V2-Lite is `n_group=topk_group=1`/softmax with no `e_score_correction_bias` and `q_lora_rank=null`. **NAMED NEXT VEHICLE: GLM-4.7-Flash** (`MODEL-TEXT-glm4-moe-lite-glm4-moe-lite-for-causal-lm`, 31.2B / 58.2 GiB, FITS GB10) — the only reachable checkpoint that closes BOTH gaps. **BLOCK NOT CLOSEABLE, nothing archived:** the DeepSeek-V2 row is `ACTIVE`, not `DONE`, so the plan/spec stay LIVE; the single open item is the named `gemvx` -> tensor-core dispatch lever. **MISTRAL FIFTH FAMILY W0-W3 LANDED 2026-07-23** ([sweep-mistral](specs/sweep-mistral.md), `MODEL-TEXT-mistral-mistral-for-causal-lm` `ACTIVE`, `CLAIM-MODEL-MISTRAL`): the closest-to-Llama dense arch (vLLM `mistral.py` = "Mistral adaptation of the LLaMA architecture") — plain rope θ1e6 (no rope_scaling) + qk-norm-optional + untied lm_head + null sliding_window, all PRE-EXISTING ⇒ NO new primitive, 3 new files + additive CMake/registry-test rows only, ZERO shared-code edit. **MODEL forward gate 30/30 greedy tokens vs vLLM 0.25.0** (tokenizer-free: fed vLLM's exact prompt tokens through our CUDA prefill; 29 STRICT token-exact + 1 near-tie, 0 forward-divergent; vLLM greedy det 4/5 K=3). W2 loader real-weights 1541 assertions. **REAL FINDING:** Mistral's SentencePiece/Metaspace tokenizer is unsupported by our ByteLevel-BPE tokenizer → the FULL paged-engine SACRED gate is BLOCKED, the pre-inventoried `LOAD-SENTENCEPIECE` row (SentencePiece tokenizer family). `-Werror` 0-warn, DSR 32, regressions UNCHANGED (Llama paged 16/16, Qwen3-dense forward 1031, registry 299; MoE/GDN gates unaffected by construction). SPEED + full paged gate both PENDING (row `ACTIVE`, not `DONE`). **OLMo-2 SPIKED 2026-07-24** ([sweep-olmo2](specs/sweep-olmo2.md), `CLAIM-SWEEP-OLMO2`): one row `INVENTORIED` → `SPIKE` (`MODEL-TEXT-olmo2-olmo2-for-causal-lm`, covering `Olmo2ForCausalLM` + its `Olmo3ForCausalLM` alias). Answers the breadth-sweep §B.3 Tier-2 rank-8 "GLM4 / Olmo2-3" item (GLM-4 + Gemma landed; OLMo-2 next). **HEADLINE: OLMo-2 is the cleanest dense bring-up yet — ZERO new compute kernels.** The two distinctive facts both reduce to WIRING over landed ops: (1) the **pure post-norm (`norm_after`) placement** is a strict SUBSET of the GLM-4/Gemma sandwich (keeps ONLY the standalone-output-norm op `glm4.cpp:174-178` — the exact primitive flagged — DROPS the pre-norms, plain residual add); (2) the **QK-norm is FULL-WIDTH not per-head** → reuses `vt::RmsNorm` at a `[T,q_size]`/`[T,kv_size]` shape but CANNOT use the fused per-head `kAttnQkNormRopeGate`. Everything else REUSES (plain RMSNorm, SiLU SwiGLU, NeoX rope, GQA paged glue, Gemma-3 sliding-window for Olmo-3, tied embeddings, packed loader, ByteLevel BPE). Gate vehicle `allenai/OLMo-2-0425-1B` (1.485B, ~2.77 GiB, fits GB10 tight ~30 GiB free); Olmo-3 rides the same row (0.25.0 oracle-support UNVERIFIED). OLMo-1 (non-parametric LayerNorm), OLMoE/FlexOlmo (MoE), OlmoHybrid (SSM) stay `INVENTORIED`. Nothing implemented. | -| 2a | `ROAD-V1-C2-LOCAL-BF16` | Close local discrete-Blackwell Qwen3.5 plain-BF16 production parity | [`MODEL-MM-qwen3-5-qwen3-5-for-conditional-generation`](model-matrix.md), [`LOAD-SAFETENSORS-DIRECT-DENSE`](engine-matrix.md), [`SERVE-CLI-BENCH`](engine-matrix.md), [`KERNEL-SSM-MAMBA`](kernel-matrix.md) | Exact `(sequence, 8-token chunk)` GDN conv dispatch is byte-identical and default ON. Rebased-main same-binary reprofile: **718.704→233.955 ms (3.072x)** and whole run **+2.272%**; sealed vLLM remains 145.421 ms (**1.609x residual**). The reviewed combined default-off K4 causal-conv + 16-token post-conv local A/B is exact and improves total throughput **0.7175%**, TTFT **1.340%**, TPOT **0.556%** and E2E **0.716%**. The #206 pretoken frontend is CPU-GREEN, but its first real A/B failed output identity (98/128 requests; 15,507/16,384 positions) because one-at-a-time AsyncLLM publication changed the execution wave; all apparent timings are VOID. Atomic wave admission matching vLLM's add-all-before-step behavior is now specified. Cross-engine throughput and TTFT remain **VOID/PENDING**; TPOT and VRAM remain OPEN | `GATING` | Fresh implementer: RED-first atomic batch queue/AsyncLLM wave admission and benchmark refill-wave dispatch; fresh mutation review; operator CPU gate; then restart counterbalanced `STRING→TOKENS→TOKENS→STRING` GPU A/B and pinned-vLLM rerun. Latency, VRAM and 27B/35B correctness remain gates; no 4B-to-gate-model extrapolation. Evidence: [campaign spec](specs/sm120-qwen35-pareto-2026-08-09.md), [production baseline](../docs/bench-evidence/qwen35-4b-sm120-main-20260807.md) | +| 2a | `ROAD-V1-C2-LOCAL-BF16` | Close local discrete-Blackwell Qwen3.5 plain-BF16 production parity | [`MODEL-MM-qwen3-5-qwen3-5-for-conditional-generation`](model-matrix.md), [`LOAD-SAFETENSORS-DIRECT-DENSE`](engine-matrix.md), [`SERVE-CLI-BENCH`](engine-matrix.md), [`KERNEL-SSM-MAMBA`](kernel-matrix.md) | Exact `(sequence, 8-token chunk)` GDN conv dispatch is byte-identical and default ON. Rebased-main same-binary reprofile: **718.704→233.955 ms (3.072x)** and whole run **+2.272%**; sealed vLLM remains 145.421 ms (**1.609x residual**). The reviewed combined default-off K4 causal-conv + 16-token post-conv local A/B is exact and improves total throughput **0.7175%**, TTFT **1.340%**, TPOT **0.556%** and E2E **0.716%**. The #206 pretoken frontend is CPU-GREEN, but its first real A/B failed output identity (98/128 requests; 15,507/16,384 positions) because one-at-a-time AsyncLLM publication changed the execution wave; all apparent timings are VOID. Atomic all-or-zero AsyncLLM wave admission matching vLLM's add-all-before-step behavior is now implemented and CPU-GREEN (12/12·433 async gate); fresh mutation review and real retry remain pending. Cross-engine throughput and TTFT remain **VOID/PENDING**; TPOT and VRAM remain OPEN | `GATING` | Fresh mutation review; operator CPU gate; then restart counterbalanced `STRING→TOKENS→TOKENS→STRING` GPU A/B and pinned-vLLM rerun. Latency, VRAM and 27B/35B correctness remain gates; no 4B-to-gate-model extrapolation. Evidence: [campaign spec](specs/sm120-qwen35-pareto-2026-08-09.md), [production baseline](../docs/bench-evidence/qwen35-4b-sm120-main-20260807.md) | | 3 | `ROAD-V1-C3` | MTP k=1 + GDN speculative path, then DFlash, DSpark and heterogeneous-vocabulary TLI | [engine matrix](engine-matrix.md), [coverage view §8](feature-matrix.md#8-speculative-decoding) | MTP and DFlash specs exist. **M-mtp-0 CLOSED 2026-07-24** - the standalone MTP draft head is oracle-parity-proven on BOTH gate checkpoints (op-level parity vs a dumped k=1 vLLM oracle, not a token-generation SACRED gate). **I2 SCHEDULER-HALF LANDED 2026-07-24** ([mtp-spec-decode §2.7](specs/mtp-spec-decode.md)): host-side spec-decode scheduler/engine plumbing + the FROZEN spec-metadata ABI that I3 (rejection sampler) and I5 (verify/propose runner) build against - `SpeculativeConfig`, `DraftTokenIds`, `Request::spec_token_ids`/`NumTokensWithSpec`, the first population of `scheduled_spec_decode_tokens`, `Scheduler::update_draft_token_ids`, the `take_draft_token_ids` seam, `EngineCore::post_step`, `InputBatch::num_accepted_tokens`/`update_req_spec_token_ids`; DEFAULT-OFF and INERT (no `SpeculativeConfig` => `num_lookahead_tokens == 0` => byte-identical engine). `SPEC-MTP` **STAYS `GATING`** because M-mtp-1..4 (greedy rejection, GDN spec slots, k>1, CUDA graphs) are still open, so spec decode remains user-invisible. DSpark is user-promoted scope with DeepSeek-V4/Qwen3 draft models, reduced-vocabulary handling and full-CUDA-graph behavior inventoried under `SPEC-DSPARK`; tokenizer-agnostic target<->draft mapping is separately inventoried as `SPEC-TLI`. Their dedicated spikes are not written **I3 GREEDY REJECTION SAMPLER LANDED 2026-07-24** (`SPEC-REJECTION` `READY` -> `ACTIVE`): per-request logits expansion to `1 + k_i` rows plus the greedy accept rule (accept a draft iff it equals the target argmax; on the first mismatch emit the target argmax and stop; bonus token when all k accept), CUDA==CPU bit-exact at vocab 248320. **I4 GDN-HALF LANDED 2026-07-24** (`SPEC-GDN-SEGMENTS` -> `ACTIVE`): the GDN spec metadata split + decode->prefill reclassification, the `T>1`/`IS_SPEC` recurrence with per-timestep snapshots, the conv sliding window honouring `num_accepted`, k+1 slot allocation - bit-exact rollback. **I5a GDN LAYER ROUTING + runner spec-metadata upload LANDED 2026-07-24** (`CLAIM-SPEC-MTP-I5A`): `GdnBlockPaged` now routes a pure-spec batch through `vt::GdnSpecDecode`/`vt::CausalConv1dSpecUpdate` and the runner uploads I4's six spec device tensors - first of the scoped M-mtp-1 sub-increments (I5a GDN wiring -> I5b prepare_prefill -> I5c MTP paged propose -> I5d config+runner-loop+27B token gate, spec §5), DEFAULT-OFF INERT, bit-exact vs the I4 ops, no e2e loop yet. **I5b `prepare_prefill_inputs` LANDED 2026-07-24** (`CLAIM-SPEC-MTP-I5B`): the drafter prefill input-prep - shift each request's `input_ids` left one within its query span, splice the just-sampled next token, `query_len -= num_rejected`, emit last-token index / query_start_loc / seq_lens into the `SpecPrefillInputs` struct; a HOST routine (no new CUDA kernel; mirrors our DEVICE-NEUTRAL `prepare_inputs`/`combine_sampled_and_draft_tokens` family), unit-gated 7 cases / 27 assertions RED-first, DEFAULT-OFF INERT, additive by construction. **I5c MTP PAGED PROPOSE + DRAFT KV LAYER LANDED 2026-07-24** (`CLAIM-SPEC-MTP-I5C`): `Qwen3_5MTPModel::ForwardPaged` runs the head + one full_attention decoder layer over the head's OWN paged draft KV layer (ReshapeAndCache + PagedAttention over the target's block table / slot mapping); `MakeQwen3_5KVCacheSpec(num_spec>0)` adds that draft KV layer (`fa_draft` FullAttentionSpec group, index num_hidden_layers); `ForwardDeviceTap` exposes the `[T,H]` post-final-norm hidden tap (INERT); and `MtpProposePrefill` is the callable k=1 propose (I5b shift-splice -> one paged forward -> argmax at last_token_indices, early-exit). CORE PROOF: the paged forward reproduces I1's standalone head logits/argmax on BOTH gate checkpoints; a two-step drive proves the draft-KV write/read (RED control diverges). DEFAULT-OFF INERT (no spec config -> draft KV layer not allocated, tap nullptr, target forward byte-identical); NOT wired into the runner step loop. **I5d-pre REGISTRY/FORWARD-SEAM ENABLING REFACTOR LANDED 2026-07-25** (`CLAIM-SPEC-MTP-I5D-PRE`): a scoping pass found the model seam is fully TYPE-ERASED, so the runner cannot reach the concrete target weights / hidden-state tap / loaded MTP weights the I5d loop needs. Four ADDITIVE, inert-when-spec-off access paths + one latent-bug fix - the `hidden_tap` out-field on the type-erased `ModelForwardInput` (routes to the existing `ForwardDeviceTap`), a `LoadedModel::BuildMtpDraft` virtual (typed path to the draft, null for non-MTP), MTP weight loading + shard retention in `FromModelDir` behind `EngineParams::speculative_config`, and the `GPUModelRunner` ctor widened with optional draft/draft-KV/`SpeculativeConfig`; PLUS the latent `initialize_kv_cache` fix (select the FIRST non-eagle full-attn group as the target so a third `fa_draft` group can't displace it; byte-identical at num_spec==0). DEFAULT-OFF INERT, unit-gated RED-first, spec-off SACRED gates byte-identical. **I5d CONFIG RUNTIME + VERIFY/PROPOSE RUNNER LOOP LANDED as a spec-off-byte-identical PARTIAL 2026-07-25** (`CLAIM-SPEC-MTP-I5D`): `--speculative-config` JSON parse -> `EngineParams` -> `LoadedEngine` resolution (widened KV `MakeQwen3_5KVCacheSpec(num_spec>0)`, `BuildMtpDraft`, forced sync scheduling, `MakeScheduler(spec)`, `EngineCore(check_for_draft=true)`) + the full runner loop (draft splice, hidden-tap capture, GDN builder spec-overload feed, k+1 GDN state-slot remap + widened conv cache + draft-KV alloc, `MtpProposePrefill`, `take_draft_token_ids`, acceptance telemetry). CUDA `-Werror` 0 warnings, cutlass-ON; spec-OFF SACRED byte-identical (27B 235/235, 35B 315/315, Coder 138/138 + spec unit tests ALL PASS). **The three-way 27B token gate is NOT yet passing**: the spec-ON engine RUNS the loop end to end and MEASURES the blocker (`test_qwen27_spec_decode`) - it throws on the FIRST prefill step at `gdn_state_gather: working/cache row shapes must match` (`src/vt/ops.cpp:1773`) because I4's spec conv rollback needs the conv row widened to `(K-1)+num_spec` while the non-spec GDN conv ops assume `(K-1)`. `SPEC-MTP` STAYS `GATING`. **I5e LANDED 2026-07-25 (`CLAIM-SPEC-MTP-I5E`): the non-spec GDN conv ops made widened-cache-aware (mirror vLLM `state_len=KERNEL_WIDTH-1` + physical `stride_conv_state_tok`, leading `(K-1)` sub-window, byte-identical at `num_spec==0`) AND the async input-combine forced off under spec (it overwrote the verify batch's draft position with the committed token -> 0 acceptance, RCA'd on the real 27B). **THE THREE-WAY 27B SINGLE-REQUEST GREEDY GATE PASSES**: our-spec-ON == vLLM `--speculative-config mtp` greedy == our-spec-OFF token-for-token, **acceptance 16/16 drafts accepted** (~16 target steps saved); spec-OFF SACRED byte-identical (27B 235/235, 35B 315/315, Coder 138/138), compute-sanitizer 0 on the spec step. `SPEC-MTP` LEAVES `GATING` (single-request greedy correctness PROVEN); NOT `DONE` - the MIXED `GdnBlockPaged` concurrency split/merge + the throughput A/B vs vLLM same-config are I6. **I6 LANDED 2026-07-25 (`CLAIM-SPEC-MTP-I6`, `benchmark_binding=true`): the §5 c1 THROUGHPUT GATE — ours spec-ON AT/ABOVE vLLM spec-ON on EVERY measured axis at c1** (TPOT 66.2/62.95 vs 69.1/65.3 ms prose/code, ours ~1.04x faster; output tput +4.6%/+3.9%; ITL/TTFT lower; acceptance ours 0.85/0.92 vs vLLM 0.838, within noise; spec helps both ~1.5-1.6x TPOT; ours ~4% faster spec-OFF too), via an additive example-only `--speculative-config` bench flag (NO engine code touched). STAYS `ACTIVE`: the c>1 mixed-batch path is still refused + owes a c>1 A/B, and no server-facing spec flag yet. | `ACTIVE` | M-mtp-0, I2 scheduler-half, I3 rejection sampler, I4 GDN spec slots, I5a GDN layer routing, I5b prepare_prefill, I5c MTP paged propose + draft KV, I5d-pre the registry/forward enabling seam, and I5d config runtime + verify/propose runner loop (spec-off byte-identical) are landed; next (before `SPEC-MTP` leaves `GATING`) is closing the measured I5d gate blocker - make the non-spec GDN conv ops widened-cache-aware (mirror vLLM `causal_conv1d` `state_len=width-1+(seqlen-1)`) + the MIXED `GdnBlockPaged` split/merge - then the passing M-mtp-1 27B k=1 greedy three-way token gate + acceptance, then M-mtp-2 35B, then DFlash, the DSpark spike/gates and TLI. **DFlash D0-redo + D1 LANDED 2026-07-26 (`CLAIM-DFLASH-D0D1`, [dflash-spec-decode §0](specs/dflash-spec-decode.md)): `SPEC-DFLASH` UNBLOCKED + `ACTIVE` on the advanced pin `555967922`/vLLM 0.26.0.dev0.** The prior 0.25.0 ORACLE-BLOCKED verdict is SUPERSEDED — under `VLLM_USE_V2_MODEL_RUNNER=1` (vllm#40898 resolved) the mixed-SWA/full z-lab 27B draft CONSTRUCTS and the drafter is ALIVE (acceptance 2.21/8.80/4.75/4.57 > 1, `num_spec=16`, flashinfer-native fp8-KV; goldens committed). Gate FORM measured STRICT MODE-MATCHED (vLLM-ON run-deterministic K>=3 but != vLLM-OFF at k=16 near-ties — NOT the MTP three-way identity). **D1 `DF-AUX-TAPS` DONE:** the single hidden tap is generalized to the multi-tap `[T,H×taps]` (`ForwardDeviceMultiTap` capturing `(hidden+res)` at `target_layer_ids`), config-gated byte-identical off; unit gate 598 assertions (RED-first), CUDA 697/697 + sanitizer 0, 27B MTP e2e 9/9 + 27B SACRED 235/235 byte-identical (inertness). **D2-D5 LANDED 2026-07-26 (`CLAIM-DFLASH-D2`/`D3`/`D4D5`/`D5`):** the drafter model + the project's first non-causal in-block attention (D2, GPU parity vs the real vLLM draft), context-KV precompute + `prepare_dflash_inputs` (D3, GPU numeric-parity 61/61), the non-autoregressive whole-block propose brick + `dflash` config-select (D4), and the RUNNER-LOOP INTEGRATION + 27B e2e (D5): the full verify/propose loop is wired (separate z-lab draft load + target-shared bf16 embed/lm_head, aux-tap capture, per-request combined-feature context accumulation honoring num_rejected, `propose_drafts_dflash`) and RUNS end to end - `test_qwen27_dflash_spec_decode` 2/4 STRICT token-exact vs the vLLM-DFlash-ON golden + acceptance ~ vLLM on ALL 4 (19/39/29/25 vs 17/39/30/25). The 2 divergences are SINGLE bf16 near-tie flips (ratified near-tie ROOT = the D3 inline context-KV recompute envelope), NOT a wiring bug; inertness SACRED 235/235 + MTP 9/9 byte-identical; CUDA `-Werror` clean, no new kernel. NOT a clean strict-4/4 pass - STRICT 4/4 token-identity + the speed A/B = D6 (persistent paged draft-KV bit-matching vLLM's fused projections + the uniform-1+k FULL CG). Capture tool + goldens: `scripts/spec/d{0,2,3}_dflash_*.py`, `tests/parity/goldens/dflash_27b{,_draft,_kvprep}/`. **D6-D9 SPEED CAMPAIGN 2026-07-27:** D6 c1 A/B + STRICT-4/4 bf16-irreducibility RCA; D7 device-resident within-step forward (bit-identical); D8 acceptance RCA + FINAL golden A/B (ours 0.69× vLLM). **D9 (`CLAIM-DFLASH-D9`) PERSISTENT PAGED DRAFT-KV LANDED (bit-identical, +22.7%): `AppendContextKVHost` + `ForwardBlockLogitsWithPrecomputedKV` replace the O(context²) per-step recompute with an append-only per-request store; ours-ON 20.99→25.75 tok/s = 0.917× vLLM-ON (28.09, was 0.69×); e2e 27/27 SAME tokens, SACRED 235/235 + MTP 9/9 byte-identical, CUDA `-Werror` clean, no new kernel. D8's "bf16 acceptance ceiling" REFUTED — same-trajectory per-step acceptance == vLLM (ratio 1.00) and ours realized acceptance (3.68/step) > vLLM (3.31); the SOLE residual (~8%) is the FULL uniform-(1+k) CUDA graph (eager-vs-graphed), a closeable increment. SPEC-DFLASH stays `ACTIVE` (speed not yet ≥ vLLM).** **SPEC-MTP → `DONE` 2026-07-26 (`CLAIM-SPEC-MTP-DONE`, records-only, closing commit I7 `72f9fb1`):** the user ratified the c>1 near-tie+SPEED criterion, closing both I6-owed items (mixed-batch concurrency + server/CLI/C-ABI `--speculative-config`); MTP k=1 is COMPLETE + gated. **M-mtp-2 CLOSED 2026-07-26 (`CLAIM-SPEC-MTP-M-MTP-2`): the 35B `Qwen3_5MoeMTP` full e2e three-way token gate PASSES** — our spec-ON == our spec-OFF == vLLM 0.25.0 `--speculative-config mtp` greedy == vLLM spec-OFF, 16/16 vs the `greedy_ids` anchor (STRICT, c1), acceptance 16/16 both sides; c1 spec-ON 1.19x TPOT / +16.3% output-tput vs spec-OFF (acceptance 0.908) — MoE speedup transfers; spec-OFF byte-identical (test+docs-only). MTP is now `DONE` on BOTH gate models (`MODEL-SPEC-qwen3-5-mtp-qwen3-5-moe-mtp` `GATING`→`DONE`). **DFlash D11+D12 2026-07-27 — the FULL uniform-(1+k) CUDA graph is being built in three parts:** D11 (`CLAIM-DFLASH-D11`) landed Part A (the device-store primitive, CPU-gated); **D12 (`CLAIM-DFLASH-D12`) landed A-wire (the D11 device store is now the PRODUCTION path; GPU-gated e2e 27/27 all-exact acceptance 19/39/29/25 + SACRED 235/235 + MTP 9/9 byte-identical) + Part B (`vt::DFlashPagedBlockAttention`, the capture-safe paged kernel; `test_ops_dflash_paged_block_attn` 795648/795648 CPU==CUDA + cross-check vs materialized `DFlashBlockAttention` + compute-sanitizer 0; NO function-local host `cu_seqlens` upload = capture-UAF fixed).** Speed UNCHANGED 0.917× (A-wire eager + Part B not yet wired). The SOLE remaining piece is Part C (static-shape capture + device mask-scatter + `BeginCapture`/replay + the ≥vLLM c1 A/B); if ours-ON-graphed ≥ vLLM-ON → SPEC-DFLASH DONE → C3 complete. C3 stays `ACTIVE` (DFlash Part C + DSpark/TLI remain) | | 4 | `ROAD-V1-C4` | Quantization: llama.cpp breadth/speed, NVFP4/FP8/MX, MLX native | [quantization matrix](quantization-matrix.md) | coverage spike merged; `QUANT-GGUF-CPU-THREADPOOL` W1-W3 implemented and correctness-gated, still `GATING` (its reproduction now exists — same-binary 1-vs-20-thread A/B is prefill 12.47x / decode 8.05x / RSS 1.000x, so **decode misses the >=10x bar**). **GGUF COMPUTE-IN-QUANT IS NOW LIVE AND DEFAULT-ON (2026-07-22, `CLAIM-QUANT-GGUF-CIQ-G4-1`):** [compute-in-quant GEMM](specs/gguf-compute-in-quant-gemm.md) **G1-G4** — block dtypes + traits, the Q8_0/Q8_K activation quantizers, the six generic `vec_dot`, `kMatmulBTQuant`, and now the ROUTING (`vt::MatmulBT` dispatches a block-dtype weight to it) — plus [keep-quant loader](specs/gguf-keep-quant-loader.md) **L1-L4**, whose master switch defaults ON wherever that op is registered for the running device (CPU today; a CUDA build still expands). Six encodings (Q4_0, Q8_0, Q3_K, Q4_K, Q5_K, Q6_K) now carry `C` = `Y`. **Correctness held exactly**: the 35B GGUF gate is 16/16 token-exact vs the same-file llama.cpp oracle with the quant path on, and the bench model's output tokens are byte-identical across the pre-G4, post-G4 and `VT_CPU_REF=1` arms — no golden regenerated. **Binding CPU A/B** (idle dgx aarch64, one flock, same binary, 3 reps): decode **3.45x**, prefill **4.16x**, peak RSS **1.16x less**; vs llama.cpp we went from 11.7x / 34.1x / 2.66x behind to **3.38x / 8.20x / 2.29x**. The projected 9-17x did **not** hold, for a measured reason: 60 % of that file's weight bytes are `f16`, which no block encoding covers. Keep-quant loader **L4** is therefore MEASURED-and-NOT-MET on RSS; other leaf specs open. **THAT #1 LEVER IS NOW LANDED (2026-07-22, `CLAIM-KERNEL-CPU-ELEM-GEMM-1`, new row [`KERNEL-GEMM-CPU-ELEM`](kernel-matrix.md)):** [the elementwise CPU GEMM](specs/cpu-elementwise-gemm.md) **E1-E4** — per-dtype specialization out of the K loop, 16 independent accumulators instead of one, AArch64 NEON + x86-64 SSE2/F16C tiers behind a runtime probe, and M-blocking — all **BYTE-IDENTICAL** to the historical kernel (`memcmp` gate, exhaustive 65,536-pattern widening check, same token md5), so nothing was regenerated. Binding same-binary A/B: prefill **3.41x**, decode **3.11x**; op-level bf16 18-24 -> 69-351 GFLOP/s. **vs llama.cpp: decode AT PARITY (1.03x), prefill 2.34x behind, RSS 2.29x worse. THEN loader L5 LANDED (2026-07-23, `CLAIM-QUANT-GGUF-KEEPQ-L5-1`):** [keep-quant loader](specs/gguf-keep-quant-loader.md) **L5** — mmap in-place residency (borrow kept blocks out of the read-only mapping, refcounted), tied-head sharing (one bf16 vocab matrix for embed+lm_head), and a read-once page release (port of llama.cpp `unmap_fragment`) — took **peak RSS 6.401 -> 3.884 GiB (2.29x -> 1.39x llama.cpp)** with decode UNCHANGED and output tokens byte-identical (md5 `d235db12f2cd304007530286a1755c95`). The remaining ~1.09 GiB over llama.cpp is the f16 expansion (no block encoding covers f16). | `PARTIAL` | **THE OWED FRESH PROFILE IS DONE (2026-07-23) and it re-ranks the plan.** A `vt::GetOp` hook (100% of wall time, reverted before binding) on the CURRENT binary: prefill is **no longer GEMM-bound** — kMatmulBTQuant 37%, **kGdnPrefill 25%**, kMatmul 12%, kMatmulBT 10%, **kPagedAttention 10%**; the two non-GEMM kernels (GDN linear-attention recurrence + paged attention) run **SINGLE-THREADED** on the CPU and are now the top prefill levers. Decode is memory-bound matmul at parity, no kernel work owed. **RE-RANK: G5/G6/G7 all only speed the already-fast quant GEMM and rank BELOW the two serial non-GEMM kernels; the new #1 CPU lever is threading kGdnPrefill + kPagedAttention.** **THAT #1 LEVER IS NOW LANDED (2026-07-23, `CLAIM-CPU-THREAD-GDN-PAGED-1`, [two-kernel threading](specs/cpu-thread-gdn-paged-2026-07-23.md)):** kGdnPrefill chunks over the (sequence, value-head) axis and kPagedAttention over query-token rows, both via the existing `ParallelForRows`, both **byte-identical** (qwen35 output-token md5 `d235db12f2cd304007530286a1755c95` unchanged at threads 1/4/20 + `VT_CPU_REF=1`, determinism battery extended, CPU ctest 158/158). **Binding dgx aarch64 (idle): prefill 1.382× same-binary (73.0→100.9 t/s), 2.43×→1.76× behind llama.cpp pp128; decode at parity; op-scaling 1→20 GdnPrefill 7.08× / PagedAttention 8.96×; fresh profile shows the two kernels 35%→8.6% of prefill and re-ranks the NEW bottleneck to the GEMMs (kMatmulBTQuant 50% + kMatmul 16% + kMatmulBT 14% = 80%) ⇒ next CPU lever is the SIMD/repack GEMM tiers (G5/G6/G7).** **THE FIRST SUCH TIER IS NOW LANDED (2026-07-23, `CLAIM-QUANT-GGUF-CIQ-G6-1`, [compute-in-quant GEMM](specs/gguf-compute-in-quant-gemm.md) G6):** the Arm **i8mm mmla `nrc==2`** `vec_dot` 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 (decode M=1 → portable, unchanged); runtime `HWCAP2_I8MM` probe + `VT_CPU_QUANT_MMLA` defeat + per-file `+i8mm`. **BYTE-IDENTICAL** where the math allows (q8_0/q4_0 bit-exact to the scalar tier, q4_K/q6_K within NMSE ≤ 5e-4), bit-identical across threads 1/2/4/20, e2e token md5 `d235db12f2cd304007530286a1755c95` byte-identical (mmla on/off/`VT_CPU_REF=1`), 35B GGUF gate 16/16 vs llama.cpp on both files. **Op-level portable→i8mm: q4_K 7–8.4×, q6_K 3.8–4.5×, q8_0 ~1.2×**; e2e prefill same-binary **1.084×** on the q8_0-dominant bench file (1.56×→**1.44× behind** llama.cpp pp128, Amdahl-bounded — the big k-quant win lands on the APEX 35B files). Fresh bottleneck: the elementwise f16/f32 GEMM (~30%, unchanged) is now co-dominant on this mixed file. CUDA `-Werror` 0-warn, regression set UNCHANGED. `docs/BENCHMARKS.md` ACCEPTED. RSS deficit closed to 1.39x by L5; the last RSS lever is an f16 keep-as-is compute path, not this loader. **THEN the GDN split-projection orientation LANDED (2026-07-23, `CLAIM-CPU-GDN-ORIENT-1`, [GDN projection orientation](specs/cpu-gdn-proj-orientation-2026-07-23.md)):** a fresh op-dispatch profile of the current binary (warm prefill, `vt::GetOp` hook + per-GEMM shape histogram, reverted before binding) found the four GDN input projections (`in_proj_qkv/z/b/a`, 72 GEMMs, **17.9%** of prefill: `kMatmulBTQuant 50.7% / kMatmul 17.9% / kMatmulBT 14.9%`) were the LAST weight family `LoadGdnGguf` still transposed into [K,N] (nk=false → the N-striding `kMatmul`, no M-blocking) after G4's `expand_nk` gave every other expanded weight the file's own [N,K] order. New `GgufLoadPolicy::gdn_expand_nk` + `MakeGdnProj` keep them [N,K] nk=true → the M-blocked `kMatmulBT`; **BYTE-IDENTICAL** (same sequential f32 K-reduction, only the weight offset differs — token md5 `d235db12f2cd304007530286a1755c95` unchanged across on/`VT_GGUF_GDN_NK=0`/`VT_CPU_REF=1` and threads 1/4/20), `test_qwen36_gguf_engine` 2/2·28/28·16/16 on APEX. **Binding same-binary prefill 1.090× / decode 1.09× (44.1→40.4 ms TPOT = 1.01× llama tg32, at parity), 1.44×→1.32× behind llama.cpp pp128, RSS unchanged.** Fresh post-change profile: `kMatmul` **17.9%→0% (72→0 calls, ELIMINATED)**, absorbed into `kMatmulBT` (14.9%→27.7%); **next CPU prefill lever = the quant GEMM (kMatmulBTQuant, now 55%): G7 repack-at-load.** **G7 LANDED 2026-07-23 (`CLAIM-QUANT-GGUF-CIQ-G7-1`, [compute-in-quant GEMM](specs/gguf-compute-in-quant-gemm.md) G7):** q8_0 repacked once at load into the i8mm `block_q8_0x4` interleave (ported llama.cpp `repack.cpp` `q8_0_4x8`), `kMatmulBTQuant` dispatches a pre-shuffled gemm/gemv with no per-block register shuffles. **BIT-IDENTICAL** (byte-permute weight + non-fused `vmlaq_f32`, 305-assertion memcmp across decode/prefill/bf16-out/threads, token md5 `d235db12f2cd304007530286a1755c95` unchanged on/`VT_CPU_QUANT_REPACK=0`/`VT_CPU_REF=1`; a `ResidentWeight`/`MakeTensor` flag-drop that produced all-zero tokens was caught by the E2E gate and fixed). Op-level q8_0 **3.7–5.9×** (518→2401 / 583→3456 / 514→1902 GFLOP/s); **E2E prefill 1.92× same-binary (1096→572 ms), 223.8 t/s vs llama.cpp pp128 177.3 = 1.26× — AT/BEYOND PARITY** (was ~1.5× behind), decode at parity, RSS 3.884 GiB unchanged. Fresh profile: q8_0 GEMM 55%→~21%; **the CPU prefill-lever search is CLOSED — the sole remaining gap to llama.cpp is peak RSS (1.39×), not prefill.** CUDA-inert (gated off any non-CPU-quant device), CUDA `-Werror` 0-warn, regression set UNCHANGED. **THE RSS GAP IS NOW CORRECTLY ATTRIBUTED (2026-07-23, `CLAIM-QUANT-GGUF-KEEPF16-L6-1`, [keep-quant loader](specs/gguf-keep-quant-loader.md) L6): it is NOT the f16 expansion.** L6 implemented keep-f16 residency (keep the file's 56 F16 weights + tied head resident as F16 and compute on them, mirroring llama.cpp `ggml_vec_dot_f16`) and MEASURED it **RSS-NEUTRAL** (3.884 → 3.832 GiB, −52 MB) and prefill-regressive (TTFT 577 → ~1000 ms, from 1.25× ahead of llama.cpp to 0.72× behind) — because L5's page-release had ALREADY dropped the f16 file pages, so keep-f16 only swaps an anonymous bf16 buffer for equal-size file-backed f16 pages. smaps attribution proves our weight residency is at llama.cpp parity (file-backed 2.63 ≈ 2.68 GiB); **the residual ~1.08 GiB is the engine's ANONYMOUS activation/KV workspace, not weights — the real, separate CPU RSS lever.** keep-f16 ships DEFAULT OFF (`VT_GGUF_KEEP_F16=1` opt-in), tokens byte-identical (md5 `d235db1…`), `test_gguf_keep_quant` 35/35 (x86+aarch64), regressions UNCHANGED (27B 235/235, 35B 315/315, Coder 138, dense 184, OPT 63, DeepSeek 223, Llama 92, GGUF engine 28/28). **NEXT CPU RSS lever: profile + shrink the engine's activation/KV working set, NOT the weight loader** | | 5 | `ROAD-V1-C5` | Sliding window, local attention, YaRN/long context | [engine matrix](engine-matrix.md), [coverage view §§2,11](feature-matrix.md#2-kv-cache--memory), [joint spike](specs/sliding-local-yarn-long-context.md) | **CUDA GPU CLOSURE 2026-07-27 (`CLAIM-ROADMAP-C5`, dgx GB10 sm_121a, clean build of `489f7771`, oracle vLLM 0.26.0.dev0):** the shared scaled-RoPE + local-mask CUDA path COMPILES `-Werror`-clean and RUNS on GB10; the C5 feature-positive correctness gates that were the stated `GATING` blocker now PASS — SWA (Gemma-2/Gemma-3 48/48), LongRoPE (Phi-4-mini 16/16, RED-first), llama3-rope (Llama-3.2-1B 16/16), dynamic-NTK (InternLM2 16/16); both RoPE 0.26-oracle recaptures BIT-IDENTICAL to goldens (zero drift). Leaves `ATTN-SLIDING-WINDOW`/`ATTN-ROPE-{LLAMA3,LONGROPE,DYNAMIC-NTK}`/`ATTN-YARN` → `ACTIVE` | `PARTIAL` | (RI) **Honest residual (vehicle-blocked, not skipped):** YaRN model e2e (no cached Nomic/gpt-oss consumer) + chunked-local model e2e (no Llama4 row) are REACHABLE-BLOCKED — operator/formula stay GPU/G3-gated; long-context positive-mask (prompt > W) SWA model e2e + the KV memory-optimization G8; and the roadmap-wide every-axis SPEED tail (all C5 leaves correctness-complete, speed-pending, mirroring their model consumers). Not row-DONE until speed + the blocked vehicles close | diff --git a/.agents/specs/sm120-qwen35-pareto-2026-08-09.md b/.agents/specs/sm120-qwen35-pareto-2026-08-09.md index 572b58ed9..2b9954d5e 100644 --- a/.agents/specs/sm120-qwen35-pareto-2026-08-09.md +++ b/.agents/specs/sm120-qwen35-pareto-2026-08-09.md @@ -8,8 +8,8 @@ VRAM gaps vs vLLM](https://github.com/mudler/vllm.cpp/issues/206) **Lifecycle:** combined prefill opt-in accepted locally; frontend pretokenization is CPU-gated but its first real GPU A/B failed token identity because sequential -enqueue changed the first scheduling wave; atomic wave admission is specified -below and pending implementation/review +enqueue changed the first scheduling wave; atomic wave admission is implemented +and CPU-green, with fresh mutation review and the restarted GPU A/B pending **Target:** local NVIDIA GeForce RTX 5070 Ti (`sm_120`), Qwen3.5-4B plain BF16. This vehicle does not substitute for the hardware-unavailable 27B/35B @@ -219,6 +219,38 @@ then restarts the complete counterbalanced `STRINGa -> TOKENSa -> TOKENSb -> STRINGb` series; all prompt/output IDs must match before any latency or throughput value is interpreted. +### Atomic wave-admission implementation checkpoint (2026-08-09) + +The RED compile failed on the deliberately absent queue-batch, AsyncLLM-wave, +and benchmark-wave APIs. The additive implementation now prepares the complete +ordered string or token-ID wave before taking the output-processor admission +lock, registers every collector under that lock, and publishes all core ADDs +through one queue operation whose mutex spans the full append. A failed element +move rolls back only the appended deque suffix; duplicate/preparation/enqueue +or shutdown failure removes only frontend states created by the call and emits +no terminal output. Existing single-request APIs are unchanged. + +The benchmark records every selected arrival before exactly one string/token +wave publish for the initial concurrency and every refill. Initial mutation +review found three routing checks that were observational rather than binding: +both benchmark callbacks could be replaced with per-request adds, the client +could replace its atomic put-many with per-item puts, and the lock-owned life +recheck could be removed while the probabilistic shutdown test stayed green. +The RED-first hardening tests then failed to compile because production exposed +none of the required injectable ownership seams. `RunBench` now delegates the +actual `add_request_wave` call to a fake-engine-tested seam, the client delegates +the actual queue publish to a put-many-only fake-queue-tested seam, and the +late-life guard seam itself acquires the admission mutex before evaluating the +alive predicate and invoking publish. A deterministic fake lock proves a +submitter can pass the outer check, lose to shutdown, and neither register nor +publish. + +Focused `-Werror` green evidence: complete `test_engine_core_proc` 14/14 cases +and 111 assertions, complete `test_async_llm` 13/13 and 435 assertions, and +complete `test_bench` 11/11 and 80 assertions. This is an implementation +checkpoint only: fresh mutation re-review, operator gate, exact real-GPU token +identity and performance acceptance remain pending. + ## Correctness, tests, review and acceptance Before product code, a fresh implementer ports or restores the smallest tests diff --git a/docs/BENCHMARKS.md b/docs/BENCHMARKS.md index d3eb43938..c847ddb97 100644 --- a/docs/BENCHMARKS.md +++ b/docs/BENCHMARKS.md @@ -41,7 +41,7 @@ The binding comparison. vLLM runs its **production graphed config**, never | Qwen3.6-27B | NVFP4 (`nvidia` @`0893e160`, ModelOpt `modelopt_mixed`) | 0.25.0 | 0/4 | **BEHIND, uniformly 0.85x** on decode throughput (was 0.72x before the FP8 tower fix); greedy continuation IDENTICAL to vLLM. A different model from the `unsloth` row (NVFP4 MLP + FP8 W8A8 GDN/attn tower) | | Qwen3.6-35B-A3B | NVFP4 `modelopt_mixed` | 0.25.0 | 2/18 | 3-rep grid 2026-08-05 @`1ea26427`: 0.93-1.03x (c4 wins), c16 0.93x. Both c16 levers A/B'd NEG: drain event -1.9%, mirror 0.999x. ★ probe found a prod async batch-1 greedy DEGENERATION bug the mirror fixes | | DeepSeek-V2-Lite | bf16 MLA | 0.25.0 | 4/25 | Attributed miss, row stays `ACTIVE` | -| Qwen3.5-4B | bf16 direct-load | 0.26.0.dev0 | host PSS; frontend exactness failed | #206 pretoken A/B FAILED (98/128); timings VOID, atomic-wave retry pending. Opt-ins: +0.718% tput/-0.556% TPOT; cross-engine tput/TTFT VOID, TPOT/VRAM OPEN ([data](bench-evidence/qwen35-4b-sm120-main-20260807.md)) | +| Qwen3.5-4B | bf16 direct-load | 0.26.0.dev0 | host PSS; frontend exactness failed | #206 pretoken A/B FAILED 98/128; timings VOID. Atomic-wave retry pending (CPU-green). Opt-ins: +0.718% tput/-0.556% TPOT; tput/TTFT VOID, TPOT/VRAM OPEN ([data](bench-evidence/qwen35-4b-sm120-main-20260807.md)) | ### GDN prefill causal-convolution by GPU diff --git a/docs/STATUS.md b/docs/STATUS.md index e5068a5c8..e9d0e5286 100644 --- a/docs/STATUS.md +++ b/docs/STATUS.md @@ -78,7 +78,7 @@ token-for-token correctness against the pinned oracle. | Qwen3.6-27B (NVFP4) text generation | Correctness-complete; speed is CHECKPOINT-dependent | Token-exact GB10 on both. `unsloth` @`890bdef7` beats vLLM 0.25.0 every c (1.007-1.045x), 115/124; `nvidia` @`0893e160` (ModelOpt FP8 tower) is **0.85x BEHIND**, decode ~100% GPU-busy | | Qwen3.6-35B-A3B (NVFP4, GDN MoE) | Correctness-complete; binding grid @`a0fa12c7` FLAT 0.935x-0.979x over c1-c32 (CoV <0.81%); the prior 0.87x c2 / 0.92x c8 "weak cells" were harness mismatch, not code; memory PSS 3.81x, GPU 1.40x | Token-exact SYNC+ASYNC; `VT_ASYNC_DEVICE_MIRROR` ON fixes async batch-1 token-0 degeneration; `VT_ASYNC_EXECUTOR` Option A NEUTRAL → OFF | | Qwen3 / Qwen2 dense (BF16) | Correctness-complete, speed-pending. Async-serving P0 FIXED (`ROW-SERVE-ASYNC-DENSE-MIRROR`): classic-dense `Qwen3ForCausalLM` now honors the async device token-ids mirror; CPU-only -Werror test-guard fixes x2 | Near-tie-robust token-exact vs vLLM (Qwen3-0.6B, Qwen3-4B); c1 effective parity, c8 decode residual. **Async device-mirror (`ROW-SERVE-ASYNC-DENSE-MIRROR`, `f9c969ae`): the #31 fix ported to the classic dense family, dgx-VERIFIED.** The shared dense `EmbedInto` (qwen3.cpp) raced the async combine's device input-ids write against a stale host upload → token-0 degeneration on the depth-2 AsyncLLM serving path (quant-independent). `EmbedInto` now consumes the device override published by `ForwardQwen3ForCausalLM`'s `DeviceTokenIdsScope` (27B-dense template); gate `test_qwen3_dense_async_serving` RED on `VT_ASYNC_DEVICE_MIRROR=0`, GREEN default, byte-identical mirror-off. dgx GB10: async gate RED→GREEN 0.6B+4B, SACRED 0.6B+4B 184/184 unchanged (byte-neutral sync path), memcheck 0 errors; Yi30/Qwen3-8B-MXFP4 default-config e2e coherent + 3/4 token-exact (p2 = oracle-ratified near-tie, gap 0.0000), closing the QUANT-CT-MXFP4 async-default residual. RESIDUAL: sibling InternLM2/Mistral/Llama scope one-liner; W4 bench RAN; FA2 GQA-swap default-ON, c2-c8 <1.0x. `FLASH-PTXAS` #82: codegen at PARITY (no ptxas lever); gap=engine context. **D1 (2026-07-31, `CLAIM-D1-BF16-MERGED-QKV`): the bf16 merged-QKV path (`Qwen3QkvMergeEnabled`/`VT_QWEN3_QKV_MERGE`) is now default-ON** — one `vt::MatmulBT` over the merged `[qdim+2kdim,H]` owner + a contiguous `vt::QkvSplit` (OLMo-2 exemplar), replacing three per-shard GEMMs. Bit-exact GEMM math (A/B unit `test_ops_qkv_merge` byte-identical, RED-first); the wider-N cuBLASLt K-reduction flips the 0.6B genuine bf16 near-tie so the SACRED 0.6B golden was regenerated (all tokens within the near-tie band, max 0.125 nats), while Qwen3-4B is byte-neutral (0 diffs, stays STRICT). Re-gated 0.6B 16/16 + 4B 16/16; consistency/launch-count fold (measured NEUTRAL on 4B decode), no new throughput owed | -| Qwen3.5-4B plain BF16 direct loading on discrete CUDA | Correctness-complete; atomic frontend wave pending; TPOT/VRAM open | #206 real pretoken A/B FAILED exactness (98/128); timings VOID. Atomic wave admission pending; TPOT 1.025x, VRAM +233 MiB OPEN ([data](bench-evidence/qwen35-4b-sm120-main-20260807.md)) | +| Qwen3.5-4B BF16 direct-load on discrete CUDA | Correctness-complete; atomic frontend CPU-green; TPOT/VRAM open | #206 pretoken A/B exactness FAIL 98/128; timings VOID. Atomic: async 13/13·435, core 14/14·111, bench 11/11·80; re-review/retry pending. TPOT 1.025x; VRAM +233 MiB ([data](bench-evidence/qwen35-4b-sm120-main-20260807.md)) | | Qwen3-Coder-30B-A3B MoE (BF16) | Correctness-complete, speed-pending | Near-tie-robust token-exact 6/6; 11 of 16 binding grid cells at or above vLLM. **D1 (2026-07-31): inherits the default-ON bf16 merged-QKV via the shared dense `AttnBlock` — byte-neutral (0 token diffs, golden UNCHANGED); re-gated 6/6** | | Llama-3.x dense (BF16) | Correctness-complete, speed-pending | Near-tie-robust token-exact 16/16 (Llama-3.2-1B); llama3 RoPE scaling | | Mistral dense (BF16) | Correctness-complete, speed-pending | Paged-engine token-exact 16/16 (Mistral-7B-v0.3) | diff --git a/docs/USAGE.md b/docs/USAGE.md index 0d78d159a..fd6fe4bd7 100644 --- a/docs/USAGE.md +++ b/docs/USAGE.md @@ -118,7 +118,10 @@ Two more example binaries ship alongside it: - `vllm-bench` ([`examples/bench/main.cpp`](../examples/bench/main.cpp)), a throughput/latency harness taking `--model`, `--dataset-path`, `--num-prompts`, `--input-len`, `--output-len`, `--concurrency`, - `--max-num-batched-tokens`, and `--num-blocks`. + `--max-num-batched-tokens`, and `--num-blocks`. It pretokenizes before timing + and atomically publishes each concurrency wave. Set + `VT_BENCH_PRETOKENIZE=0` for the timed-string rollback; the report names the + resolved mode. - `tokenize` ([`examples/tokenize/main.cpp`](../examples/tokenize/main.cpp)), a tokenizer smoke tool taking ` `. diff --git a/examples/bench/bench_core.h b/examples/bench/bench_core.h index f3937c24b..091cb41b0 100644 --- a/examples/bench/bench_core.h +++ b/examples/bench/bench_core.h @@ -204,6 +204,28 @@ inline decltype(auto) DispatchBenchPromptAdmission( return std::forward(timed_string_callback)(); } +// One benchmark refill is one observable engine wave. Record every selected +// request's arrival before invoking exactly one engine batch publish; in rollback +// mode, string tokenization therefore remains timed but begins only after the +// complete wave has the same arrival boundary as the pretokenized path. +template +inline decltype(auto) DispatchBenchPromptWaveAdmission( + Engine& engine, const char* env_value, std::size_t wave_size, + ArrivalCallback&& arrival_callback, + PretokenizedWaveFactory&& pretokenized_wave_factory, + TimedStringWaveFactory&& timed_string_wave_factory) { + for (std::size_t offset = 0; offset < wave_size; ++offset) { + std::forward(arrival_callback)(offset); + } + if (ResolveBenchPretokenizedAdmission(env_value)) { + return engine.add_request_wave( + std::forward(pretokenized_wave_factory)()); + } + return engine.add_request_wave( + std::forward(timed_string_wave_factory)()); +} + // Own the ordering boundary between workload preparation and measurement. // The clock callback is invoked exactly once and only after every default-path // prompt has been encoded with the same special-token policy as InputProcessor. @@ -596,27 +618,56 @@ inline BenchResult RunBench(const BenchConfig& cfg) { int done = 0; auto admit = [&]() { - while (next < cfg.num_prompts && in_flight < cfg.concurrency) { - const std::string rid = std::to_string(next); - RequestRecord rec; - rec.arrival_s = now_s(); - records[rid] = rec; - const size_t prompt_index = static_cast(next); - active.emplace( - rid, DispatchBenchPromptAdmission( - pretokenize_env, - [&]() { - return engine.add_request( - rid, std::move(pretokenized_prompts[prompt_index]), - MakeSampling(cfg, next)); - }, - [&]() { - return engine.add_request(rid, prompts[prompt_index], - MakeSampling(cfg, next)); - })); - ++next; - ++in_flight; + const int available = cfg.concurrency - in_flight; + const int wave_size = std::min(cfg.num_prompts - next, available); + if (wave_size <= 0) return; + const int wave_begin = next; + + std::vector admitted = + DispatchBenchPromptWaveAdmission( + engine, pretokenize_env, static_cast(wave_size), + [&](std::size_t offset) { + const int request_index = + wave_begin + static_cast(offset); + RequestRecord record; + record.arrival_s = now_s(); + records[std::to_string(request_index)] = std::move(record); + }, + [&]() { + std::vector wave; + wave.reserve(static_cast(wave_size)); + for (int offset = 0; offset < wave_size; ++offset) { + const int request_index = wave_begin + offset; + const std::size_t prompt_index = + static_cast(request_index); + wave.push_back(vllm::v1::AsyncTokensRequestInput{ + std::to_string(request_index), + std::move(pretokenized_prompts[prompt_index]), + MakeSampling(cfg, request_index), 0}); + } + return wave; + }, + [&]() { + std::vector wave; + wave.reserve(static_cast(wave_size)); + for (int offset = 0; offset < wave_size; ++offset) { + const int request_index = wave_begin + offset; + const std::size_t prompt_index = + static_cast(request_index); + wave.push_back(vllm::v1::AsyncStringRequestInput{ + std::to_string(request_index), prompts[prompt_index], + MakeSampling(cfg, request_index), 0}); + } + return wave; + }); + if (admitted.size() != static_cast(wave_size)) { + throw std::runtime_error("AsyncLLM admitted an incomplete bench wave"); + } + for (vllm::v1::AsyncRequest& request : admitted) { + active.emplace(request.request_id, std::move(request)); } + next += wave_size; + in_flight += wave_size; }; admit(); diff --git a/include/vllm/v1/engine/async_llm.h b/include/vllm/v1/engine/async_llm.h index 7019eaff0..344e96388 100644 --- a/include/vllm/v1/engine/async_llm.h +++ b/include/vllm/v1/engine/async_llm.h @@ -24,6 +24,7 @@ #include #include #include +#include #include #include "vllm/multimodal/inputs.h" // multimodal::MultiModalInputs (mm request) @@ -43,6 +44,38 @@ struct AsyncRequest { std::shared_ptr collector; }; +// Ordered-wave admission inputs. Separate types make it impossible for one +// item to carry both prompt forms. The batch API preserves these objects' +// order in both core publication and returned collectors. +struct AsyncStringRequestInput { + std::string request_id; + std::string prompt; + SamplingParams params; + int priority = 0; +}; + +struct AsyncTokensRequestInput { + std::string request_id; + std::vector prompt_token_ids; + SamplingParams params; + int priority = 0; +}; + +// The second admission check lives on the lock-owned publish route, rather +// than only at the start of input preparation. This injectable seam lets a +// deterministic test model shutdown winning after the outer fast check and +// prove that no registration or core publish callback can run afterward. +template +inline decltype(auto) PublishAsyncRequestWaveIfAlive( + Lockable& admission_mutex, AlivePredicate&& alive_predicate, + PublishCallback&& publish_callback) { + std::lock_guard lock(admission_mutex); + if (!std::forward(alive_predicate)()) { + throw EngineDeadError("request wave submitted to a stopped AsyncLLM"); + } + return std::forward(publish_callback)(); +} + class AsyncLLM { public: // Mirrors AsyncLLM.__init__ wiring, using caller-owned collaborators just as @@ -94,6 +127,16 @@ class AsyncLLM { std::vector prompt_token_ids, SamplingParams params, int priority = 0); + // Prepare and atomically publish one complete string or TokensPrompt wave. + // Every input/core Request and collector is built before any frontend state + // is registered. On duplicate, preparation/enqueue failure, or shutdown, + // this call publishes no core prefix and removes only state created by this + // call, without producing synthetic terminal outputs. + std::vector add_request_wave( + std::vector requests); + std::vector add_request_wave( + std::vector requests); + // add_request for a MULTIMODAL prompt (ROAD-V1-MM MM-SERVE-ENGINE). Strictly // ADDITIVE overload mirroring LLMEngine::add_request(MultiModalInputs): builds // the request from the placeholder-EXPANDED prompt ids + mm_features via @@ -151,6 +194,15 @@ class AsyncLLM { void shutdown(); private: + struct PreparedRequest { + EngineCoreRequest request; + std::optional prompt; + std::shared_ptr collector; + std::unique_ptr core_request; + }; + + std::vector PublishPreparedWave( + std::vector prepared); void RunOutputHandler(); InputProcessor& input_processor_; diff --git a/include/vllm/v1/engine/core_client.h b/include/vllm/v1/engine/core_client.h index 496e75d54..d8c0b4635 100644 --- a/include/vllm/v1/engine/core_client.h +++ b/include/vllm/v1/engine/core_client.h @@ -45,6 +45,7 @@ #include #include #include +#include #include #include "vllm/v1/engine/core_proc.h" // EngineCoreProc + queue/message types @@ -64,6 +65,16 @@ class EngineDeadError : public std::runtime_error { (detail.empty() ? std::string() : " [" + detail + "]")) {} }; +// The client-to-core wave route owns the queue's atomic batch primitive. Keep +// this seam injectable so focused tests can use a queue that deliberately has +// no per-item API: replacing the batch publish with an item loop then fails to +// compile instead of silently weakening wave visibility. +template +inline void PublishEngineCoreInputWaveAtomically(Queue& queue, + std::vector items) { + queue.put_many_nowait(std::move(items)); +} + // The in-process EngineCore client over the busy-loop queue split. // Collaborator lifetimes as EngineCore: scheduler/executor (and the optional // structured-output manager) are caller-owned and must outlive the client. @@ -95,6 +106,11 @@ class InprocClient { // engines_running flip (:887-888) is deferred. void add_request_async(std::unique_ptr request); + // Add one ordered frontend wave atomically. The input queue holds its mutex + // across the full append and wakes the engine thread once, so its first drain + // cannot observe only a prefix of this call. Empty waves are no-ops. + void add_requests_async(std::vector> requests); + // abort_requests (core_client.py:891-893): enqueue an ABORT; no-op for an // empty list or a dead engine. void abort_requests_async(const std::vector& request_ids); diff --git a/include/vllm/v1/engine/core_proc.h b/include/vllm/v1/engine/core_proc.h index 6e75ed05b..ea1da288e 100644 --- a/include/vllm/v1/engine/core_proc.h +++ b/include/vllm/v1/engine/core_proc.h @@ -86,7 +86,7 @@ enum class EngineShutdownState : int { // BlockingQueue: in-proc analog of Python's unbounded queue.Queue — the type // of EngineCoreProc.input_queue / output_queue (core.py:915-916). Blocking // get(), non-blocking put_nowait()/try_get(), thread-safe. -template +template class BlockingQueue { public: void put_nowait(T item) { @@ -97,6 +97,26 @@ class BlockingQueue { cv_.notify_one(); } + // Publish one ordered request wave while holding the queue mutex across the + // complete append. A consumer therefore observes either the previous queue + // or the previous queue followed by the whole batch, never a prefix. If an + // element move/allocation throws, remove only the suffix appended by this + // call before releasing the mutex. A successful non-empty wave emits one + // wakeup regardless of its size. + void put_many_nowait(std::vector items) { + if (items.empty()) return; + std::unique_lock lock(mutex_); + const std::size_t original_size = items_.size(); + try { + for (T& item : items) items_.push_back(std::move(item)); + } catch (...) { + while (items_.size() > original_size) items_.pop_back(); + throw; + } + lock.unlock(); + cv_.notify_one(); + } + // Blocking pop (queue.Queue.get(block=True)). T get() { std::unique_lock lock(mutex_); @@ -131,7 +151,7 @@ class BlockingQueue { private: mutable std::mutex mutex_; - std::condition_variable cv_; + ConditionVariable cv_; std::deque items_; }; diff --git a/include/vllm/v1/engine/output_processor.h b/include/vllm/v1/engine/output_processor.h index 53f29c40a..c9a280fe5 100644 --- a/include/vllm/v1/engine/output_processor.h +++ b/include/vllm/v1/engine/output_processor.h @@ -202,6 +202,9 @@ class OutputProcessor { return static_cast(request_states_.size()); } bool has_unfinished_requests() const { return !request_states_.empty(); } + bool has_request(const std::string& request_id) const { + return request_states_.find(request_id) != request_states_.end(); + } // add_request (:512): build + register a RequestState. `parent_req` is non-null // only for a child of an n>1 parallel-sampling request (SAMPLE-N); it is stored @@ -237,6 +240,11 @@ class OutputProcessor { const std::vector& request_ids, bool produce_final_output = false); + // Admission-transaction rollback: erase only the named frontend states, + // without allocating a return list or manufacturing terminal output. The + // caller serializes this with add/process/abort using its existing lock. + void rollback_requests(const std::vector& request_ids) noexcept; + // AsyncLLM teardown helper: abort every tracked internal request and return // the IDs that still need forwarding to EngineCore. The caller provides the // same external serialization as abort_requests/process_outputs. diff --git a/src/vllm/v1/engine/async_llm.cpp b/src/vllm/v1/engine/async_llm.cpp index c46d097ab..3b546ab0b 100644 --- a/src/vllm/v1/engine/async_llm.cpp +++ b/src/vllm/v1/engine/async_llm.cpp @@ -7,6 +7,7 @@ #include #include #include +#include #include #include @@ -123,6 +124,123 @@ AsyncRequest AsyncLLM::add_request(const std::string& request_id, return AsyncRequest{request.request_id, std::move(collector)}; } +std::vector AsyncLLM::add_request_wave( + std::vector requests) { + if (requests.empty()) return {}; + if (shutdown_started_.load() || errored_.load() || + engine_core_.engine_dead()) { + throw EngineDeadError("request wave submitted to a stopped AsyncLLM"); + } + + std::set request_ids; + for (const AsyncStringRequestInput& input : requests) { + if (!request_ids.insert(input.request_id).second) { + throw std::invalid_argument("duplicate request id in wave: " + + input.request_id); + } + } + + std::vector prepared; + prepared.reserve(requests.size()); + for (AsyncStringRequestInput& input : requests) { + EngineCoreRequest request = input_processor_.process_inputs( + input.request_id, input.prompt, std::move(input.params), + /*arrival_time=*/std::nullopt, input.priority); + auto collector = std::make_shared( + request.sampling_params.output_kind, request.request_id); + auto core_request = std::make_unique( + Request::FromEngineCoreRequest(request, block_hasher_)); + prepared.push_back(PreparedRequest{ + std::move(request), std::move(input.prompt), std::move(collector), + std::move(core_request)}); + } + return PublishPreparedWave(std::move(prepared)); +} + +std::vector AsyncLLM::add_request_wave( + std::vector requests) { + if (requests.empty()) return {}; + if (shutdown_started_.load() || errored_.load() || + engine_core_.engine_dead()) { + throw EngineDeadError("request wave submitted to a stopped AsyncLLM"); + } + + std::set request_ids; + for (const AsyncTokensRequestInput& input : requests) { + if (!request_ids.insert(input.request_id).second) { + throw std::invalid_argument("duplicate request id in wave: " + + input.request_id); + } + } + + std::vector prepared; + prepared.reserve(requests.size()); + for (AsyncTokensRequestInput& input : requests) { + EngineCoreRequest request = input_processor_.process_inputs_tokens( + input.request_id, std::move(input.prompt_token_ids), + std::move(input.params), /*arrival_time=*/std::nullopt, input.priority); + auto collector = std::make_shared( + request.sampling_params.output_kind, request.request_id); + auto core_request = std::make_unique( + Request::FromEngineCoreRequest(request, block_hasher_)); + prepared.push_back(PreparedRequest{ + std::move(request), std::nullopt, std::move(collector), + std::move(core_request)}); + } + return PublishPreparedWave(std::move(prepared)); +} + +std::vector AsyncLLM::PublishPreparedWave( + std::vector prepared) { + std::vector result; + std::vector rollback_ids; + std::vector> core_requests; + result.reserve(prepared.size()); + rollback_ids.reserve(prepared.size()); + core_requests.reserve(prepared.size()); + for (PreparedRequest& item : prepared) { + result.push_back(AsyncRequest{item.request.request_id, item.collector}); + rollback_ids.push_back(item.request.request_id); + core_requests.push_back(std::move(item.core_request)); + } + + return PublishAsyncRequestWaveIfAlive( + output_processor_mutex_, + [&]() { + return !shutdown_started_.load() && !errored_.load() && + !engine_core_.engine_dead(); + }, + [&]() -> std::vector { + // Reject every collision before creating the first new frontend state. + // This keeps a colliding pre-existing request outside the rollback set. + for (const PreparedRequest& item : prepared) { + if (output_processor_.has_request(item.request.request_id)) { + throw std::invalid_argument("duplicate live request id: " + + item.request.request_id); + } + } + + std::size_t rollback_count = 0; + try { + for (std::size_t i = 0; i < prepared.size(); ++i) { + // Include the current id in rollback before registration: + // add_request may have inserted one of its maps before a later + // allocation fails. + rollback_count = i + 1; + PreparedRequest& item = prepared[i]; + output_processor_.add_request(item.request, item.prompt, + /*request_index=*/0, item.collector); + } + engine_core_.add_requests_async(std::move(core_requests)); + } catch (...) { + rollback_ids.resize(rollback_count); + output_processor_.rollback_requests(rollback_ids); + throw; + } + return std::move(result); + }); +} + AsyncRequest AsyncLLM::add_request(const std::string& request_id, multimodal::MultiModalInputs mm_inputs, SamplingParams params, int priority) { diff --git a/src/vllm/v1/engine/core_client.cpp b/src/vllm/v1/engine/core_client.cpp index 95027679a..cc6075a3d 100644 --- a/src/vllm/v1/engine/core_client.cpp +++ b/src/vllm/v1/engine/core_client.cpp @@ -74,6 +74,23 @@ void InprocClient::add_request_async(std::unique_ptr request) { proc_.input_queue.put_nowait(std::move(item)); } +void InprocClient::add_requests_async( + std::vector> requests) { + if (requests.empty()) return; + + std::vector items; + items.reserve(requests.size()); + static const bool kLoopTrace = std::getenv("VT_LOOP_TRACE") != nullptr; + for (std::unique_ptr& request : requests) { + EngineCoreInputItem item; + item.type = EngineCoreRequestType::kAdd; + item.request = std::move(request); + if (kLoopTrace) item.enqueue_ts = MonotonicSeconds(); + items.push_back(std::move(item)); + } + PublishEngineCoreInputWaveAtomically(proc_.input_queue, std::move(items)); +} + void InprocClient::abort_requests_async( const std::vector& request_ids) { // core_client.py:891-893: "if request_ids and not resources.engine_dead". diff --git a/src/vllm/v1/engine/output_processor.cpp b/src/vllm/v1/engine/output_processor.cpp index f9b845632..e9c52c039 100644 --- a/src/vllm/v1/engine/output_processor.cpp +++ b/src/vllm/v1/engine/output_processor.cpp @@ -575,6 +575,28 @@ std::vector OutputProcessor::abort_requests( return request_ids_to_abort; } +void OutputProcessor::rollback_requests( + const std::vector& request_ids) noexcept { + for (const std::string& request_id : request_ids) { + auto request_it = request_states_.find(request_id); + if (request_it == request_states_.end()) continue; + + const RequestState& state = *request_it->second; + auto external_it = external_req_ids_.find(state.external_req_id); + if (external_it != external_req_ids_.end()) { + std::vector& internal_ids = external_it->second; + internal_ids.erase( + std::remove(internal_ids.begin(), internal_ids.end(), request_id), + internal_ids.end()); + if (internal_ids.empty()) external_req_ids_.erase(external_it); + } + if (state.parent_req != nullptr) { + parent_requests_.erase(state.parent_req->request_id()); + } + request_states_.erase(request_it); + } +} + std::vector OutputProcessor::abort_all_requests( bool produce_final_output) { std::vector request_ids; diff --git a/tests/examples/test_bench.cpp b/tests/examples/test_bench.cpp index 3b3a4c264..518505bd3 100644 --- a/tests/examples/test_bench.cpp +++ b/tests/examples/test_bench.cpp @@ -21,6 +21,7 @@ using vllm::bench::BenchConfig; using vllm::bench::BenchResult; using vllm::bench::DispatchBenchPromptAdmission; +using vllm::bench::DispatchBenchPromptWaveAdmission; using vllm::bench::PretokenizeBenchPromptsThenStartClock; using vllm::bench::RunBench; @@ -118,6 +119,35 @@ struct RecordingTokenizer { } }; +struct WaveOnlyEngine { + std::vector* events = nullptr; + int token_publishes = 0; + int string_publishes = 0; + std::vector published_ids; + + std::vector add_request_wave( + std::vector wave) { + ++token_publishes; + events->push_back("token-publish"); + published_ids.clear(); + for (const auto& request : wave) { + published_ids.push_back(request.request_id); + } + return published_ids; + } + + std::vector add_request_wave( + std::vector wave) { + ++string_publishes; + events->push_back("string-publish"); + published_ids.clear(); + for (const auto& request : wave) { + published_ids.push_back(request.request_id); + } + return published_ids; + } +}; + } // namespace TEST_CASE("bench: pretokenized admission dispatch is default-on with exact rollback") { @@ -144,6 +174,60 @@ TEST_CASE("bench: pretokenized admission dispatch is default-on with exact rollb CHECK(invalid.string_calls == 0); } +TEST_CASE("bench: wave admission owns one ordered engine publish") { + std::vector events; + WaveOnlyEngine engine; + engine.events = &events; + const std::vector result = DispatchBenchPromptWaveAdmission( + engine, /*env_value=*/nullptr, /*wave_size=*/3, + [&](size_t offset) { + events.push_back("arrival-" + std::to_string(offset)); + }, + [&]() { + events.push_back("token-build"); + std::vector wave; + for (int id = 4; id < 7; ++id) { + wave.push_back({std::to_string(id), {id}, {}, 0}); + } + return wave; + }, + [&]() { + events.push_back("string-build"); + return std::vector{ + {"wrong-arm", "prompt", {}, 0}}; + }); + CHECK(result == std::vector{"4", "5", "6"}); + CHECK(engine.token_publishes == 1); + CHECK(engine.string_publishes == 0); + CHECK(engine.published_ids == std::vector{"4", "5", "6"}); + CHECK(events == std::vector{"arrival-0", "arrival-1", + "arrival-2", "token-build", + "token-publish"}); + + events.clear(); + const std::vector rollback = DispatchBenchPromptWaveAdmission( + engine, /*env_value=*/"0", /*wave_size=*/2, + [&](size_t offset) { + events.push_back("arrival-" + std::to_string(offset)); + }, + [&]() { + events.push_back("token-build"); + return std::vector{ + {"wrong-arm", {99}, {}, 0}}; + }, + [&]() { + events.push_back("string-build"); + return std::vector{ + {"9", "first", {}, 0}, {"10", "second", {}, 0}}; + }); + CHECK(rollback == std::vector{"9", "10"}); + CHECK(engine.token_publishes == 1); + CHECK(engine.string_publishes == 1); + CHECK(engine.published_ids == std::vector{"9", "10"}); + CHECK(events == std::vector{"arrival-0", "arrival-1", + "string-build", "string-publish"}); +} + TEST_CASE("bench: report exposes the resolved pretokenized admission mode") { BenchResult result; result.pretokenized_admission = true; diff --git a/tests/vllm/v1/test_async_llm.cpp b/tests/vllm/v1/test_async_llm.cpp index da40df9e5..d0aca4184 100644 --- a/tests/vllm/v1/test_async_llm.cpp +++ b/tests/vllm/v1/test_async_llm.cpp @@ -56,6 +56,8 @@ using vllm::tok::MapBytesToUnicode; using vllm::tok::Tokenizer; using vllm::v1::AsyncLLM; using vllm::v1::AsyncRequest; +using vllm::v1::AsyncStringRequestInput; +using vllm::v1::AsyncTokensRequestInput; using vllm::v1::Executor; using vllm::v1::FullAttentionSpec; using vllm::v1::get_request_block_hasher; @@ -65,6 +67,7 @@ using vllm::v1::KVCacheConfig; using vllm::v1::ModelRunnerBase; using vllm::v1::ModelRunnerOutput; using vllm::v1::OutputProcessor; +using vllm::v1::PublishAsyncRequestWaveIfAlive; using vllm::v1::Scheduler; using vllm::v1::SchedulerOutput; using vllm::v1::sha256_cbor; @@ -137,6 +140,21 @@ class CerrRedirect { std::streambuf* previous_; }; +struct ObservableMutex { + void lock() { + ++lock_calls; + locked = true; + } + void unlock() { + ++unlock_calls; + locked = false; + } + + bool locked = false; + int lock_calls = 0; + int unlock_calls = 0; +}; + std::unique_ptr CreateScheduler() { SchedulerConfig cfg; cfg.max_num_seqs = 128; @@ -162,7 +180,8 @@ Tokenizer BuildFixture() { .string(); json doc; doc["version"] = "1.0"; - doc["added_tokens"] = json::array(); + doc["added_tokens"] = json::array( + {{{"id", 21}, {"content", ""}, {"special", true}}}); doc["normalizer"] = nullptr; doc["pre_tokenizer"] = { {"type", "Sequence"}, @@ -178,6 +197,17 @@ Tokenizer BuildFixture() { {"add_prefix_space", false}, {"trim_offsets", false}, {"use_regex", false}}})}}; + doc["post_processor"] = json::parse(R"json({ + "type": "TemplateProcessing", + "single": [ + {"SpecialToken": {"id": "", "type_id": 0}}, + {"Sequence": {"id": "A", "type_id": 0}} + ], + "pair": [], + "special_tokens": { + "": {"id": "", "ids": [21], "tokens": [""]} + } + })json"); json vocab = {{"h", 0}, {"e", 1}, {"l", 2}, {"o", 3}, {"w", 4}, {"r", 5}, {"d", 6}, {"Ġ", 7}, {"1", 8}, {"2", 9}, {"ll", 10}, {"he", 11}, @@ -241,6 +271,35 @@ void InitHash() { } // namespace +TEST_CASE("async_llm wave guard rejects shutdown after the outer check") { + bool alive = true; + bool published = false; + bool predicate_observed_lock = false; + ObservableMutex admission_mutex; + + // Model a submitter that passed the lock-free fast check, then lost the race + // to shutdown before entering the admission critical section. + REQUIRE(alive); + alive = false; + CHECK_THROWS_AS( + PublishAsyncRequestWaveIfAlive( + admission_mutex, + [&]() { + predicate_observed_lock = admission_mutex.locked; + return alive; + }, + [&]() { + published = true; + return 17; + }), + vllm::v1::EngineDeadError); + CHECK(predicate_observed_lock); + CHECK_FALSE(published); + CHECK_FALSE(admission_mutex.locked); + CHECK(admission_mutex.lock_calls == 1); + CHECK(admission_mutex.unlock_calls == 1); +} + TEST_CASE("async_llm test_load: concurrent requests all finish with unique ids") { InitHash(); Tokenizer tokenizer = BuildFixture(); @@ -270,6 +329,205 @@ TEST_CASE("async_llm test_load: concurrent requests all finish with unique ids") CHECK_FALSE(engine.has_unfinished_requests()); } +TEST_CASE( + "async_llm ordered waves preserve string/token inputs and one-item behavior") { + InitHash(); + Tokenizer tokenizer = BuildFixture(); + HfConfig config = MakeConfig(); + auto scheduler = CreateScheduler(); + RunnerStub runner; + Executor executor(runner); + InputProcessor input(tokenizer, config); + OutputProcessor output(&tokenizer); + AsyncLLM engine(input, *scheduler, executor, output, + get_request_block_hasher(16, sha256_cbor)); + + const std::vector prompts = {"hello", "hello world", "world"}; + CHECK(tokenizer.EncodeWithSpecialTokens(prompts.front()) != + tokenizer.Encode(prompts.front())); + + std::vector strings; + for (size_t i = 0; i < prompts.size(); ++i) { + strings.push_back({"string-" + std::to_string(i), prompts[i], + Params(4, RequestOutputKind::kDelta), 0}); + } + std::vector string_requests = + engine.add_request_wave(std::move(strings)); + REQUIRE(string_requests.size() == prompts.size()); + + std::vector> string_prompt_ids; + std::vector> string_output_ids; + for (size_t i = 0; i < string_requests.size(); ++i) { + CHECK(string_requests[i].request_id == "string-" + std::to_string(i)); + std::vector frames; + CHECK(Drain(engine, string_requests[i], &frames) == 4); + REQUIRE_FALSE(frames.empty()); + string_prompt_ids.push_back(frames.front().prompt_token_ids); + std::vector generated; + for (const RequestOutput& frame : frames) { + REQUIRE(frame.outputs.size() == 1); + generated.insert(generated.end(), frame.outputs[0].token_ids.begin(), + frame.outputs[0].token_ids.end()); + } + string_output_ids.push_back(std::move(generated)); + } + + std::vector tokens; + for (size_t i = 0; i < prompts.size(); ++i) { + tokens.push_back({"tokens-" + std::to_string(i), + tokenizer.EncodeWithSpecialTokens(prompts[i]), + Params(4, RequestOutputKind::kDelta), 0}); + } + std::vector token_requests = + engine.add_request_wave(std::move(tokens)); + REQUIRE(token_requests.size() == prompts.size()); + + std::vector> token_prompt_ids; + std::vector> token_output_ids; + for (size_t i = 0; i < token_requests.size(); ++i) { + CHECK(token_requests[i].request_id == "tokens-" + std::to_string(i)); + std::vector frames; + CHECK(Drain(engine, token_requests[i], &frames) == 4); + REQUIRE_FALSE(frames.empty()); + token_prompt_ids.push_back(frames.front().prompt_token_ids); + std::vector generated; + for (const RequestOutput& frame : frames) { + REQUIRE(frame.outputs.size() == 1); + generated.insert(generated.end(), frame.outputs[0].token_ids.begin(), + frame.outputs[0].token_ids.end()); + } + token_output_ids.push_back(std::move(generated)); + } + CHECK(token_prompt_ids == string_prompt_ids); + CHECK(token_output_ids == string_output_ids); + + std::vector singleton; + singleton.push_back({"singleton", tokenizer.EncodeWithSpecialTokens("hello"), + Params(3, RequestOutputKind::kDelta), 0}); + std::vector one = engine.add_request_wave(std::move(singleton)); + REQUIRE(one.size() == 1); + CHECK(one.front().request_id == "singleton"); + CHECK(Drain(engine, one.front()) == 3); + CHECK_FALSE(engine.has_unfinished_requests()); +} + +TEST_CASE("async_llm wave preparation failure publishes no request") { + InitHash(); + Tokenizer tokenizer = BuildFixture(); + HfConfig config = MakeConfig(); + auto scheduler = CreateScheduler(); + RunnerStub runner; + Executor executor(runner); + InputProcessor input(tokenizer, config); + OutputProcessor output(&tokenizer); + AsyncLLM engine(input, *scheduler, executor, output, + get_request_block_hasher(16, sha256_cbor)); + + SamplingParams invalid = Params(2, RequestOutputKind::kDelta); + invalid.max_tokens = 0; + std::vector wave; + wave.push_back( + {"prepared-before-failure", "hello", + Params(2, RequestOutputKind::kDelta), 0}); + wave.push_back({"invalid", "hello", std::move(invalid), 0}); + wave.push_back( + {"never-prepared", "hello", Params(2, RequestOutputKind::kDelta), 0}); + CHECK_THROWS_AS(engine.add_request_wave(std::move(wave)), std::runtime_error); + CHECK_FALSE(engine.has_unfinished_requests()); + + // The first id was prepared but never registered/published and is reusable. + AsyncRequest reused = engine.add_request( + "prepared-before-failure", "hello", + Params(2, RequestOutputKind::kDelta)); + CHECK(Drain(engine, reused) == 2); +} + +TEST_CASE("async_llm duplicate wave rollback preserves the existing request") { + InitHash(); + Tokenizer tokenizer = BuildFixture(); + HfConfig config = MakeConfig(); + auto scheduler = CreateScheduler(); + RunnerStub runner(std::chrono::milliseconds(1)); + Executor executor(runner); + InputProcessor input(tokenizer, config); + OutputProcessor output(&tokenizer); + AsyncLLM engine(input, *scheduler, executor, output, + get_request_block_hasher(16, sha256_cbor)); + + AsyncRequest existing = engine.add_request( + "existing", "hello", Params(100000, RequestOutputKind::kDelta)); + std::vector wave; + wave.push_back( + {"new-a", "hello", Params(2, RequestOutputKind::kDelta), 0}); + wave.push_back( + {"existing", "hello", Params(2, RequestOutputKind::kDelta), 0}); + wave.push_back( + {"new-b", "hello", Params(2, RequestOutputKind::kDelta), 0}); + CHECK_THROWS_AS(engine.add_request_wave(std::move(wave)), + std::invalid_argument); + CHECK(engine.get_num_unfinished_requests() == 1); + + // Rollback must not strand new-a, publish any failed-wave prefix, or abort + // the colliding pre-existing request. + std::vector retry; + retry.push_back( + {"new-a", "hello", Params(2, RequestOutputKind::kDelta), 0}); + std::vector admitted = + engine.add_request_wave(std::move(retry)); + REQUIRE(admitted.size() == 1); + CHECK(Drain(engine, admitted.front()) == 2); + engine.abort(existing.request_id); + RequestOutput terminal = engine.get_output(existing); + CHECK(terminal.finished); + CHECK_FALSE(engine.has_unfinished_requests()); +} + +TEST_CASE("async_llm concurrent shutdown accepts or rejects a complete wave") { + InitHash(); + Tokenizer tokenizer = BuildFixture(); + HfConfig config = MakeConfig(); + + for (int attempt = 0; attempt < 16; ++attempt) { + auto scheduler = CreateScheduler(); + RunnerStub runner(std::chrono::milliseconds(1)); + Executor executor(runner); + InputProcessor input(tokenizer, config); + OutputProcessor output(&tokenizer); + AsyncLLM engine(input, *scheduler, executor, output, + get_request_block_hasher(16, sha256_cbor)); + + std::optional> accepted; + std::exception_ptr submit_error; + std::thread submitter([&] { + try { + std::vector wave; + for (int i = 0; i < 3; ++i) { + wave.push_back({"shutdown-wave-" + std::to_string(i), "hello", + Params(1000, RequestOutputKind::kDelta), 0}); + } + accepted = engine.add_request_wave(std::move(wave)); + } catch (...) { + submit_error = std::current_exception(); + } + }); + + engine.shutdown(); + submitter.join(); + if (accepted.has_value()) { + REQUIRE(accepted->size() == 3); + for (const AsyncRequest& request : *accepted) { + RequestOutput terminal = request.collector->get(); + CHECK(terminal.finished); + } + } else { + REQUIRE(submit_error != nullptr); + CHECK_THROWS_AS(std::rethrow_exception(submit_error), + vllm::v1::EngineDeadError); + } + CHECK_FALSE(output.has_unfinished_requests()); + } +} + TEST_CASE("async_llm test_abort and test_multi_abort leave other requests healthy") { InitHash(); Tokenizer tokenizer = BuildFixture(); diff --git a/tests/vllm/v1/test_engine_core_proc.cpp b/tests/vllm/v1/test_engine_core_proc.cpp index 6133a88ce..8bd687480 100644 --- a/tests/vllm/v1/test_engine_core_proc.cpp +++ b/tests/vllm/v1/test_engine_core_proc.cpp @@ -30,12 +30,16 @@ // deferred (core.py:196-223 selection; ENG-ASYNC-SCHED W3). #include +#include +#include +#include #include #include #include #include #include #include +#include #include #include "vllm/config/scheduler.h" @@ -65,12 +69,14 @@ using vllm::v1::EngineShutdownState; using vllm::v1::Executor; using vllm::v1::FinishReason; using vllm::v1::FullAttentionSpec; +using vllm::v1::BlockingQueue; using vllm::v1::get_request_block_hasher; using vllm::v1::init_none_hash; using vllm::v1::InprocClient; using vllm::v1::KVCacheConfig; using vllm::v1::ModelRunnerBase; using vllm::v1::ModelRunnerOutput; +using vllm::v1::PublishEngineCoreInputWaveAtomically; using vllm::v1::Request; using vllm::v1::Scheduler; using vllm::v1::SchedulerOutput; @@ -79,6 +85,60 @@ using vt::DType; namespace { +struct MoveGate { + std::atomic armed{false}; + std::atomic entered{false}; + std::atomic release{false}; + std::atomic throw_on_move{false}; +}; + +struct GatedMoveItem { + GatedMoveItem() = default; + GatedMoveItem(int value, std::shared_ptr gate) + : value(value), gate(std::move(gate)) {} + GatedMoveItem(const GatedMoveItem&) = delete; + GatedMoveItem& operator=(const GatedMoveItem&) = delete; + + GatedMoveItem(GatedMoveItem&& other) + : value(other.value), gate(std::move(other.gate)) { + if (gate != nullptr && value == 2 && gate->armed.load()) { + gate->entered.store(true); + while (!gate->release.load()) std::this_thread::yield(); + if (gate->throw_on_move.load()) { + throw std::runtime_error("injected batch element move failure"); + } + } + } + + GatedMoveItem& operator=(GatedMoveItem&& other) noexcept { + value = other.value; + gate = std::move(other.gate); + return *this; + } + + int value = 0; + std::shared_ptr gate; +}; + +struct CountingConditionVariable { + void notify_one() { ++notifications; } + + template + void wait(Lock&, Predicate&&) {} + + static inline int notifications = 0; +}; + +struct PutManyOnlyQueue { + int publish_calls = 0; + std::vector published; + + void put_many_nowait(std::vector items) { + ++publish_calls; + published = std::move(items); + } +}; + // The canned token the stub "samples" for every scheduled request (same seam // as test_engine_core.cpp; upstream drives a real tiny model). constexpr int32_t kCannedToken = 42; @@ -181,6 +241,82 @@ void LoopUntilDone(InprocClient& client, const std::set& want, } // namespace +TEST_CASE("engine core wave route owns one atomic put-many publish") { + PutManyOnlyQueue queue; + PublishEngineCoreInputWaveAtomically(queue, std::vector{4, 5, 6}); + + CHECK(queue.publish_calls == 1); + CHECK(queue.published == std::vector{4, 5, 6}); +} + +TEST_CASE("BlockingQueue batch publish is all-at-once and ordered") { + BlockingQueue queue; + const auto gate = std::make_shared(); + std::vector batch; + batch.emplace_back(1, gate); + batch.emplace_back(2, gate); + batch.emplace_back(3, gate); + gate->armed.store(true); + + auto producer = std::async(std::launch::async, [&] { + queue.put_many_nowait(std::move(batch)); + }); + for (int i = 0; i < 100 && !gate->entered.load(); ++i) { + std::this_thread::sleep_for(std::chrono::milliseconds(1)); + } + REQUIRE(gate->entered.load()); + + // The first element has already been appended, but a consumer must not be + // able to acquire the queue mutex until the complete wave is present. + auto consumer = std::async(std::launch::async, [&] { + GatedMoveItem item; + const bool found = queue.try_get(item); + return std::make_pair(found, item.value); + }); + CHECK(consumer.wait_for(std::chrono::milliseconds(20)) == + std::future_status::timeout); + + gate->release.store(true); + producer.get(); + const auto first = consumer.get(); + CHECK(first == std::make_pair(true, 1)); + + GatedMoveItem item; + REQUIRE(queue.try_get(item)); + CHECK(item.value == 2); + REQUIRE(queue.try_get(item)); + CHECK(item.value == 3); + CHECK_FALSE(queue.try_get(item)); +} + +TEST_CASE("BlockingQueue batch publish rolls back an appended suffix on failure") { + BlockingQueue queue; + const auto gate = std::make_shared(); + queue.put_nowait(GatedMoveItem(9, gate)); + + std::vector batch; + batch.emplace_back(1, gate); + batch.emplace_back(2, gate); + batch.emplace_back(3, gate); + gate->armed.store(true); + gate->release.store(true); + gate->throw_on_move.store(true); + CHECK_THROWS_AS(queue.put_many_nowait(std::move(batch)), std::runtime_error); + + gate->armed.store(false); + GatedMoveItem item; + REQUIRE(queue.try_get(item)); + CHECK(item.value == 9); + CHECK_FALSE(queue.try_get(item)); +} + +TEST_CASE("BlockingQueue successful batch emits one notification") { + CountingConditionVariable::notifications = 0; + BlockingQueue queue; + queue.put_many_nowait(std::vector{1, 2, 3}); + CHECK(CountingConditionVariable::notifications == 1); +} + // --------------------------------------------------------------------------- // Normal request cycle (test_engine_core_client.py:548-562): add N requests, // pull outputs off the busy loop until all finish; every request produced From 70f01f3009311c2c978c153cf20d9ecdf6796313 Mon Sep 17 00:00:00 2001 From: Richard Palethorpe Date: Sun, 9 Aug 2026 22:13:15 +0100 Subject: [PATCH 09/17] spec(cuda): bound argmax scratch growth (#206) FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: Codex:gpt-5.6-sol [Codex] --- ...qwen35-argmax-scratch-growth-2026-08-09.md | 252 ++++++++++++++++++ 1 file changed, 252 insertions(+) create mode 100644 .agents/specs/sm120-qwen35-argmax-scratch-growth-2026-08-09.md diff --git a/.agents/specs/sm120-qwen35-argmax-scratch-growth-2026-08-09.md b/.agents/specs/sm120-qwen35-argmax-scratch-growth-2026-08-09.md new file mode 100644 index 000000000..cc8910d13 --- /dev/null +++ b/.agents/specs/sm120-qwen35-argmax-scratch-growth-2026-08-09.md @@ -0,0 +1,252 @@ +# RTX 5070 Ti Qwen3.5-4B argmax scratch-growth spike + +**Issue:** [#206 — RTX 5070 Ti: close Qwen3.5-4B TTFT, TPOT, and +VRAM gaps vs vLLM](https://github.com/mudler/vllm.cpp/issues/206) + +**Row:** `KERNEL-SAMPLING`, feeding `ROAD-V1-C2-LOCAL-BF16` + +**Parent campaign:** +[RTX 5070 Ti Qwen3.5-4B Pareto campaign](sm120-qwen35-pareto-2026-08-09.md) + +**Lifecycle:** `SPIKED`; no implementation or performance credit + +## Decision and objective + +The first implementation candidate is an opt-in replacement for the current +exact-size, process-global argmax scratch allocator. It will use a per-device, +per-stream owner, round capacity geometrically to the next power of two, +allocate the new value/index pair with `cudaMallocAsync` on the submitting +stream, and pass replaced blocks to the existing `RetireGraphScratch` helper. +It must make no synchronous allocation or free call after the first sufficient +capacity is available. A strict selector preserves the incumbent path exactly +when unset, `0`, or invalid. + +This spike is deliberately narrower than changing the two-pass argmax kernel. +The objective is to remove the measured host-blocking scratch growth storm, +improve Qwen3.5-4B TTFT without degrading throughput or TPOT, and retain exact +lowest-index argmax behavior and CUDA-graph pointer validity. It does not change +the default, the sampler ABI, scheduling, the normal device pool, or another +kernel family. + +## Exact attribution and source ground + +An `LD_PRELOAD` trace captured **506** total `cudaFree` calls. Exactly **30 hot +calls (15 pairs)** resolve through `src/vt/cuda/cuda_sample.cu:159-166` +`EnsureArgmaxScratch`, called from `GreedyArgmaxCuda` at `:193-217`; the other +476 have different teardown stacks and are excluded. The first free in each hot +pair blocks for about 191--200 ms while its immediately following sibling is +only microseconds. Independently, `VT_POOL_STATS` reports the normal pool +uncapped with a 99.14% hit rate. These facts reject the earlier normal-pool-cap +hypothesis and identify argmax scratch ownership as the next falsifiable lever; +they do not attribute the teardown calls to argmax. + +The local path is: + +- `src/vt/cuda/cuda_sample.cu:83-89` defines the exact reduction: higher value, + then lower global index on ties. +- `ArgmaxPartialKernel` at `:102-128` and `ArgmaxFinalKernel` at `:130-151` + implement the two-pass reduction, including the all-`-inf` sentinel rule. +- Process-global `g_argmax_val`, `g_argmax_idx`, and `g_argmax_cap` at `:153-157` + own the current scratch. `EnsureArgmaxScratch` frees both old blocks + synchronously, allocates exact requested sizes with `cudaMalloc`, and publishes + them at `:159-166`. +- `GreedyArgmaxCuda` computes + `blocks_per_row=min(ceil(vocab/256),256)`, requests + `rows*blocks_per_row` elements, and launches both kernels at `:193-217`. +- `VT_FAST_ARGMAX=0` at `:168-201` is a slow-kernel comparison, not an allocator + rollback, and must remain independent of this candidate. + +The pinned vLLM oracle at `5559679229bc` routes all-greedy sampling through +`${VLLM_SOURCE}/vllm/v1/sample/sampler.py:239-271`, where +`greedy_sample` is `logits.argmax(dim=-1).view(-1)`. The pinned Torch reduction +defines the observable tie rule in +`.venv-vllm-pin/lib/python3.12/site-packages/torch/include/ATen/native/SharedReduceOps.h:437-447,493-496`: +equal values select the lower index. `ATen/native/ReduceOps.h:19-30` owns the +dispatched argmax stub. The matched vLLM trace uses this production path and has +only 9 allocation and 22 free calls after its warm-up/graph-capture interval; +that count is a denominator observation, not a claim about unobserved Torch +allocator internals. + +Existing local parity gates are +`tests/vt/test_ops_sample.cpp:81-103` (basic and lowest-index tie), `:608-632` +(CUDA versus CPU random rows), and `:675-715` (151,936-wide cross-block tie, +high-index maximum, and all-`-inf`). The pinned sampler fixture construction is +at `${VLLM_SOURCE}/tests/v1/sample/test_sampler.py:128-166`. Commit +`5da39b0cf` introduced the fast two-pass argmax and its scratch; this work must +not restore the serialized slow scan. + +## Binding behavior and ownership contract + +1. Add a pure, CPU-testable selector such as + `ArgmaxGeometricScratchEnabled(const char*)`. Only the exact string `"1"` + selects the candidate. Unset, `"0"`, empty, `"true"`, `"01"`, and every + invalid value invoke only the unchanged incumbent allocator. The working + selector name is `VT_ARGMAX_GEOMETRIC_SCRATCH`; the implementer may choose a + clearer name before RED only if this spec is amended first. +2. Key candidate scratch by `(device index, native CUDA stream)`. Each owner has + one value pointer, one index pointer, an element capacity, and a submit mutex. + A stream never consumes another stream's partials, and a stream handle reused + on another device cannot alias the entry. The mutex covers growth decision, + publication, and both kernel submissions so another host thread cannot + replace or concurrently overwrite that owner's pair between launches. +3. For `need>capacity`, validate overflow and choose + `new_capacity=next_power_of_two(need)`. Allocate both candidate blocks with + `cudaMallocAsync` on that same stream before publishing either pointer or the + capacity. If the second allocation fails, enqueue cleanup of the first on the + same stream and leave the old pair and capacity unchanged before propagating + the error. +4. After both allocations succeed, atomically publish the pair and capacity, + then call `RetireGraphScratch` for both replaced non-null blocks. Never call + `cudaFree` or `cudaFreeAsync` on an old published block: a previously captured + graph may contain that address. Same/smaller requests allocate, free, and + retire nothing. +5. Query `cudaStreamIsCapturing` before a capacity miss. A hit is capture-safe + because its pointers are stable. A miss during capture fails before allocation, + retirement, publication, or kernel launch and says that scratch must be + warmed/reserved before capture. Growth outside capture may occur after an + older graph was recorded because retirement keeps every baked pointer valid. +6. `RetireGraphScratch` remains the process-lifetime graph owner already used by + other grow-only CUDA scratch. Do not add an unsafe queue-destruction free of + retired blocks. Per-stream map cleanup is outside this small spike unless an + independent test proves no graph can outlive it; owner cardinality and retained + bytes must nevertheless be reported. +7. `VT_FAST_ARGMAX=0` continues to select the existing slow kernel without + allocating candidate scratch. The new selector changes allocation ownership + only; it may not change kernel geometry, comparator, output type, sampling + metadata, or public ABI. + +For the bound benchmark (`vocab ~= 248,320`, maximum batched rows 2,048), +`blocks_per_row=256`, so the maximum request is 524,288 elements. The exact-size +final pair uses +`524288*(sizeof(float)+sizeof(int64_t)) = 6,291,456` bytes (6 MiB). The generic +next-power-of-two ratio gives a coarse active bound below 12 MiB; at this exact +power-of-two workload ceiling it is at most 6 MiB, and the sum of all preceding +geometric capacities is less than the active capacity. The acceptance gate uses +the deliberately looser **18 MiB total active-plus-retired** bound per benchmark +owner to cover transient allocation and allocator-accounting effects. Overflow, +owner count, and total retained bytes are explicit diagnostics outside timed +execution; multiple owners are reported separately rather than hidden inside +the per-owner bound. + +## Alternatives and discriminators + +| Strategy | Benefit | Rejection risk / discriminator | +|---|---|---| +| Exact-size synchronous grow/free (incumbent) | Exact rollback, bounded live allocation | Observed 30 hot `cudaFree` calls and repeated 191--200 ms host stalls | +| Power-of-two capacity only | Reduces growths from distinct batch high-water marks to logarithmic growth | Still performs a synchronizing `cudaFree` on every remaining growth; insufficient if any measured hot free remains | +| Power-of-two plus `RetireGraphScratch` and same-stream `cudaMallocAsync` (**first candidate**) | Removes hot synchronous free, preserves graph pointers, and bounds geometric retention | Retained bytes and stream-owner cardinality must meet the hard memory gates; async allocation must occur only outside capture | +| Stream-ordered async free | Bounds live bytes without a process-lifetime retire list | Rejected for a published pointer that a graph may replay; reconsider only with proven graph-lifetime ownership and destruction ordering | +| Explicit pre-serve reserve | Can make the entire serving interval allocation-free with the tightest known capacity | Requires a runner-to-sampler lifecycle surface and complete coverage of every caller; attempt only if the small candidate still grows in the measured interval | + +If the first candidate eliminates the attributed calls but misses the latency +gate, the next experiment is explicit pre-serve reserve using +`max_rows*min(ceil(vocab/256),256)`. Do not combine it silently with the first +measurement. If retained memory fails its bound, do not substitute async free +without first proving graph lifetime with a dedicated spec amendment. + +## RED-first implementation and focused tests + +A fresh implementer adds the smallest portable ownership/selector test before +product code and records the intended compile failure because the seam does not +exist. Prefer a small header-level state machine with injected allocate, retire, +capture-query, and submit callbacks; CUDA adapter code supplies the actual runtime +calls. The tests must prove: + +- strict selector truth table and exactly one incumbent/candidate callback; +- requests `1,3,5,9,17,32`, shrink, and regrow produce capacities + `1,4,8,16,32,32,32`, with no calls on hits; +- both new allocations precede publication, both previous pointers retire once, + and no free callback exists on successful hot growth; +- failure of either allocation preserves the previous pair/capacity, cleans only + a partially created candidate through the same-stream failure path, and does + not retire or launch; +- device/stream keys isolate owners and a concurrent submit cannot observe a + half-published pair or interleave between partial and final launches; +- a warmed capture hit allocates nothing; a capture miss fails before changing + state or launching; +- overflow is rejected and the geometric byte bound holds across every tested + growth sequence; +- unset, `0`, and invalid selectors execute byte-for-byte incumbent callback + identity, while `VT_FAST_ARGMAX=0` executes neither fast allocator; +- the existing CUDA exactness fixtures remain exact, plus a growth sequence over + several row counts preserves CPU token IDs, cross-block lowest-index ties, + unique maxima, and all-`-inf` index zero. + +The fresh reviewer mutates the immutable implementation in a scratch copy and +must kill at least these defects: permissive/inverted selector; invoking both +callbacks; exact rather than power-of-two growth; freeing instead of retiring; +publishing after only one successful allocation; dropping the device or stream +from the owner key; releasing the submit lock between kernels; allowing a +capture miss; leaking/double-cleaning a failed first allocation; wrong +`float`/`int64_t` byte sizes; overflow in rounding; changing `blocks_per_row`; +last-index tie selection; and an all-`-inf` sentinel that returns a nonzero index. +The tree is restored byte-for-byte after every mutation campaign. + +## Correctness, trace, memory, and timing gates + +Correctness precedes timing. On the exact immutable binary used for A/B: + +1. Run the focused pure selector/owner tests, complete sampling CPU tests, full + CUDA `test_ops_sample`, clean CUDA `-Werror`, and `compute-sanitizer memcheck` + on the growth/tie/capture fixture. Port the applicable pinned sampler + parameters rather than weakening local coverage. +2. Run Qwen3.5-4B BF16 with the parent campaign's cached model and ShareGPT + prompts: 128 requests, 128 output tokens, concurrency 32, + `max_num_batched_tokens=2048`, 1,280 KV blocks, greedy sampling, the accepted + production frontend, and every other accepted toggle held fixed. Prompt IDs, + every request's output IDs, output order, and their hashes must match across + all six arms before interpreting performance. +3. Use one binary and an idle, cgroup-contained RTX 5070 Ti. Run the + counterbalanced sequence + `INC-a -> GEO-a -> GEO-b -> INC-b -> INC-c -> GEO-c`, where the new selector + is the only difference. Record exact command, binary SHA-256, git revision, + model snapshot/hash, driver/runtime, clocks/power, ambient contention, and + raw outputs. +4. Record total and output throughput, mean/median/p95 TTFT, mean TPOT/ITL, + mean/p95 end-to-end latency, GPU peak/resident memory, host PSS, and load time + for every arm. Report each value and candidate/incumbent ratio, not only a + headline mean. +5. Trace both local arms with the same `nsys` command and full request interval. + The candidate passes attribution only if zero `cudaFree` calls resolve to + `EnsureArgmaxScratch`, zero synchronous frees occur on candidate growth, and + the remaining allocation/retire count agrees with the geometric high-water + sequence. Record CUDA API time and argmax kernel counts/durations. Re-run the + pinned production vLLM trace with the identical workload if its existing + trace cannot be matched to the same host state and profiler options. +6. Hard memory gates: the final active argmax pair is at most 6 MiB for the + declared `n<=2048`, `blocks_per_row<=256` workload, and accounted + active-plus-retired scratch remains below the conservative 18 MiB bound per + benchmark owner. Observed peak VRAM is no more than incumbent plus 18 MiB and + no more than the historical 13,053.3 MiB local peak. Host PSS and owner + cardinality may not grow unbounded across repeated queue create/destroy cycles. +7. Accept only if all three candidate legs are exact and stable, at least two of + three paired TTFT observations improve, aggregate mean TTFT improves by at + least 1%, and total/output throughput, TPOT/ITL, E2E, load time, peak VRAM, + and host PSS do not regress outside the campaign's calibrated noise band. + A trace improvement without end-to-end improvement is diagnostic evidence, + not a default flip. + +The operator reruns the full row gate after fresh implementation and mutation +review. Any later default change requires this spec's `## Outcome`, accepted raw +evidence, and the parent campaign records; this spike alone authorizes none. + +## Rollback, risks, and stop conditions + +Rollback is exact and same-binary: unset, `VT_ARGMAX_GEOMETRIC_SCRATCH=0`, or +an invalid value takes only the unchanged process-global exact-size +`cudaFree`/`cudaMalloc` path. The slow-kernel `VT_FAST_ARGMAX=0` remains a +separate diagnostic and receives no performance credit. + +Stop and return `NEEDS_DECISION` rather than widening scope if the change needs +a public API, scheduler behavior, sampler arithmetic, default flip, normal pool +change, or keyed roadmap/matrix edit. Reject the candidate on any token mismatch, +tie/all-`-inf` semantic drift, capture allocation, dangling graph pointer, +cross-stream race, unbounded owner growth, memory-gate failure, sanitizer error, +or Pareto regression. If all attributed frees disappear but TTFT does not move, +retain the trace as a falsified end-to-end hypothesis and profile the next +largest same-tool interval; never declare a ceiling. + +## Outcome + +Pending fresh RED-first implementation, mutation review, operator verification, +and same-binary GPU evidence. No default or performance claim is made. From 876033614f708f93c56ad66d9d539f9347eaf5d4 Mon Sep 17 00:00:00 2001 From: Richard Palethorpe Date: Sun, 9 Aug 2026 22:27:55 +0100 Subject: [PATCH 10/17] KERNEL-SAMPLING add opt-in geometric argmax scratch Keep argmax scratch per CUDA device and stream, grow it geometrically with stream-ordered allocation, and retire graph-visible blocks. Preserve the exact incumbent allocator unless the strict experimental selector is enabled. Add portable ownership, failure, capture, isolation, and serialization tests. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: Codex:gpt-5.6-sol [Codex] --- include/vt/cuda/argmax_scratch.h | 186 +++++++++++++++++++++++ scripts/env-doc-allowlist.txt | 1 + src/vt/cuda/cuda_sample.cu | 64 +++++++- tests/CMakeLists.txt | 1 + tests/vt/test_argmax_scratch.cpp | 247 +++++++++++++++++++++++++++++++ 5 files changed, 492 insertions(+), 7 deletions(-) create mode 100644 include/vt/cuda/argmax_scratch.h create mode 100644 tests/vt/test_argmax_scratch.cpp diff --git a/include/vt/cuda/argmax_scratch.h b/include/vt/cuda/argmax_scratch.h new file mode 100644 index 000000000..af88207b5 --- /dev/null +++ b/include/vt/cuda/argmax_scratch.h @@ -0,0 +1,186 @@ +// vllm.cpp original -- portable ownership policy for CUDA argmax scratch. +#ifndef VT_CUDA_ARGMAX_SCRATCH_H_ +#define VT_CUDA_ARGMAX_SCRATCH_H_ + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +namespace vt::cuda { + +inline bool ArgmaxGeometricScratchEnabled(const char* value) noexcept { + return value != nullptr && value[0] == '1' && value[1] == '\0'; +} + +template +decltype(auto) DispatchArgmaxScratch(const char* value, Incumbent&& incumbent, + Candidate&& candidate) { + if (ArgmaxGeometricScratchEnabled(value)) { + return std::forward(candidate)(); + } + return std::forward(incumbent)(); +} + +inline std::size_t RoundArgmaxScratchCapacity(std::size_t need) { + if (need <= 1) return need; + std::size_t capacity = 1; + while (capacity < need) { + if (capacity > std::numeric_limits::max() / 2) { + throw std::overflow_error("argmax geometric scratch capacity overflow"); + } + capacity *= 2; + } + return capacity; +} + +struct ArgmaxScratchKey { + int device = 0; + std::uintptr_t stream = 0; + + friend bool operator==(const ArgmaxScratchKey&, const ArgmaxScratchKey&) = default; +}; + +struct ArgmaxScratchView { + float* values = nullptr; + std::int64_t* indices = nullptr; + std::size_t capacity = 0; +}; + +struct ArgmaxScratchDiagnostics { + void* values = nullptr; + void* indices = nullptr; + std::size_t capacity = 0; + std::size_t active_bytes = 0; + std::size_t retired_bytes = 0; + std::size_t growths = 0; +}; + +class ArgmaxScratchOwners { + public: + // The caller's submit callback launches both reduction kernels. It runs while + // the per-(device, stream) owner mutex is held, preventing another host thread + // from replacing or overwriting the pair between those submissions. + template + void Submit(ArgmaxScratchKey key, std::size_t need, CaptureQuery&& is_capturing, + Allocate&& allocate, CleanupPartial&& cleanup_partial, Retire&& retire, + Launch&& submit) { + Owner* owner = nullptr; + bool capture_already_checked = false; + { + std::lock_guard lock(owners_mu_); + auto it = owners_.find(key); + if (it == owners_.end()) { + if (std::forward(is_capturing)()) ThrowCaptureMiss(); + capture_already_checked = true; + it = owners_.emplace(key, std::make_unique()).first; + } + owner = it->second.get(); + } + + std::lock_guard submit_lock(owner->submit_mu); + if (need > owner->capacity) { + if (!capture_already_checked && std::forward(is_capturing)()) { + ThrowCaptureMiss(); + } + const std::size_t new_capacity = RoundArgmaxScratchCapacity(need); + if (new_capacity > std::numeric_limits::max() / sizeof(std::int64_t)) { + throw std::overflow_error("argmax geometric scratch byte-size overflow"); + } + + void* new_values = nullptr; + void* new_indices = nullptr; + new_values = std::forward(allocate)(new_capacity * sizeof(float)); + if (new_values == nullptr) { + throw std::runtime_error("argmax geometric scratch value allocation returned null"); + } + try { + new_indices = std::forward(allocate)(new_capacity * sizeof(std::int64_t)); + if (new_indices == nullptr) { + throw std::runtime_error("argmax geometric scratch index allocation returned null"); + } + } catch (...) { + std::forward(cleanup_partial)(new_values); + throw; + } + + void* old_values = owner->values; + void* old_indices = owner->indices; + const std::size_t old_capacity = owner->capacity; + owner->values = new_values; + owner->indices = new_indices; + owner->capacity = new_capacity; + ++owner->growths; + if (old_values != nullptr) { + owner->retired_bytes += old_capacity * sizeof(float); + std::forward(retire)(old_values); + } + if (old_indices != nullptr) { + owner->retired_bytes += old_capacity * sizeof(std::int64_t); + std::forward(retire)(old_indices); + } + } + + std::forward(submit)(ArgmaxScratchView{static_cast(owner->values), + static_cast(owner->indices), + owner->capacity}); + } + + std::optional Diagnostics(ArgmaxScratchKey key) const { + Owner* owner = nullptr; + { + std::lock_guard lock(owners_mu_); + const auto it = owners_.find(key); + if (it == owners_.end()) return std::nullopt; + owner = it->second.get(); + } + std::lock_guard lock(owner->submit_mu); + return ArgmaxScratchDiagnostics{ + owner->values, owner->indices, + owner->capacity, owner->capacity * (sizeof(float) + sizeof(std::int64_t)), + owner->retired_bytes, owner->growths}; + } + + std::size_t OwnerCount() const { + std::lock_guard lock(owners_mu_); + return owners_.size(); + } + + private: + struct Owner { + mutable std::mutex submit_mu; + void* values = nullptr; + void* indices = nullptr; + std::size_t capacity = 0; + std::size_t retired_bytes = 0; + std::size_t growths = 0; + }; + + struct KeyHash { + std::size_t operator()(ArgmaxScratchKey key) const noexcept { + std::size_t hash = static_cast(key.stream); + hash ^= static_cast(static_cast(key.device)) + 0x9e3779b97f4a7c15ULL + + (hash << 6) + (hash >> 2); + return hash; + } + }; + + [[noreturn]] static void ThrowCaptureMiss() { + throw std::runtime_error( + "argmax geometric scratch capacity miss during CUDA graph capture; " + "warm or reserve scratch before capture"); + } + + mutable std::mutex owners_mu_; + std::unordered_map, KeyHash> owners_; +}; + +} // namespace vt::cuda + +#endif // VT_CUDA_ARGMAX_SCRATCH_H_ diff --git a/scripts/env-doc-allowlist.txt b/scripts/env-doc-allowlist.txt index 5abe69189..f742ba6e6 100644 --- a/scripts/env-doc-allowlist.txt +++ b/scripts/env-doc-allowlist.txt @@ -16,6 +16,7 @@ VT_ROCM_GEMM_COMPUTE VT_ROCM_GEMV VT_ROCM_HIPBLASLT VLLM_MM_TOWER_PROFILE +VT_ARGMAX_GEOMETRIC_SCRATCH VT_ARCH_TACTIC_STATS VT_ATTN_DECODE_GQA VT_ATTN_DECODE_OPT diff --git a/src/vt/cuda/cuda_sample.cu b/src/vt/cuda/cuda_sample.cu index 652988503..1aa2ed7a0 100644 --- a/src/vt/cuda/cuda_sample.cu +++ b/src/vt/cuda/cuda_sample.cu @@ -18,10 +18,13 @@ #include #include +#include #include #include #include "vt/backend.h" +#include "vt/cuda/argmax_scratch.h" +#include "vt/cuda/graph_safe_scratch.h" #include "vt/ops.h" namespace vt::cuda { @@ -155,6 +158,7 @@ __global__ void ArgmaxFinalKernel(int64_t* out, const float* part_val, const int float* g_argmax_val = nullptr; int64_t* g_argmax_idx = nullptr; size_t g_argmax_cap = 0; // capacity in elements +ArgmaxScratchOwners g_argmax_geometric_owners; void EnsureArgmaxScratch(size_t elems) { if (elems <= g_argmax_cap) return; @@ -190,6 +194,14 @@ bool FastArgmaxEnabled() { return on; } +const char* ArgmaxGeometricScratchSelector() { + static const char* selector = [] { + const char* value = std::getenv("VT_ARGMAX_GEOMETRIC_SCRATCH"); + return ArgmaxGeometricScratchEnabled(value) ? "1" : nullptr; + }(); + return selector; +} + void GreedyArgmaxCuda(Queue& q, Tensor& token_ids, const Tensor& logits) { const int64_t n = logits.shape[0], v = logits.shape[1]; if (n == 0 || v == 0) return; @@ -207,13 +219,51 @@ void GreedyArgmaxCuda(Queue& q, Tensor& token_ids, const Tensor& logits) { if (bpr > kBlock) bpr = kBlock; // pass 2 reduces bpr partials with kBlock threads if (bpr < 1) bpr = 1; - EnsureArgmaxScratch(static_cast(n) * bpr); - dim3 grid1(static_cast(bpr), static_cast(n)); - ArgmaxPartialKernel<<>>(g_argmax_val, g_argmax_idx, logits.Ptr(), v, - bpr); - ArgmaxFinalKernel<<(n), kBlock, 0, s>>>(token_ids.Ptr(), - g_argmax_val, g_argmax_idx, bpr); - Check(cudaGetLastError(), "greedy_argmax launch"); + DispatchArgmaxScratch( + ArgmaxGeometricScratchSelector(), + [&] { + // Exact rollback: retain the original process-global exact-size + // cudaFree/cudaMalloc allocator and launch sequence byte-for-byte. + EnsureArgmaxScratch(static_cast(n) * bpr); + dim3 grid1(static_cast(bpr), static_cast(n)); + ArgmaxPartialKernel<<>>(g_argmax_val, g_argmax_idx, + logits.Ptr(), v, bpr); + ArgmaxFinalKernel<<(n), kBlock, 0, s>>>( + token_ids.Ptr(), g_argmax_val, g_argmax_idx, bpr); + Check(cudaGetLastError(), "greedy_argmax launch"); + }, + [&] { + if (static_cast(n) > + std::numeric_limits::max() / static_cast(bpr)) { + throw std::overflow_error("argmax scratch element-count overflow"); + } + const size_t need = static_cast(n) * static_cast(bpr); + const ArgmaxScratchKey key{q.device.index, reinterpret_cast(s)}; + g_argmax_geometric_owners.Submit( + key, need, + [&] { + cudaStreamCaptureStatus status = cudaStreamCaptureStatusNone; + Check(cudaStreamIsCapturing(s, &status), "argmax scratch capture query"); + return status != cudaStreamCaptureStatusNone; + }, + [&](size_t bytes) -> void* { + void* p = nullptr; + Check(cudaMallocAsync(&p, bytes, s), "argmax geometric scratch allocation"); + return p; + }, + [&](void* p) { + Check(cudaFreeAsync(p, s), "argmax geometric scratch partial cleanup"); + }, + [](void* p) { RetireGraphScratch(p); }, + [&](ArgmaxScratchView scratch) { + dim3 grid1(static_cast(bpr), static_cast(n)); + ArgmaxPartialKernel<<>>(scratch.values, scratch.indices, + logits.Ptr(), v, bpr); + ArgmaxFinalKernel<<(n), kBlock, 0, s>>>( + token_ids.Ptr(), scratch.values, scratch.indices, bpr); + Check(cudaGetLastError(), "greedy_argmax launch"); + }); + }); } // --- compute_probs / compute_logprobs (block-per-row softmax) --------------- diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt index 3ed1ec824..455610d73 100644 --- a/tests/CMakeLists.txt +++ b/tests/CMakeLists.txt @@ -1047,6 +1047,7 @@ target_include_directories(test_ops_conv1d_depthwise PRIVATE ${CMAKE_SOURCE_DIR} vllm_cpp_add_test(test_ops_attn_relpos vt/test_ops_attn_relpos.cpp) target_include_directories(test_ops_attn_relpos PRIVATE ${CMAKE_SOURCE_DIR}/src) vllm_cpp_add_test(test_ops_paged_attn vt/test_ops_paged_attn.cpp) +vllm_cpp_add_test(test_argmax_scratch vt/test_argmax_scratch.cpp) vllm_cpp_add_test(test_ops_sample vt/test_ops_sample.cpp) vllm_cpp_add_test(test_ops_penalties vt/test_ops_penalties.cpp) diff --git a/tests/vt/test_argmax_scratch.cpp b/tests/vt/test_argmax_scratch.cpp new file mode 100644 index 000000000..2e59c7459 --- /dev/null +++ b/tests/vt/test_argmax_scratch.cpp @@ -0,0 +1,247 @@ +#include + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#include "vt/cuda/argmax_scratch.h" + +namespace { +using vt::cuda::ArgmaxScratchKey; +using vt::cuda::ArgmaxScratchOwners; +using vt::cuda::ArgmaxScratchView; + +void* FakePtr(std::uintptr_t value) { + return reinterpret_cast(value); +} + +TEST_CASE("argmax geometric scratch selector is strict and dispatches exactly once") { + const char* disabled[] = {nullptr, "", "0", "true", "01", "2", " 1"}; + for (const char* value : disabled) { + int incumbent = 0; + int candidate = 0; + const int selected = vt::cuda::DispatchArgmaxScratch( + value, + [&] { + ++incumbent; + return 17; + }, + [&] { + ++candidate; + return 29; + }); + CHECK(selected == 17); + CHECK(incumbent == 1); + CHECK(candidate == 0); + } + int incumbent = 0; + int candidate = 0; + CHECK(vt::cuda::DispatchArgmaxScratch( + "1", + [&] { + ++incumbent; + return 17; + }, + [&] { + ++candidate; + return 29; + }) == 29); + CHECK(incumbent == 0); + CHECK(candidate == 1); +} + +TEST_CASE("argmax geometric scratch grows transactionally by powers of two") { + ArgmaxScratchOwners owners; + const ArgmaxScratchKey key{3, 0x1234}; + std::uintptr_t next = 0x1000; + std::vector allocations; + std::vector cleanups; + std::vector retired; + std::vector seen_caps; + auto allocate = [&](std::size_t bytes) { + allocations.push_back(bytes); + next += 0x100; + return FakePtr(next); + }; + auto submit = [&](ArgmaxScratchView view) { seen_caps.push_back(view.capacity); }; + + for (std::size_t need : {1U, 3U, 5U, 9U, 17U, 32U, 7U, 31U}) { + owners.Submit( + key, need, [] { return false; }, allocate, [&](void* p) { cleanups.push_back(p); }, + [&](void* p) { retired.push_back(p); }, submit); + } + CHECK(seen_caps == std::vector{1, 4, 8, 16, 32, 32, 32, 32}); + CHECK(allocations == std::vector{4, 8, 16, 32, 32, 64, 64, 128, 128, 256}); + CHECK(cleanups.empty()); + CHECK(retired.size() == 8); // both blocks from each replaced capacity + const auto diag = owners.Diagnostics(key); + REQUIRE(diag.has_value()); + CHECK(diag->capacity == 32); + CHECK(diag->active_bytes == 32 * (sizeof(float) + sizeof(std::int64_t))); + CHECK(diag->retired_bytes == (1 + 4 + 8 + 16) * (sizeof(float) + sizeof(std::int64_t))); + CHECK(diag->active_bytes + diag->retired_bytes < 2 * diag->active_bytes); +} + +TEST_CASE("argmax geometric scratch preserves the published pair on allocation failure") { + ArgmaxScratchOwners owners; + const ArgmaxScratchKey key{0, 0x44}; + std::uintptr_t next = 0x2000; + std::vector cleanups; + std::vector retired; + int calls = 0; + auto ok_alloc = [&](std::size_t) { + next += 0x100; + return FakePtr(next); + }; + owners.Submit( + key, 3, [] { return false; }, ok_alloc, [&](void* p) { cleanups.push_back(p); }, + [&](void* p) { retired.push_back(p); }, [&](ArgmaxScratchView) { ++calls; }); + const auto before = *owners.Diagnostics(key); + + int allocation_number = 0; + CHECK_THROWS_AS(owners.Submit( + key, 5, [] { return false; }, + [&](std::size_t) -> void* { + ++allocation_number; + if (allocation_number == 2) + throw std::runtime_error("second allocation failed"); + next += 0x100; + return FakePtr(next); + }, + [&](void* p) { cleanups.push_back(p); }, + [&](void* p) { retired.push_back(p); }, [&](ArgmaxScratchView) { ++calls; }), + std::runtime_error); + const auto after = *owners.Diagnostics(key); + CHECK(after.values == before.values); + CHECK(after.indices == before.indices); + CHECK(after.capacity == before.capacity); + CHECK(cleanups == std::vector{FakePtr(0x2300)}); + CHECK(retired.empty()); + CHECK(calls == 1); +} + +TEST_CASE("argmax geometric scratch isolates keys and fails capture misses closed") { + ArgmaxScratchOwners owners; + std::uintptr_t next = 0x3000; + int allocations = 0; + int submits = 0; + auto allocate = [&](std::size_t) { + ++allocations; + next += 0x100; + return FakePtr(next); + }; + auto cleanup = [](void*) {}; + auto retire = [](void*) {}; + auto submit = [&](ArgmaxScratchView) { ++submits; }; + const ArgmaxScratchKey a{0, 0x99}; + const ArgmaxScratchKey b{1, 0x99}; + const ArgmaxScratchKey c{0, 0x100}; + owners.Submit(a, 4, [] { return false; }, allocate, cleanup, retire, submit); + owners.Submit(b, 4, [] { return false; }, allocate, cleanup, retire, submit); + owners.Submit(c, 4, [] { return false; }, allocate, cleanup, retire, submit); + CHECK(owners.OwnerCount() == 3); + CHECK(owners.Diagnostics(a)->values != owners.Diagnostics(b)->values); + CHECK(owners.Diagnostics(a)->values != owners.Diagnostics(c)->values); + + int capture_queries = 0; + owners.Submit( + a, 4, + [&] { + ++capture_queries; + return true; + }, + allocate, cleanup, retire, submit); + CHECK(capture_queries == 0); // a warmed hit is capture-safe + CHECK_THROWS_AS(owners.Submit( + a, 5, + [&] { + ++capture_queries; + return true; + }, + allocate, cleanup, retire, submit), + std::runtime_error); + CHECK(capture_queries == 1); + CHECK(allocations == 6); + CHECK(submits == 4); + CHECK(owners.Diagnostics(a)->capacity == 4); + + CHECK_THROWS_AS( + owners.Submit( + ArgmaxScratchKey{7, 0x777}, 1, [] { return true; }, allocate, cleanup, retire, submit), + std::runtime_error); + CHECK(owners.OwnerCount() == 3); // an absent-owner miss changes no state +} + +TEST_CASE("argmax geometric scratch submit lock spans the complete launch callback") { + ArgmaxScratchOwners owners; + const ArgmaxScratchKey key{0, 0x55}; + std::uintptr_t next = 0x4000; + auto allocate = [&](std::size_t) { + next += 0x100; + return FakePtr(next); + }; + owners.Submit( + key, 4, [] { return false; }, allocate, [](void*) {}, [](void*) {}, [](ArgmaxScratchView) {}); + + std::mutex mu; + std::condition_variable cv; + bool first_between_kernels = false; + bool release_first = false; + std::atomic second_entered{false}; + std::thread first([&] { + owners.Submit( + key, 4, [] { return false; }, allocate, [](void*) {}, [](void*) {}, + [&](ArgmaxScratchView) { + std::unique_lock lock(mu); + first_between_kernels = true; + cv.notify_all(); + cv.wait(lock, [&] { return release_first; }); + }); + }); + { + std::unique_lock lock(mu); + cv.wait(lock, [&] { return first_between_kernels; }); + } + std::thread second([&] { + owners.Submit( + key, 4, [] { return false; }, allocate, [](void*) {}, [](void*) {}, + [&](ArgmaxScratchView) { second_entered = true; }); + }); + std::this_thread::sleep_for(std::chrono::milliseconds(20)); + CHECK_FALSE(second_entered.load()); + { + std::lock_guard lock(mu); + release_first = true; + } + cv.notify_all(); + first.join(); + second.join(); + CHECK(second_entered.load()); +} + +TEST_CASE("argmax geometric scratch rejects capacity and byte overflow") { + CHECK(vt::cuda::RoundArgmaxScratchCapacity(17) == 32); + CHECK(vt::cuda::RoundArgmaxScratchCapacity(32) == 32); + CHECK_THROWS_AS(vt::cuda::RoundArgmaxScratchCapacity(std::numeric_limits::max()), + std::overflow_error); + ArgmaxScratchOwners owners; + int allocations = 0; + CHECK_THROWS_AS(owners.Submit( + ArgmaxScratchKey{0, 1}, std::numeric_limits::max() / 4 + 1, + [] { return false; }, + [&](std::size_t) { + ++allocations; + return FakePtr(1); + }, + [](void*) {}, [](void*) {}, [](ArgmaxScratchView) {}), + std::overflow_error); + CHECK(allocations == 0); +} +} // namespace From 89b930ce9bf1d9154f451e844dd66443dcf2f4d4 Mon Sep 17 00:00:00 2001 From: Richard Palethorpe Date: Sun, 9 Aug 2026 22:53:55 +0100 Subject: [PATCH 11/17] fix(KERNEL-SAMPLING): bind argmax scratch adapter Route cuda_sample through the capability-limited ownership seam and harden transactional first allocation plus reviewer mutation coverage. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: Codex:gpt-5.6-sol [Codex] --- include/vt/cuda/argmax_scratch.h | 55 +++++++++---- src/vt/cuda/cuda_sample.cu | 41 +++++----- tests/CMakeLists.txt | 1 + tests/vt/test_argmax_scratch.cpp | 132 ++++++++++++++++++++++++++++++- 4 files changed, 194 insertions(+), 35 deletions(-) diff --git a/include/vt/cuda/argmax_scratch.h b/include/vt/cuda/argmax_scratch.h index af88207b5..5f178a311 100644 --- a/include/vt/cuda/argmax_scratch.h +++ b/include/vt/cuda/argmax_scratch.h @@ -71,21 +71,8 @@ class ArgmaxScratchOwners { void Submit(ArgmaxScratchKey key, std::size_t need, CaptureQuery&& is_capturing, Allocate&& allocate, CleanupPartial&& cleanup_partial, Retire&& retire, Launch&& submit) { - Owner* owner = nullptr; - bool capture_already_checked = false; - { - std::lock_guard lock(owners_mu_); - auto it = owners_.find(key); - if (it == owners_.end()) { - if (std::forward(is_capturing)()) ThrowCaptureMiss(); - capture_already_checked = true; - it = owners_.emplace(key, std::make_unique()).first; - } - owner = it->second.get(); - } - - std::lock_guard submit_lock(owner->submit_mu); - if (need > owner->capacity) { + auto grow = [&](Owner* owner, bool capture_already_checked) { + if (need <= owner->capacity) return; if (!capture_already_checked && std::forward(is_capturing)()) { ThrowCaptureMiss(); } @@ -125,8 +112,31 @@ class ArgmaxScratchOwners { owner->retired_bytes += old_capacity * sizeof(std::int64_t); std::forward(retire)(old_indices); } + }; + + std::unique_lock owners_lock(owners_mu_); + auto it = owners_.find(key); + if (it == owners_.end()) { + if (std::forward(is_capturing)()) ThrowCaptureMiss(); + auto new_owner = std::make_unique(); + Owner* owner = new_owner.get(); + std::unique_lock submit_lock(owner->submit_mu); + // Keep the map lock until the first pair is complete. A failed first + // allocation therefore cannot publish an empty owner or race a second + // creator for the same key. + grow(owner, true); + owners_.emplace(key, std::move(new_owner)); + owners_lock.unlock(); + std::forward(submit)(ArgmaxScratchView{static_cast(owner->values), + static_cast(owner->indices), + owner->capacity}); + return; } + Owner* owner = it->second.get(); + owners_lock.unlock(); + std::lock_guard submit_lock(owner->submit_mu); + grow(owner, false); std::forward(submit)(ArgmaxScratchView{static_cast(owner->values), static_cast(owner->indices), owner->capacity}); @@ -181,6 +191,21 @@ class ArgmaxScratchOwners { std::unordered_map, KeyHash> owners_; }; +// This is the single production routing seam between cuda_sample and the pure +// ownership state machine. Runtime is intentionally capability-limited: the +// old published blocks can only be retired, while only a never-published +// partial allocation can use stream-ordered cleanup. +template +void SubmitCudaArgmaxScratch(ArgmaxScratchOwners& owners, const QueueLike& queue, Stream stream, + std::size_t need, Runtime& runtime, Launch&& launch) { + const ArgmaxScratchKey key{queue.device.index, reinterpret_cast(stream)}; + owners.Submit( + key, need, [&] { return runtime.IsCapturing(stream); }, + [&](std::size_t bytes) { return runtime.Allocate(bytes, stream); }, + [&](void* p) { runtime.CleanupPartial(p, stream); }, [&](void* p) { runtime.Retire(p); }, + std::forward(launch)); +} + } // namespace vt::cuda #endif // VT_CUDA_ARGMAX_SCRATCH_H_ diff --git a/src/vt/cuda/cuda_sample.cu b/src/vt/cuda/cuda_sample.cu index 1aa2ed7a0..86685e169 100644 --- a/src/vt/cuda/cuda_sample.cu +++ b/src/vt/cuda/cuda_sample.cu @@ -202,6 +202,26 @@ const char* ArgmaxGeometricScratchSelector() { return selector; } +struct CudaArgmaxScratchRuntime { + bool IsCapturing(cudaStream_t stream) const { + cudaStreamCaptureStatus status = cudaStreamCaptureStatusNone; + Check(cudaStreamIsCapturing(stream, &status), "argmax scratch capture query"); + return status != cudaStreamCaptureStatusNone; + } + + void* Allocate(size_t bytes, cudaStream_t stream) const { + void* p = nullptr; + Check(cudaMallocAsync(&p, bytes, stream), "argmax geometric scratch allocation"); + return p; + } + + void CleanupPartial(void* p, cudaStream_t stream) const { + Check(cudaFreeAsync(p, stream), "argmax geometric scratch partial cleanup"); + } + + void Retire(void* p) const { RetireGraphScratch(p); } +}; + void GreedyArgmaxCuda(Queue& q, Tensor& token_ids, const Tensor& logits) { const int64_t n = logits.shape[0], v = logits.shape[1]; if (n == 0 || v == 0) return; @@ -238,24 +258,9 @@ void GreedyArgmaxCuda(Queue& q, Tensor& token_ids, const Tensor& logits) { throw std::overflow_error("argmax scratch element-count overflow"); } const size_t need = static_cast(n) * static_cast(bpr); - const ArgmaxScratchKey key{q.device.index, reinterpret_cast(s)}; - g_argmax_geometric_owners.Submit( - key, need, - [&] { - cudaStreamCaptureStatus status = cudaStreamCaptureStatusNone; - Check(cudaStreamIsCapturing(s, &status), "argmax scratch capture query"); - return status != cudaStreamCaptureStatusNone; - }, - [&](size_t bytes) -> void* { - void* p = nullptr; - Check(cudaMallocAsync(&p, bytes, s), "argmax geometric scratch allocation"); - return p; - }, - [&](void* p) { - Check(cudaFreeAsync(p, s), "argmax geometric scratch partial cleanup"); - }, - [](void* p) { RetireGraphScratch(p); }, - [&](ArgmaxScratchView scratch) { + CudaArgmaxScratchRuntime runtime; + SubmitCudaArgmaxScratch( + g_argmax_geometric_owners, q, s, need, runtime, [&](ArgmaxScratchView scratch) { dim3 grid1(static_cast(bpr), static_cast(n)); ArgmaxPartialKernel<<>>(scratch.values, scratch.indices, logits.Ptr(), v, bpr); diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt index 455610d73..f02c9de74 100644 --- a/tests/CMakeLists.txt +++ b/tests/CMakeLists.txt @@ -1048,6 +1048,7 @@ vllm_cpp_add_test(test_ops_attn_relpos vt/test_ops_attn_relpos.cpp) target_include_directories(test_ops_attn_relpos PRIVATE ${CMAKE_SOURCE_DIR}/src) vllm_cpp_add_test(test_ops_paged_attn vt/test_ops_paged_attn.cpp) vllm_cpp_add_test(test_argmax_scratch vt/test_argmax_scratch.cpp) +target_compile_definitions(test_argmax_scratch PRIVATE VLLM_CPP_SOURCE_DIR="${CMAKE_SOURCE_DIR}") vllm_cpp_add_test(test_ops_sample vt/test_ops_sample.cpp) vllm_cpp_add_test(test_ops_penalties vt/test_ops_penalties.cpp) diff --git a/tests/vt/test_argmax_scratch.cpp b/tests/vt/test_argmax_scratch.cpp index 2e59c7459..52e8967e2 100644 --- a/tests/vt/test_argmax_scratch.cpp +++ b/tests/vt/test_argmax_scratch.cpp @@ -5,9 +5,11 @@ #include #include #include +#include #include #include #include +#include #include #include @@ -23,7 +25,7 @@ void* FakePtr(std::uintptr_t value) { } TEST_CASE("argmax geometric scratch selector is strict and dispatches exactly once") { - const char* disabled[] = {nullptr, "", "0", "true", "01", "2", " 1"}; + const char* disabled[] = {nullptr, "", "0", "true", "01", "2", " 1", "10", "1x", "1 "}; for (const char* value : disabled) { int incumbent = 0; int candidate = 0; @@ -80,7 +82,9 @@ TEST_CASE("argmax geometric scratch grows transactionally by powers of two") { CHECK(seen_caps == std::vector{1, 4, 8, 16, 32, 32, 32, 32}); CHECK(allocations == std::vector{4, 8, 16, 32, 32, 64, 64, 128, 128, 256}); CHECK(cleanups.empty()); - CHECK(retired.size() == 8); // both blocks from each replaced capacity + CHECK(retired == std::vector{FakePtr(0x1100), FakePtr(0x1200), FakePtr(0x1300), + FakePtr(0x1400), FakePtr(0x1500), FakePtr(0x1600), + FakePtr(0x1700), FakePtr(0x1800)}); const auto diag = owners.Diagnostics(key); REQUIRE(diag.has_value()); CHECK(diag->capacity == 32); @@ -127,6 +131,130 @@ TEST_CASE("argmax geometric scratch preserves the published pair on allocation f CHECK(calls == 1); } +TEST_CASE("argmax geometric scratch rejects a null first allocation transactionally") { + ArgmaxScratchOwners owners; + const ArgmaxScratchKey absent{4, 0x404}; + int launches = 0; + int retirements = 0; + CHECK_THROWS_AS(owners.Submit( + absent, 3, [] { return false; }, [](std::size_t) -> void* { return nullptr; }, + [](void*) { FAIL("null first allocation must not need cleanup"); }, + [&](void*) { ++retirements; }, [&](ArgmaxScratchView) { ++launches; }), + std::runtime_error); + CHECK(owners.OwnerCount() == 0); + CHECK_FALSE(owners.Diagnostics(absent).has_value()); + CHECK(retirements == 0); + CHECK(launches == 0); + + const ArgmaxScratchKey existing{4, 0x405}; + std::uintptr_t next = 0x5000; + owners.Submit( + existing, 3, [] { return false; }, + [&](std::size_t) { + next += 0x100; + return FakePtr(next); + }, + [](void*) {}, [&](void*) { ++retirements; }, [&](ArgmaxScratchView) { ++launches; }); + const auto before = *owners.Diagnostics(existing); + CHECK_THROWS_AS( + owners.Submit( + existing, 5, [] { return false; }, [](std::size_t) -> void* { return nullptr; }, + [](void*) { FAIL("null first allocation must not need cleanup"); }, + [&](void*) { ++retirements; }, [&](ArgmaxScratchView) { ++launches; }), + std::runtime_error); + const auto after = *owners.Diagnostics(existing); + CHECK(after.values == before.values); + CHECK(after.indices == before.indices); + CHECK(after.capacity == before.capacity); + CHECK(after.growths == before.growths); + CHECK(retirements == 0); + CHECK(launches == 1); +} + +TEST_CASE("CUDA argmax adapter binds device stream capture allocation and retirement") { + struct FakeQueue { + struct { + int index; + } device; + } queue{{7}}; + struct CapabilityLimitedRuntime { + bool capturing = false; + std::uintptr_t next = 0x6000; + std::vector capture_streams; + std::vector> allocations; + std::vector> partial_cleanups; + std::vector retirements; + + bool IsCapturing(void* stream) { + capture_streams.push_back(reinterpret_cast(stream)); + return capturing; + } + void* Allocate(std::size_t bytes, void* stream) { + allocations.emplace_back(bytes, reinterpret_cast(stream)); + next += 0x100; + return FakePtr(next); + } + void CleanupPartial(void* p, void* stream) { + partial_cleanups.emplace_back(p, reinterpret_cast(stream)); + } + void Retire(void* p) { retirements.push_back(p); } + } runtime; + + ArgmaxScratchOwners owners; + void* const stream = FakePtr(0x777); + std::vector launches; + vt::cuda::SubmitCudaArgmaxScratch(owners, queue, stream, 3, runtime, + [&](ArgmaxScratchView view) { launches.push_back(view); }); + vt::cuda::SubmitCudaArgmaxScratch(owners, queue, stream, 5, runtime, + [&](ArgmaxScratchView view) { launches.push_back(view); }); + CHECK(owners.OwnerCount() == 1); + CHECK(owners.Diagnostics(ArgmaxScratchKey{7, 0x777})->capacity == 8); + CHECK_FALSE(owners.Diagnostics(ArgmaxScratchKey{0, 0x777}).has_value()); + CHECK_FALSE(owners.Diagnostics(ArgmaxScratchKey{7, 0}).has_value()); + CHECK(runtime.capture_streams == std::vector{0x777, 0x777}); + CHECK(runtime.allocations == std::vector>{ + {16, 0x777}, {32, 0x777}, {32, 0x777}, {64, 0x777}}); + CHECK(runtime.partial_cleanups.empty()); + CHECK(runtime.retirements == std::vector{FakePtr(0x6100), FakePtr(0x6200)}); + REQUIRE(launches.size() == 2); + CHECK(launches[0].capacity == 4); + CHECK(launches[1].capacity == 8); + + CapabilityLimitedRuntime captured; + captured.capturing = true; + CHECK_THROWS_AS(vt::cuda::SubmitCudaArgmaxScratch( + owners, queue, FakePtr(0x778), 1, captured, + [](ArgmaxScratchView) { FAIL("capture miss must not launch"); }), + std::runtime_error); + CHECK(captured.capture_streams == std::vector{0x778}); + CHECK(captured.allocations.empty()); + CHECK(captured.retirements.empty()); + CHECK_FALSE(owners.Diagnostics(ArgmaxScratchKey{7, 0x778}).has_value()); +} + +TEST_CASE("CUDA argmax call site binds the capability adapter to graph-safe runtime operations") { + std::ifstream input(std::string(VLLM_CPP_SOURCE_DIR) + "/src/vt/cuda/cuda_sample.cu"); + REQUIRE(input.good()); + const std::string source((std::istreambuf_iterator(input)), + std::istreambuf_iterator()); + const auto runtime_begin = source.find("struct CudaArgmaxScratchRuntime {"); + const auto runtime_end = source.find("\n};", runtime_begin); + REQUIRE(runtime_begin != std::string::npos); + REQUIRE(runtime_end != std::string::npos); + const std::string runtime = source.substr(runtime_begin, runtime_end - runtime_begin); + CHECK(runtime.find("cudaStreamIsCapturing(stream, &status)") != std::string::npos); + CHECK(runtime.find("return status != cudaStreamCaptureStatusNone") != std::string::npos); + CHECK(runtime.find("cudaMallocAsync(&p, bytes, stream)") != std::string::npos); + CHECK(runtime.find("cudaFreeAsync(p, stream)") != std::string::npos); + CHECK(runtime.find("void Retire(void* p) const { RetireGraphScratch(p); }") != + std::string::npos); + + const auto route = source.find("SubmitCudaArgmaxScratch(", runtime_end); + REQUIRE(route != std::string::npos); + const std::string invocation = source.substr(route, 180); + CHECK(invocation.find("g_argmax_geometric_owners, q, s, need, runtime") != std::string::npos); +} + TEST_CASE("argmax geometric scratch isolates keys and fails capture misses closed") { ArgmaxScratchOwners owners; std::uintptr_t next = 0x3000; From 98878c67f50bb22bff421da9013b9c035c0bec85 Mon Sep 17 00:00:00 2001 From: Richard Palethorpe Date: Sun, 9 Aug 2026 23:18:09 +0100 Subject: [PATCH 12/17] perf(KERNEL-SSM-MAMBA): add BF16 vector writeback discriminator FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: Codex:gpt-5.6-sol [Codex] --- scripts/env-doc-allowlist.txt | 1 + src/vt/cuda/cuda_gdn.cu | 141 +++++++++++++++++++++++------ src/vt/cuda/gdn_decode_fused.h | 49 ++++++++++ tests/vt/test_gdn_decode_fused.cpp | 122 +++++++++++++++++++++++++ tests/vt/test_ops_gdn.cpp | 64 +++++++++++-- 5 files changed, 337 insertions(+), 40 deletions(-) diff --git a/scripts/env-doc-allowlist.txt b/scripts/env-doc-allowlist.txt index f742ba6e6..58ce81cb9 100644 --- a/scripts/env-doc-allowlist.txt +++ b/scripts/env-doc-allowlist.txt @@ -76,6 +76,7 @@ VT_GDN_BF16 VT_GDN_CHUNKED VT_GDN_CHUNKO_TRITON VT_GDN_CHUNK_VEC +VT_GDN_DECODE_BF16_VECSTORE VT_GDN_DECODE_NW VT_GDN_DELTAH_TRITON VT_GDN_DELTAH_VEC diff --git a/src/vt/cuda/cuda_gdn.cu b/src/vt/cuda/cuda_gdn.cu index 0ad4eff66..872eb7093 100644 --- a/src/vt/cuda/cuda_gdn.cu +++ b/src/vt/cuda/cuda_gdn.cu @@ -270,6 +270,27 @@ __device__ inline void Store(__nv_bfloat16* p, int64_t i, float v) { p[i] = __float2bfloat16(v); // round-to-nearest-even, same as host F32ToBF16 } +// Pack exactly eight scalar BF16 RN conversions into one aligned 16-byte +// transaction. The aggregate is alignment-safe; no potentially unaligned +// scalar object is type-punned. +__device__ inline void StoreBf16Vec8(__nv_bfloat16* p, float v0, float v1, + float v2, float v3, float v4, float v5, + float v6, float v7) { + static_assert(sizeof(__nv_bfloat16) * 8 == sizeof(int4)); + static_assert(alignof(int4) == 16); + int4 packed; + auto* elements = reinterpret_cast<__nv_bfloat16*>(&packed); + elements[0] = __float2bfloat16(v0); + elements[1] = __float2bfloat16(v1); + elements[2] = __float2bfloat16(v2); + elements[3] = __float2bfloat16(v3); + elements[4] = __float2bfloat16(v4); + elements[5] = __float2bfloat16(v5); + elements[6] = __float2bfloat16(v6); + elements[7] = __float2bfloat16(v7); + *reinterpret_cast(p) = packed; +} + template __device__ inline float RoundToStorage(float v); template <> @@ -2778,13 +2799,17 @@ void GdnPackedDecodeKernelCuda(Queue& q, Tensor& out, // a power of two dividing 32 and the launcher only uses NW>1 when BV==32, so a // block is always a whole number of warps) — the xor butterfly stays in-warp. template + bool REGSTATE, bool BF16_VECSTORE = false> __global__ void GdnDecodeFusedKernel(Tout* out, const Tin* q, const Tin* k, const Tin* v, const float* g, const float* beta, TState* state, const int32_t* state_idx, int64_t hk_n, int64_t dk, int64_t hv_n, int64_t dv, int64_t bv, float scale) { static_assert(!REGSTATE || (SWIZZLED && NW == 8), "register state requires the production swizzled NW8 layout"); + static_assert(!BF16_VECSTORE || + (REGSTATE && SWIZZLED && NW == 8 && + std::is_same::value), + "BF16 vector writeback requires BF16 REGSTATE BV16+SWIZZLE+NW8"); const int64_t i_v = blockIdx.x; // value-dim tile const int64_t i_nh = blockIdx.y; // fused (sequence, v-head) const int64_t i_n = i_nh / hv_n; // sequence == decode token index @@ -2900,19 +2925,35 @@ __global__ void GdnDecodeFusedKernel(Tout* out, const Tin* q, const Tin* k, cons #pragma unroll for (int off = 1; off < NW; off <<= 1) o += __shfl_xor_sync(0xffffffffu, o, off); if (vrow < dv && wk == 0) Store(out, (i_n * hv_n + hv) * dv + vrow, o); - if constexpr (REGSTATE) { + if constexpr (REGSTATE && !BF16_VECSTORE) { for (int64_t j = 0; j < 16; ++j) { r[GdnDecodeRegisterSharedColumn(wk, j)] = rr[j]; } } - __syncthreads(); + if constexpr (BF16_VECSTORE) { + // Eligibility proves Dv=Dk=128 and BV16: every lane owns one complete, + // valid 16-element logical slice. Each destination begins at a 32-byte + // boundary, so both 16-byte stores are naturally aligned. + static_assert(sizeof(TState) == 2); +#pragma unroll + for (int pack = 0; pack < 2; ++pack) { + const int64_t column = + GdnDecodeBf16PackLogicalColumn(wk, pack, 0); + const int base = pack * 8; + StoreBf16Vec8(s_head + static_cast(vi) * 128 + column, + rr[base], rr[base + 1], rr[base + 2], rr[base + 3], + rr[base + 4], rr[base + 5], rr[base + 6], rr[base + 7]); + } + } else { + __syncthreads(); - // Coalesced write-back of the updated slice from f32 registers to the - // configured fp16/bf16/fp32 temporal cache. - for (int64_t e = tid; e < tile; e += blockDim.x) { - const int64_t c = e % dk; - const int64_t sc = SWIZZLED ? (c % 16) * 8 + c / 16 : c; - Store(s_head, e, sbh[(e / dk) * sdk + sc]); + // Coalesced write-back of the updated slice from f32 registers to the + // configured fp16/bf16/fp32 temporal cache. + for (int64_t e = tid; e < tile; e += blockDim.x) { + const int64_t c = e % dk; + const int64_t sc = SWIZZLED ? (c % 16) * 8 + c / 16 : c; + Store(s_head, e, sbh[(e / dk) * sdk + sc]); + } } } @@ -2920,7 +2961,8 @@ template void LaunchGdnDecodeFusedNW(cudaStream_t s, Tensor& out, const Tensor& q_in, const Tensor& k, const Tensor& v, const Tensor& g, const Tensor& beta, Tensor& state, const int32_t* state_idx, int64_t n, const GdnArgs& args, - const GdnDecodeLaunchContract& contract) { + const GdnDecodeLaunchContract& contract, + bool bf16_vecstore) { const int64_t hk_n = q_in.shape[1], dk = q_in.shape[2]; const int64_t hv_n = v.shape[1], dv = v.shape[2]; const int64_t bv = contract.value_tile; @@ -2940,13 +2982,38 @@ void LaunchGdnDecodeFusedNW(cudaStream_t s, Tensor& out, const Tensor& q_in, con args.scale); }, [&](const GdnDecodeLaunchContract&) { - GdnDecodeFusedKernel - <<(contract.block_threads), - contract.shared_bytes, s>>>( - out.Ptr(), q_in.Ptr(), k.Ptr(), - v.Ptr(), g.Ptr(), beta.Ptr(), - state.Ptr(), state_idx, hk_n, dk, hv_n, dv, bv, - args.scale); + if constexpr (std::is_same::value) { + DispatchGdnDecodeBf16Vecstore( + bf16_vecstore, + [&] { + GdnDecodeFusedKernel + <<(contract.block_threads), + contract.shared_bytes, s>>>( + out.Ptr(), q_in.Ptr(), k.Ptr(), + v.Ptr(), g.Ptr(), beta.Ptr(), + state.Ptr(), state_idx, hk_n, dk, hv_n, dv, + bv, args.scale); + }, + [&] { + GdnDecodeFusedKernel + <<(contract.block_threads), + contract.shared_bytes, s>>>( + out.Ptr(), q_in.Ptr(), k.Ptr(), + v.Ptr(), g.Ptr(), beta.Ptr(), + state.Ptr(), state_idx, hk_n, dk, hv_n, dv, + bv, args.scale); + }); + } else { + GdnDecodeFusedKernel + <<(contract.block_threads), + contract.shared_bytes, s>>>( + out.Ptr(), q_in.Ptr(), k.Ptr(), + v.Ptr(), g.Ptr(), beta.Ptr(), + state.Ptr(), state_idx, hk_n, dk, hv_n, dv, bv, + args.scale); + } }); } else { GdnDecodeFusedKernel @@ -2971,23 +3038,28 @@ template void LaunchGdnDecodeFused(cudaStream_t s, Tensor& out, const Tensor& q_in, const Tensor& k, const Tensor& v, const Tensor& g, const Tensor& beta, Tensor& state, const int32_t* state_idx, int64_t n, const GdnArgs& args, - const GdnDecodeLaunchContract& contract) { + const GdnDecodeLaunchContract& contract, + bool bf16_vecstore) { switch (contract.lanes_per_row) { case 2: LaunchGdnDecodeFusedNW(s, out, q_in, k, v, g, beta, state, state_idx, - n, args, contract); + n, args, contract, + bf16_vecstore); break; case 4: LaunchGdnDecodeFusedNW(s, out, q_in, k, v, g, beta, state, state_idx, - n, args, contract); + n, args, contract, + bf16_vecstore); break; case 8: LaunchGdnDecodeFusedNW(s, out, q_in, k, v, g, beta, state, state_idx, - n, args, contract); + n, args, contract, + bf16_vecstore); break; default: LaunchGdnDecodeFusedNW(s, out, q_in, k, v, g, beta, state, state_idx, - n, args, contract); + n, args, contract, + bf16_vecstore); break; } } @@ -2998,16 +3070,19 @@ template void LaunchGdnDecodeFusedS(cudaStream_t s, Tensor& out, const Tensor& q_in, const Tensor& k, const Tensor& v, const Tensor& g, const Tensor& beta, Tensor& state, const int32_t* state_idx, int64_t n, const GdnArgs& args, - const GdnDecodeLaunchContract& contract) { + const GdnDecodeLaunchContract& contract, + bool bf16_vecstore) { if (state.dtype == DType::kBF16) LaunchGdnDecodeFused(s, out, q_in, k, v, g, beta, state, state_idx, - n, args, contract); + n, args, contract, + bf16_vecstore); else if (state.dtype == DType::kF16) LaunchGdnDecodeFused(s, out, q_in, k, v, g, beta, - state, state_idx, n, args, contract); + state, state_idx, n, args, contract, + bf16_vecstore); else LaunchGdnDecodeFused(s, out, q_in, k, v, g, beta, state, state_idx, n, args, - contract); + contract, bf16_vecstore); } // Decode dispatch. state_idx == nullptr: compact [n,Hv,Dv,Dk] state (row==i_n). @@ -3042,20 +3117,26 @@ void GdnDecodeFusedCuda(Queue& q, Tensor& out, const Tensor& q_in, const Tensor& args, "gdn_decode"); return; } + const bool bf16_vecstore = GdnDecodeBf16VecstoreEligible( + std::getenv("VT_GDN_DECODE_BF16_VECSTORE"), contract, dv, dk, nw, + state.dtype == DType::kBF16, true); if (q_in.dtype == DType::kF32) { if (out.dtype == DType::kF32) LaunchGdnDecodeFusedS(s, out, q_in, k, v, g, beta, state, state_idx, n, - args, contract); + args, contract, bf16_vecstore); else LaunchGdnDecodeFusedS(s, out, q_in, k, v, g, beta, state, state_idx, - n, args, contract); + n, args, contract, + bf16_vecstore); } else { if (out.dtype == DType::kF32) LaunchGdnDecodeFusedS<__nv_bfloat16, float>(s, out, q_in, k, v, g, beta, state, state_idx, - n, args, contract); + n, args, contract, + bf16_vecstore); else LaunchGdnDecodeFusedS<__nv_bfloat16, __nv_bfloat16>( - s, out, q_in, k, v, g, beta, state, state_idx, n, args, contract); + s, out, q_in, k, v, g, beta, state, state_idx, n, args, contract, + bf16_vecstore); } }; DispatchGdnDecodeValueTile(std::getenv("VT_GDN_DECODE_BV"), diff --git a/src/vt/cuda/gdn_decode_fused.h b/src/vt/cuda/gdn_decode_fused.h index c98d5a6c5..1fa93e54d 100644 --- a/src/vt/cuda/gdn_decode_fused.h +++ b/src/vt/cuda/gdn_decode_fused.h @@ -35,6 +35,13 @@ inline constexpr bool GdnDecodeRegstateFlagIsOn(const char* env_value) { return env_value != nullptr && env_value[0] == '1' && env_value[1] == '\0'; } +// Direct BF16 vector writeback is a strict, independently selectable +// discriminator. It is never implied by REGSTATE and every non-exact spelling +// leaves the incumbent shared-memory writeback intact. +inline constexpr bool GdnDecodeBf16VecstoreFlagIsOn(const char* env_value) { + return env_value != nullptr && env_value[0] == '1' && env_value[1] == '\0'; +} + inline constexpr int64_t GdnDecodeStateStride(bool swizzled, int64_t dk, int lanes_per_row) { return dk + (swizzled ? lanes_per_row : 1); @@ -54,6 +61,25 @@ inline constexpr int64_t GdnDecodeRegisterLogicalColumn(int lane, int slot) { return static_cast(lane) * 16 + slot; } +#if defined(__CUDACC__) +__host__ __device__ +#endif +inline constexpr int64_t GdnDecodeBf16PackLogicalColumn(int lane, int pack, + int element) { + return static_cast(lane) * 16 + pack * 8 + element; +} + +#if defined(__CUDACC__) +__host__ __device__ +#endif +inline constexpr int64_t GdnDecodeBf16PackByteOffset(int lane, int pack) { + return GdnDecodeBf16PackLogicalColumn(lane, pack, 0) * 2; +} + +static_assert(GdnDecodeBf16PackByteOffset(0, 0) == 0); +static_assert(GdnDecodeBf16PackByteOffset(0, 1) == 16); +static_assert(GdnDecodeBf16PackByteOffset(7, 1) == 240); + #if defined(__CUDACC__) __host__ __device__ #endif @@ -73,6 +99,21 @@ struct GdnDecodeLaunchContract { bool regstate; }; +// The vector specialization is intentionally narrower than REGSTATE. Keeping +// the dtype and regular-decode predicates explicit here makes host dispatch and +// portable mutation tests agree on every gate. +inline constexpr bool GdnDecodeBf16VecstoreEligible( + const char* env_value, const GdnDecodeLaunchContract& contract, int64_t dv, + int64_t dk, int requested_nw, bool state_is_bf16, + bool regular_fused_decode) { + return GdnDecodeBf16VecstoreFlagIsOn(env_value) && contract.regstate && + contract.swizzled && + contract.selected_tile == GdnDecodeValueTile::kBv16 && + contract.value_tile == 16 && contract.lanes_per_row == 8 && + contract.block_threads == 128 && dv == 128 && dk == 128 && + requested_nw == 8 && state_is_bf16 && regular_fused_decode; +} + inline constexpr GdnDecodeLaunchContract GdnDecodeLaunchContractFor( const char* bv_env_value, const char* swizzle_env_value, const char* regstate_env_value, int64_t dv, int64_t dk, @@ -124,6 +165,14 @@ inline decltype(auto) DispatchGdnDecodeStateStorage( return shared_launch(contract); } +template +inline decltype(auto) DispatchGdnDecodeBf16Vecstore( + bool vector_eligible, IncumbentLaunch&& incumbent_launch, + VectorLaunch&& vector_launch) { + if (vector_eligible) return vector_launch(); + return incumbent_launch(); +} + // Shared callback dispatcher used by production and portable tests. The // callbacks receive fully resolved geometry and exactly one is invoked. template diff --git a/tests/vt/test_gdn_decode_fused.cpp b/tests/vt/test_gdn_decode_fused.cpp index 468618701..8ef25ba90 100644 --- a/tests/vt/test_gdn_decode_fused.cpp +++ b/tests/vt/test_gdn_decode_fused.cpp @@ -2,6 +2,8 @@ // value tile; the CUDA recurrence remains in cuda_gdn.cu. #include +#include +#include #include #include "vt/cuda/gdn_decode_fused.h" @@ -12,6 +14,11 @@ using vt::cuda::GdnDecodeLaunchContractFor; using vt::cuda::GdnDecodeRegisterLogicalColumn; using vt::cuda::GdnDecodeRegisterSharedColumn; using vt::cuda::GdnDecodeRegstateFlagIsOn; +using vt::cuda::DispatchGdnDecodeBf16Vecstore; +using vt::cuda::GdnDecodeBf16PackByteOffset; +using vt::cuda::GdnDecodeBf16PackLogicalColumn; +using vt::cuda::GdnDecodeBf16VecstoreEligible; +using vt::cuda::GdnDecodeBf16VecstoreFlagIsOn; using vt::cuda::GdnDecodeSharedColumn; using vt::cuda::GdnDecodeStateStride; using vt::cuda::GdnDecodeValueTile; @@ -279,3 +286,118 @@ TEST_CASE("fused GDN decode register slots map bijectively to logical and shared CHECK(shared_seen[c]); } } + +namespace { + +uint16_t ReferenceBf16Rn(float value) { + uint32_t bits = 0; + static_assert(sizeof(bits) == sizeof(value)); + std::memcpy(&bits, &value, sizeof(bits)); + const uint32_t exponent = bits & 0x7f800000u; + const uint32_t mantissa = bits & 0x007fffffu; + if (exponent == 0x7f800000u && mantissa != 0) { + return static_cast((bits >> 16) | 0x0040u); + } + const uint32_t rounding_bias = 0x7fffu + ((bits >> 16) & 1u); + return static_cast((bits + rounding_bias) >> 16); +} + +} // namespace + +TEST_CASE("VT_GDN_DECODE_BF16_VECSTORE selects only exact one") { + CHECK(GdnDecodeBf16VecstoreFlagIsOn("1")); + constexpr const char* invalid[] = { + nullptr, "", "0", "01", "10", "1 ", " 1", "+1", "1x", "true", "yes"}; + for (const char* value : invalid) { + CAPTURE(value == nullptr ? "" : value); + CHECK_FALSE(GdnDecodeBf16VecstoreFlagIsOn(value)); + } +} + +TEST_CASE("GDN BF16 vector writeback requires every production predicate") { + const auto production = + GdnDecodeLaunchContractFor("16", "1", "1", 128, 128, 8); + auto eligible = [&](const char* vecstore, const auto& contract, int64_t dv, + int64_t dk, int nw, bool state_is_bf16, + bool regular_fused_decode) { + return GdnDecodeBf16VecstoreEligible( + vecstore, contract, dv, dk, nw, state_is_bf16, + regular_fused_decode); + }; + CHECK(eligible("1", production, 128, 128, 8, true, true)); + CHECK_FALSE(eligible(nullptr, production, 128, 128, 8, true, true)); + CHECK_FALSE(eligible("1x", production, 128, 128, 8, true, true)); + CHECK_FALSE(eligible("1", GdnDecodeLaunchContractFor("32", "1", "1", 128, 128, 8), + 128, 128, 8, true, true)); + CHECK_FALSE(eligible("1", GdnDecodeLaunchContractFor("16", "0", "1", 128, 128, 8), + 128, 128, 8, true, true)); + CHECK_FALSE(eligible("1", GdnDecodeLaunchContractFor("16", "1", "0", 128, 128, 8), + 128, 128, 8, true, true)); + CHECK_FALSE(eligible("1", production, 127, 128, 8, true, true)); + CHECK_FALSE(eligible("1", production, 128, 127, 8, true, true)); + CHECK_FALSE(eligible("1", production, 128, 128, 4, true, true)); + CHECK_FALSE(eligible("1", production, 128, 128, 8, false, true)); + CHECK_FALSE(eligible("1", production, 128, 128, 8, true, false)); +} + +TEST_CASE("GDN BF16 vector writeback dispatches exactly one callback") { + auto select = [](bool vector_eligible) { + int incumbent_calls = 0; + int vector_calls = 0; + const int selected = DispatchGdnDecodeBf16Vecstore( + vector_eligible, + [&] { + ++incumbent_calls; + return 0; + }, + [&] { + ++vector_calls; + return 1; + }); + CHECK(incumbent_calls + vector_calls == 1); + return selected; + }; + CHECK(select(false) == 0); + CHECK(select(true) == 1); +} + +TEST_CASE("GDN BF16 vector packs cover the aligned 128-column state row") { + bool seen[128] = {}; + for (int lane = 0; lane < 8; ++lane) { + const int64_t first = GdnDecodeBf16PackLogicalColumn(lane, 0, 0); + const int64_t second = GdnDecodeBf16PackLogicalColumn(lane, 1, 0); + CHECK(GdnDecodeBf16PackByteOffset(lane, 0) % 16 == 0); + CHECK(GdnDecodeBf16PackByteOffset(lane, 1) % 16 == 0); + CHECK(second == first + 8); + for (int pack = 0; pack < 2; ++pack) { + for (int element = 0; element < 8; ++element) { + const int64_t column = + GdnDecodeBf16PackLogicalColumn(lane, pack, element); + REQUIRE(column >= 0); + REQUIRE(column < 128); + CHECK_FALSE(seen[column]); + seen[column] = true; + } + } + } + for (bool visited : seen) CHECK(visited); + + const std::array values = { + 0.0f, -0.0f, 1.0f, -2.0f, 0.333251953125f, 65504.0f, + 1.00390625f, 1.01171875f, -1.00390625f, -1.01171875f, + 1.0e-20f, -1.0e-20f, 3.14159265f, -2.71828183f, 42.5f, -99.75f}; + std::array packed{}; + for (int pack = 0; pack < 2; ++pack) { + for (int element = 0; element < 8; ++element) { + const uint16_t bits = ReferenceBf16Rn(values[pack * 8 + element]); + std::memcpy(packed.data() + pack * 16 + element * 2, &bits, + sizeof(bits)); + } + } + for (int element = 0; element < 16; ++element) { + uint16_t packed_bits = 0; + std::memcpy(&packed_bits, packed.data() + element * 2, + sizeof(packed_bits)); + CHECK(packed_bits == ReferenceBf16Rn(values[element])); + } +} diff --git a/tests/vt/test_ops_gdn.cpp b/tests/vt/test_ops_gdn.cpp index f624887c8..0b842ae93 100644 --- a/tests/vt/test_ops_gdn.cpp +++ b/tests/vt/test_ops_gdn.cpp @@ -1869,6 +1869,7 @@ struct GdnDecodeTestEnv { const char* swizzle; const char* regstate; const char* nw; + const char* bf16_vecstore = nullptr; }; void SetOrUnsetEnv(const char* name, const char* value) { @@ -1914,6 +1915,7 @@ void RunGdnDecodeExactCase(int64_t dv, DType io_dtype, DType state_dtype, SetOrUnsetEnv("VT_GDN_DECODE_BV", env.bv); SetOrUnsetEnv("VT_GDN_DECODE_SWIZZLE", env.swizzle); SetOrUnsetEnv("VT_GDN_DECODE_REGSTATE", env.regstate); + SetOrUnsetEnv("VT_GDN_DECODE_BF16_VECSTORE", env.bf16_vecstore); DeviceTensor dq(gpu, gq.q, io_dtype, {n, hk, dk}, q.data()); DeviceTensor dkt(gpu, gq.q, io_dtype, {n, hk, dk}, k.data()); DeviceTensor dvt(gpu, gq.q, io_dtype, {n, hv, dv}, v.data()); @@ -1944,6 +1946,7 @@ void RunGdnDecodeExactCase(int64_t dv, DType io_dtype, DType state_dtype, run(incumbent, out_incumbent, state_incumbent); run(candidate, out_candidate, state_candidate); unsetenv("VT_GDN_DECODE_REGSTATE"); + unsetenv("VT_GDN_DECODE_BF16_VECSTORE"); unsetenv("VT_GDN_DECODE_SWIZZLE"); unsetenv("VT_GDN_DECODE_BV"); unsetenv("VT_GDN_DECODE_NW"); @@ -3538,6 +3541,28 @@ TEST_CASE("CUDA fused GDN decode register-state schedule is byte-exact") { nw4_shared, nw4_reg, 6440); } +TEST_CASE("CUDA fused GDN decode BF16 vector writeback is byte-exact") { + if (!HasCuda()) { + MESSAGE("no CUDA backend registered; skipping"); + return; + } + constexpr GdnDecodeTestEnv incumbent{"16", "1", "1", "8", "0"}; + constexpr GdnDecodeTestEnv vectorized{"16", "1", "1", "8", "1"}; + RunGdnDecodeExactCase(128, DType::kBF16, DType::kBF16, false, + incumbent, vectorized, 6450); + RunGdnDecodeExactCase(128, DType::kBF16, DType::kBF16, true, + incumbent, vectorized, 6460); + + // The selector is ineligible for independent fp16/fp32 cache dtypes and a + // partial value dimension; those configurations must remain incumbent. + RunGdnDecodeExactCase(128, DType::kBF16, DType::kF16, true, + incumbent, vectorized, 6470); + RunGdnDecodeExactCase(128, DType::kF32, DType::kF32, true, + incumbent, vectorized, 6480); + RunGdnDecodeExactCase(37, DType::kBF16, DType::kBF16, true, + incumbent, vectorized, 6490); +} + TEST_CASE("CUDA public GdnDecode binds decode layout to exact production geometry") { if (!HasCuda()) { MESSAGE("no CUDA backend registered; skipping"); @@ -3557,8 +3582,10 @@ TEST_CASE("CUDA public GdnDecode binds decode layout to exact production geometr DType::kBF16); const auto g = RandomF32(static_cast(n * hv), 6433, -0.5f, -0.01f); const auto beta = RandomF32(static_cast(n * hv), 6434, 0.1f, 0.9f); - const auto state = RandomF32(static_cast(n * hv * dv * dk), 6435, - -0.25f, 0.25f); + const auto state = + Pack(RandomF32(static_cast(n * hv * dv * dk), 6435, -0.25f, + 0.25f), + DType::kBF16); Backend& gpu = vt::GetBackend(DeviceType::kCUDA); QueueGuard gq(gpu); @@ -3569,7 +3596,7 @@ TEST_CASE("CUDA public GdnDecode binds decode layout to exact production geometr DeviceTensor dvt(gpu, gq.q, DType::kBF16, {n, hv, dv}, v.data()); DeviceTensor dg(gpu, gq.q, DType::kF32, {n, hv}, g.data()); DeviceTensor dbeta(gpu, gq.q, DType::kF32, {n, hv}, beta.data()); - DeviceTensor dst(gpu, gq.q, DType::kF32, {n, hv, dv, dk}, state.data()); + DeviceTensor dst(gpu, gq.q, DType::kBF16, {n, hv, dv, dk}, state.data()); DeviceTensor dout(gpu, gq.q, DType::kBF16, {n, hv, dv}); gpu.Synchronize(gq.q); @@ -3577,33 +3604,40 @@ TEST_CASE("CUDA public GdnDecode binds decode layout to exact production geometr const char* bv_env; const char* swizzle_env; const char* regstate_env; + const char* vecstore_env; bool register_state; + bool vector_writeback; unsigned grid_x; unsigned block_x; unsigned shared_bytes; }; constexpr ExpectedGeometry cases[] = { - {nullptr, "1", "1", false, 4, 256, 17536}, - {"16", "0", "1", false, 8, 128, 9280}, - {"16", "1", "0", false, 8, 128, 9728}, - {"16", "1", "1", true, 8, 128, 9728}, - {"16", "1", "1x", false, 8, 128, 9728}, - {"16", "1x", "1", false, 8, 128, 9280}, - {"32", "1", "1", false, 4, 256, 17536}, + {nullptr, "1", "1", "1", false, false, 4, 256, 17536}, + {"16", "0", "1", "1", false, false, 8, 128, 9280}, + {"16", "1", "0", "1", false, false, 8, 128, 9728}, + {"16", "1", "1", "0", true, false, 8, 128, 9728}, + {"16", "1", "1", "1", true, true, 8, 128, 9728}, + {"16", "1", "1", "1x", true, false, 8, 128, 9728}, + {"16", "1", "1x", "1", false, false, 8, 128, 9728}, + {"16", "1x", "1", "1", false, false, 8, 128, 9280}, + {"32", "1", "1", "1", false, false, 4, 256, 17536}, }; const cudaStream_t stream = static_cast(gq.q.handle); setenv("VT_GDN_DECODE_NW", "8", 1); void* shared_production_function = nullptr; + void* register_production_function = nullptr; for (const auto& expected : cases) { CAPTURE(expected.bv_env == nullptr ? "" : expected.bv_env); CAPTURE(expected.swizzle_env); CAPTURE(expected.regstate_env); + CAPTURE(expected.vecstore_env); if (expected.bv_env == nullptr) unsetenv("VT_GDN_DECODE_BV"); else setenv("VT_GDN_DECODE_BV", expected.bv_env, 1); setenv("VT_GDN_DECODE_SWIZZLE", expected.swizzle_env, 1); setenv("VT_GDN_DECODE_REGSTATE", expected.regstate_env, 1); + setenv("VT_GDN_DECODE_BF16_VECSTORE", expected.vecstore_env, 1); CudaGraphGuard captured; REQUIRE(cudaStreamBeginCapture(stream, cudaStreamCaptureModeThreadLocal) == cudaSuccess); @@ -3633,6 +3667,15 @@ TEST_CASE("CUDA public GdnDecode binds decode layout to exact production geometr REQUIRE(shared_production_function != nullptr); CHECK(params.func != shared_production_function); } + if (expected.register_state && !expected.vector_writeback && + std::string(expected.regstate_env) == "1" && + std::string(expected.vecstore_env) == "0") { + register_production_function = params.func; + } + if (expected.vector_writeback) { + REQUIRE(register_production_function != nullptr); + CHECK(params.func != register_production_function); + } CHECK(params.gridDim.x == expected.grid_x); CHECK(params.gridDim.y == static_cast(n * hv)); CHECK(params.gridDim.z == 1); @@ -3642,6 +3685,7 @@ TEST_CASE("CUDA public GdnDecode binds decode layout to exact production geometr CHECK(params.sharedMemBytes == expected.shared_bytes); } unsetenv("VT_GDN_DECODE_REGSTATE"); + unsetenv("VT_GDN_DECODE_BF16_VECSTORE"); unsetenv("VT_GDN_DECODE_SWIZZLE"); unsetenv("VT_GDN_DECODE_BV"); unsetenv("VT_GDN_DECODE_NW"); From ae85186e91ab4ea4e6421fef2d9ece187514b3a5 Mon Sep 17 00:00:00 2001 From: Richard Palethorpe Date: Sun, 9 Aug 2026 23:30:52 +0100 Subject: [PATCH 13/17] fix(KERNEL-SSM-MAMBA): bind BF16 vector writeback contracts FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: Codex:gpt-5.6-sol [Codex] --- src/vt/cuda/cuda_gdn.cu | 31 ++++----- src/vt/cuda/gdn_decode_fused.h | 74 +++++++++++++++++++- tests/CMakeLists.txt | 1 + tests/vt/test_gdn_decode_fused.cpp | 104 ++++++++++++++++++++++++++--- 4 files changed, 178 insertions(+), 32 deletions(-) diff --git a/src/vt/cuda/cuda_gdn.cu b/src/vt/cuda/cuda_gdn.cu index 872eb7093..bcda6fda4 100644 --- a/src/vt/cuda/cuda_gdn.cu +++ b/src/vt/cuda/cuda_gdn.cu @@ -278,17 +278,13 @@ __device__ inline void StoreBf16Vec8(__nv_bfloat16* p, float v0, float v1, float v6, float v7) { static_assert(sizeof(__nv_bfloat16) * 8 == sizeof(int4)); static_assert(alignof(int4) == 16); - int4 packed; - auto* elements = reinterpret_cast<__nv_bfloat16*>(&packed); - elements[0] = __float2bfloat16(v0); - elements[1] = __float2bfloat16(v1); - elements[2] = __float2bfloat16(v2); - elements[3] = __float2bfloat16(v3); - elements[4] = __float2bfloat16(v4); - elements[5] = __float2bfloat16(v5); - elements[6] = __float2bfloat16(v6); - elements[7] = __float2bfloat16(v7); - *reinterpret_cast(p) = packed; + const float values[8] = {v0, v1, v2, v3, v4, v5, v6, v7}; + const GdnDecodeBf16Pack8 packed = GdnDecodePackBf16Vec8(values); + const int4 transaction{static_cast(packed.words[0]), + static_cast(packed.words[1]), + static_cast(packed.words[2]), + static_cast(packed.words[3])}; + *reinterpret_cast(p) = transaction; } template @@ -2962,7 +2958,7 @@ void LaunchGdnDecodeFusedNW(cudaStream_t s, Tensor& out, const Tensor& q_in, con const Tensor& v, const Tensor& g, const Tensor& beta, Tensor& state, const int32_t* state_idx, int64_t n, const GdnArgs& args, const GdnDecodeLaunchContract& contract, - bool bf16_vecstore) { + const GdnDecodeBf16VecstoreCapability& bf16_vecstore) { const int64_t hk_n = q_in.shape[1], dk = q_in.shape[2]; const int64_t hv_n = v.shape[1], dv = v.shape[2]; const int64_t bv = contract.value_tile; @@ -2983,8 +2979,7 @@ void LaunchGdnDecodeFusedNW(cudaStream_t s, Tensor& out, const Tensor& q_in, con }, [&](const GdnDecodeLaunchContract&) { if constexpr (std::is_same::value) { - DispatchGdnDecodeBf16Vecstore( - bf16_vecstore, + bf16_vecstore.Dispatch( [&] { GdnDecodeFusedKernel @@ -3039,7 +3034,7 @@ void LaunchGdnDecodeFused(cudaStream_t s, Tensor& out, const Tensor& q_in, const const Tensor& v, const Tensor& g, const Tensor& beta, Tensor& state, const int32_t* state_idx, int64_t n, const GdnArgs& args, const GdnDecodeLaunchContract& contract, - bool bf16_vecstore) { + const GdnDecodeBf16VecstoreCapability& bf16_vecstore) { switch (contract.lanes_per_row) { case 2: LaunchGdnDecodeFusedNW(s, out, q_in, k, v, g, beta, state, state_idx, @@ -3071,7 +3066,7 @@ void LaunchGdnDecodeFusedS(cudaStream_t s, Tensor& out, const Tensor& q_in, cons const Tensor& v, const Tensor& g, const Tensor& beta, Tensor& state, const int32_t* state_idx, int64_t n, const GdnArgs& args, const GdnDecodeLaunchContract& contract, - bool bf16_vecstore) { + const GdnDecodeBf16VecstoreCapability& bf16_vecstore) { if (state.dtype == DType::kBF16) LaunchGdnDecodeFused(s, out, q_in, k, v, g, beta, state, state_idx, n, args, contract, @@ -3117,9 +3112,9 @@ void GdnDecodeFusedCuda(Queue& q, Tensor& out, const Tensor& q_in, const Tensor& args, "gdn_decode"); return; } - const bool bf16_vecstore = GdnDecodeBf16VecstoreEligible( + const GdnDecodeBf16VecstoreCapability bf16_vecstore{ std::getenv("VT_GDN_DECODE_BF16_VECSTORE"), contract, dv, dk, nw, - state.dtype == DType::kBF16, true); + state.dtype == DType::kBF16, true}; if (q_in.dtype == DType::kF32) { if (out.dtype == DType::kF32) LaunchGdnDecodeFusedS(s, out, q_in, k, v, g, beta, state, state_idx, n, diff --git a/src/vt/cuda/gdn_decode_fused.h b/src/vt/cuda/gdn_decode_fused.h index 1fa93e54d..9b4eb2c05 100644 --- a/src/vt/cuda/gdn_decode_fused.h +++ b/src/vt/cuda/gdn_decode_fused.h @@ -5,6 +5,7 @@ #include #include +#include namespace vt::cuda { @@ -80,6 +81,48 @@ static_assert(GdnDecodeBf16PackByteOffset(0, 0) == 0); static_assert(GdnDecodeBf16PackByteOffset(0, 1) == 16); static_assert(GdnDecodeBf16PackByteOffset(7, 1) == 240); +// Portable representation of the exact 16-byte transaction emitted by the +// CUDA writeback. Keeping conversion and packing here lets the host tests +// validate the bytes that StoreBf16Vec8 actually consumes. +struct alignas(16) GdnDecodeBf16Pack8 { + uint32_t words[4]; +}; + +#if defined(__CUDACC__) +__host__ __device__ +#endif +inline uint16_t GdnDecodeBf16RnBits(float value) { + uint32_t bits = 0; +#if defined(__CUDA_ARCH__) + bits = __float_as_uint(value); +#else + std::memcpy(&bits, &value, sizeof(bits)); +#endif + const uint32_t exponent = bits & 0x7f800000u; + const uint32_t mantissa = bits & 0x007fffffu; + if (exponent == 0x7f800000u && mantissa != 0) { + return static_cast((bits >> 16) | 0x0040u); + } + const uint32_t rounding_bias = 0x7fffu + ((bits >> 16) & 1u); + return static_cast((bits + rounding_bias) >> 16); +} + +#if defined(__CUDACC__) +__host__ __device__ +#endif +inline GdnDecodeBf16Pack8 GdnDecodePackBf16Vec8(const float* values) { + GdnDecodeBf16Pack8 packed{}; + for (int pair = 0; pair < 4; ++pair) { + const uint32_t lo = GdnDecodeBf16RnBits(values[pair * 2]); + const uint32_t hi = GdnDecodeBf16RnBits(values[pair * 2 + 1]); + packed.words[pair] = lo | (hi << 16); + } + return packed; +} + +static_assert(sizeof(GdnDecodeBf16Pack8) == 16); +static_assert(alignof(GdnDecodeBf16Pack8) == 16); + #if defined(__CUDACC__) __host__ __device__ #endif @@ -109,11 +152,36 @@ inline constexpr bool GdnDecodeBf16VecstoreEligible( return GdnDecodeBf16VecstoreFlagIsOn(env_value) && contract.regstate && contract.swizzled && contract.selected_tile == GdnDecodeValueTile::kBv16 && - contract.value_tile == 16 && contract.lanes_per_row == 8 && - contract.block_threads == 128 && dv == 128 && dk == 128 && - requested_nw == 8 && state_is_bf16 && regular_fused_decode; + contract.value_tile == 16 && contract.value_tiles == 8 && + contract.lanes_per_row == 8 && contract.block_threads == 128 && + contract.shared_bytes == 9728 && contract.should_launch && dv == 128 && + dk == 128 && requested_nw == 8 && state_is_bf16 && + regular_fused_decode; } +struct GdnDecodeBf16VecstoreCapability { + const char* env_value; + GdnDecodeLaunchContract contract; + int64_t dv; + int64_t dk; + int requested_nw; + bool state_is_bf16; + bool regular_fused_decode; + + constexpr bool Eligible() const { + return GdnDecodeBf16VecstoreEligible( + env_value, contract, dv, dk, requested_nw, state_is_bf16, + regular_fused_decode); + } + + template + inline decltype(auto) Dispatch(IncumbentLaunch&& incumbent_launch, + VectorLaunch&& vector_launch) const { + if (Eligible()) return vector_launch(); + return incumbent_launch(); + } +}; + inline constexpr GdnDecodeLaunchContract GdnDecodeLaunchContractFor( const char* bv_env_value, const char* swizzle_env_value, const char* regstate_env_value, int64_t dv, int64_t dk, diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt index f02c9de74..03e6a97e1 100644 --- a/tests/CMakeLists.txt +++ b/tests/CMakeLists.txt @@ -894,6 +894,7 @@ target_include_directories(test_gdn_packed_decode_triton PRIVATE ${CMAKE_SOURCE_ # in cuda_gdn.cu and is covered by same-process byte-exact CUDA comparisons. vllm_cpp_add_test(test_gdn_decode_fused vt/test_gdn_decode_fused.cpp) target_include_directories(test_gdn_decode_fused PRIVATE ${CMAKE_SOURCE_DIR}/src) +target_compile_definitions(test_gdn_decode_fused PRIVATE VLLM_CPP_SOURCE_DIR="${CMAKE_SOURCE_DIR}") # CPU-tier contract for the two NUMERICS-NEUTRAL FP4-quant vectorized-load # fast-path flags (VT_FP4_QUANT_FAST + VT_SILU_FP4_FAST, both default OFF / # '1'-opt-in). Reaches the internal header under src/ (not public API); the diff --git a/tests/vt/test_gdn_decode_fused.cpp b/tests/vt/test_gdn_decode_fused.cpp index 8ef25ba90..bf62468fb 100644 --- a/tests/vt/test_gdn_decode_fused.cpp +++ b/tests/vt/test_gdn_decode_fused.cpp @@ -4,7 +4,10 @@ #include #include +#include #include +#include +#include #include "vt/cuda/gdn_decode_fused.h" @@ -19,6 +22,8 @@ using vt::cuda::GdnDecodeBf16PackByteOffset; using vt::cuda::GdnDecodeBf16PackLogicalColumn; using vt::cuda::GdnDecodeBf16VecstoreEligible; using vt::cuda::GdnDecodeBf16VecstoreFlagIsOn; +using vt::cuda::GdnDecodeBf16VecstoreCapability; +using vt::cuda::GdnDecodePackBf16Vec8; using vt::cuda::GdnDecodeSharedColumn; using vt::cuda::GdnDecodeStateStride; using vt::cuda::GdnDecodeValueTile; @@ -338,6 +343,37 @@ TEST_CASE("GDN BF16 vector writeback requires every production predicate") { CHECK_FALSE(eligible("1", production, 128, 128, 4, true, true)); CHECK_FALSE(eligible("1", production, 128, 128, 8, false, true)); CHECK_FALSE(eligible("1", production, 128, 128, 8, true, false)); + + // Corrupt exactly one field at a time. Tests made from another factory + // contract can leave several correlated fields false and therefore survive + // deletion of the predicate they claim to cover. + auto corrupted = production; + corrupted.regstate = false; + CHECK_FALSE(eligible("1", corrupted, 128, 128, 8, true, true)); + corrupted = production; + corrupted.swizzled = false; + CHECK_FALSE(eligible("1", corrupted, 128, 128, 8, true, true)); + corrupted = production; + corrupted.selected_tile = GdnDecodeValueTile::kBv32; + CHECK_FALSE(eligible("1", corrupted, 128, 128, 8, true, true)); + corrupted = production; + corrupted.value_tile = 32; + CHECK_FALSE(eligible("1", corrupted, 128, 128, 8, true, true)); + corrupted = production; + corrupted.value_tiles = 4; + CHECK_FALSE(eligible("1", corrupted, 128, 128, 8, true, true)); + corrupted = production; + corrupted.lanes_per_row = 4; + CHECK_FALSE(eligible("1", corrupted, 128, 128, 8, true, true)); + corrupted = production; + corrupted.block_threads = 64; + CHECK_FALSE(eligible("1", corrupted, 128, 128, 8, true, true)); + corrupted = production; + corrupted.shared_bytes = 9727; + CHECK_FALSE(eligible("1", corrupted, 128, 128, 8, true, true)); + corrupted = production; + corrupted.should_launch = false; + CHECK_FALSE(eligible("1", corrupted, 128, 128, 8, true, true)); } TEST_CASE("GDN BF16 vector writeback dispatches exactly one callback") { @@ -386,18 +422,64 @@ TEST_CASE("GDN BF16 vector packs cover the aligned 128-column state row") { 0.0f, -0.0f, 1.0f, -2.0f, 0.333251953125f, 65504.0f, 1.00390625f, 1.01171875f, -1.00390625f, -1.01171875f, 1.0e-20f, -1.0e-20f, 3.14159265f, -2.71828183f, 42.5f, -99.75f}; - std::array packed{}; for (int pack = 0; pack < 2; ++pack) { - for (int element = 0; element < 8; ++element) { - const uint16_t bits = ReferenceBf16Rn(values[pack * 8 + element]); - std::memcpy(packed.data() + pack * 16 + element * 2, &bits, - sizeof(bits)); + const auto packed = GdnDecodePackBf16Vec8(values.data() + pack * 8); + for (int pair = 0; pair < 4; ++pair) { + const uint16_t lo = static_cast(packed.words[pair]); + const uint16_t hi = static_cast(packed.words[pair] >> 16); + CHECK(lo == ReferenceBf16Rn(values[pack * 8 + pair * 2])); + CHECK(hi == ReferenceBf16Rn(values[pack * 8 + pair * 2 + 1])); } } - for (int element = 0; element < 16; ++element) { - uint16_t packed_bits = 0; - std::memcpy(&packed_bits, packed.data() + element * 2, - sizeof(packed_bits)); - CHECK(packed_bits == ReferenceBf16Rn(values[element])); - } +} + +TEST_CASE("GDN BF16 vector capability binds eligibility and launch") { + const auto production = + GdnDecodeLaunchContractFor("16", "1", "1", 128, 128, 8); + auto select = [&](const char* env, bool state_is_bf16) { + int incumbent = 0; + int vector = 0; + const GdnDecodeBf16VecstoreCapability capability{ + env, production, 128, 128, 8, state_is_bf16, true}; + const int result = capability.Dispatch( + [&] { return ++incumbent, 0; }, [&] { return ++vector, 1; }); + CHECK(incumbent + vector == 1); + return result; + }; + CHECK(select("1", true) == 1); + CHECK(select("0", true) == 0); + CHECK(select("1", false) == 0); +} + +TEST_CASE("CUDA GDN call site uses the tested vector capability and pack seam") { + std::ifstream input(std::string(VLLM_CPP_SOURCE_DIR) + + "/src/vt/cuda/cuda_gdn.cu"); + REQUIRE(input.good()); + const std::string source((std::istreambuf_iterator(input)), + std::istreambuf_iterator()); + + const auto store_begin = source.find("__device__ inline void StoreBf16Vec8("); + const auto store_end = source.find("\n}", store_begin); + REQUIRE(store_begin != std::string::npos); + REQUIRE(store_end != std::string::npos); + const std::string store = source.substr(store_begin, store_end - store_begin); + CHECK(store.find("GdnDecodePackBf16Vec8(values)") != std::string::npos); + CHECK(store.find("*reinterpret_cast(p) = transaction") != + std::string::npos); + + const auto kernel = source.find("if constexpr (BF16_VECSTORE) {"); + REQUIRE(kernel != std::string::npos); + const std::string writeback = source.substr(kernel, 900); + CHECK(writeback.find("for (int pack = 0; pack < 2; ++pack)") != + std::string::npos); + CHECK(writeback.find("StoreBf16Vec8(") != std::string::npos); + + const auto route = source.find( + "const GdnDecodeBf16VecstoreCapability bf16_vecstore{"); + REQUIRE(route != std::string::npos); + const std::string invocation = source.substr(route, 260); + CHECK(invocation.find("VT_GDN_DECODE_BF16_VECSTORE") != std::string::npos); + CHECK(invocation.find("state.dtype == DType::kBF16, true") != + std::string::npos); + CHECK(source.find("bf16_vecstore.Dispatch(") != std::string::npos); } From 96919cc5aca5041785a682cb04cce89945d55f48 Mon Sep 17 00:00:00 2001 From: Richard Palethorpe Date: Sun, 9 Aug 2026 23:49:33 +0100 Subject: [PATCH 14/17] perf(KERNEL-SSM-MAMBA): close sm120 Pareto experiments Remove the geometric argmax scratch and BF16 GDN vector writeback candidates after same-binary profiles miss their Pareto gates. Record the corrected atomic-wave baseline, retained local opt-ins, and falsified trace evidence. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: Codex:gpt-5.6-sol [Codex] --- .agents/NOW.md | 4 +- .agents/benchmark-record.md | 73 ++++ .agents/kernel-matrix.md | 23 +- .agents/roadmap_v1.md | 2 +- ...qwen35-argmax-scratch-growth-2026-08-09.md | 28 +- ...decode-bf16-vector-writeback-2026-08-09.md | 37 +- docs/BENCHMARKS.md | 2 +- docs/STATUS.md | 2 +- include/vt/cuda/argmax_scratch.h | 211 ---------- scripts/check-public-doc-tables.py | 7 +- scripts/env-doc-allowlist.txt | 2 - src/vt/cuda/cuda_gdn.cu | 136 ++----- src/vt/cuda/cuda_sample.cu | 69 +--- src/vt/cuda/gdn_decode_fused.h | 117 ------ tests/CMakeLists.txt | 3 - tests/vt/test_argmax_scratch.cpp | 375 ------------------ tests/vt/test_gdn_decode_fused.cpp | 204 ---------- tests/vt/test_ops_gdn.cpp | 64 +-- 18 files changed, 204 insertions(+), 1155 deletions(-) delete mode 100644 include/vt/cuda/argmax_scratch.h delete mode 100644 tests/vt/test_argmax_scratch.cpp diff --git a/.agents/NOW.md b/.agents/NOW.md index 2e6bc0e55..3462bf4f4 100644 --- a/.agents/NOW.md +++ b/.agents/NOW.md @@ -21,7 +21,7 @@ Work: exact-chunks on main `1ce0d662b`; sm_120 measured at `3d2581551`. | MiniMax-H3 | **PRUNED ckpts RUN (#241): Q8_0 renders, seam 0.9941** | same-binary A/B | | Kimi-Linear-48B | 122/128 held; grouped router parallelised, e2e NOT ESTABLISHED | ckpt is tiktoken-only: no warm server | | 35B binding grid | @`a0fa12c7`: **flat 0.935-0.979x, NO c2/c8 weak cell** (CoV <0.81%); mem PSS 3.81x | Attribute the flat ~5% mid-band; TTFT c2 0.872x | -| Qwen3.5-4B sm_120 | #206 pretoken A/B FAIL: 98/128 exact; timings VOID. Atomic wave CPU-green | Review; GPU A/B | +| Qwen3.5-4B sm_120 | Exact: tput **1.0283x**; TTFT/TPOT/E2E **1.0853/1.0165/1.0288x slower**; VRAM +118.7 MiB | Profile TTFT split | | RPi5 A76 CPU | **R5 asm GREEN; llama NOT MET**: 0.461x pf, 0.653x dec | W6: BF16 GEMM | | MXFP4 parity | c1 1.020, c2-c8 0.962-0.969. **#82 CLOSED: ptxas-lineage REFUTED** | TERMINAL: at parity | | SERVE-ASYNC-DENSE-MIRROR | **LANDED+VERIFIED** (`f9c969ae`): async mirror, dense Qwen3; SACRED 184/184 | Sibling scope one-liner | @@ -53,7 +53,7 @@ latency/memory on every axis, both gate models, reproduced 2–3x idle. See of roof. Dense-marlin +0.5%; Triton-AOT GDN a WASH. 2. **Spike the Parakeet encoder row** (vLLM: `nano_nemotron_vl.py`; the transducer half is NOT in vLLM: separate call). -3. **Qwen3.5-4B #206:** atomic admission CPU-green; review, GPU/vLLM rerun. +3. **Qwen3.5-4B #206:** tput +2.83%; latency/VRAM open. Profile TTFT split. 2. **Merge the invocation-parity prevention** (CI guard + AGENTS.md checklist); CUDA build-verify the byte-exact `kGemvHeuristicAlgos` refactor on dgx. 4. **Restore `local-ai-worker`** on dgx at campaign end (`--restart=always`). diff --git a/.agents/benchmark-record.md b/.agents/benchmark-record.md index 037402374..f207c2430 100644 --- a/.agents/benchmark-record.md +++ b/.agents/benchmark-record.md @@ -18165,3 +18165,76 @@ arrivals for a whole refill wave and publish it once, so string rollback still times tokenization while both arms execute identical batches. The full RED-first, rollback and mutation gates are binding in the [campaign spec](specs/sm120-qwen35-pareto-2026-08-09.md). + +## 2026-08-09 — Qwen3.5 corrected frontend comparison and retained local decode stack + +Atomic all-at-once request-wave admission repaired the pretokenized benchmark's +batching confound at reviewed head `da449a88f0a06710839a0a6568f18426cb44fc19`. +All repeated local and pinned-vLLM legs generated the same 128 x 128-token +workload; the local token SHA-256 is +`be20ffbceb61f0264ca21d972bfc5fc51e855e64f2b945de71669cae666aa702`. +The corrected, non-profiled three-repetition comparison is: + +| Engine | total tok/s | output tok/s | mean TTFT | mean TPOT / ITL | mean E2E | peak VRAM | +|---|---:|---:|---:|---:|---:|---:| +| vllm.cpp direct-load | 6831.7100 | 755.4300 | 1016.5133 ms | 34.4767 ms | 5395.1333 ms | 12,950.7 MiB | +| pinned vLLM `555967922` | 6643.4025 | 734.6087 | 936.5893 ms | 33.9171 ms | 5244.0659 ms | 12,832.0 MiB | +| ratio / gap | **1.028345x** | **1.028345x** | **1.085335x slower** | **1.016497x slower** | **1.028807x slower** | **+118.7 MiB** | + +Host peak/stable PSS is 2.338/0.770 GiB locally versus 7.868/4.496 GiB for +vLLM, so host memory remains a large win. Raw comparison root is +`/tmp/qwen35-compare-da449a88f/`; same-tool traces are +`/tmp/qwen35-profile-da449a88f-{ours,vllm}.{nsys-rep,sqlite}`. Throughput is +now a valid win; TTFT, TPOT/ITL, E2E and VRAM remain open. + +Two later exact, counterbalanced local stacks were retained as default-OFF +opt-ins. `VT_GDN_DECODE_BV=16`, `VT_GDN_DECODE_SWIZZLE=1` and +`VT_GDN_DECODE_REGSTATE=1` moved the direct-load local mean from +6838.4400/756.1767 tok/s, 1014.0233 ms TTFT, 34.4533 ms TPOT and 5389.7900 ms +E2E to **6852.1200/757.6900**, **1011.8500**, **34.3867** and **5378.9500** +(+0.200% throughput, -0.214% TTFT, -0.193% TPOT, -0.201% E2E). Adding +`VT_GDN_SLACK_MEMSET=1` produced **6856.8633/758.2133**, **1010.4800 ms**, +**34.3700 ms** and **5375.1900 ms** against its adjacent decode-stack mean +6852.5700/757.7367, 1011.5233, 34.3900 and 5378.6167. These local gains have +not been consolidated into a new pinned-vLLM ratio; do not project them as a +cross-engine pass. Roots are `/tmp/qwen35-ab-decode-stack-da449a88f/` and +`/tmp/qwen35-ab-slack-memset-da449a88f/`. + +## 2026-08-09 — geometric argmax scratch rejected and removed + +Reviewed head `06db3bbb3454f85e20f75b96eeb299afba7036ae` passed 9 portable +cases / 105 assertions plus CUDA sampling. All six legs were token-exact. + +| Arm mean, 3 reps | total / output tok/s | TTFT | TPOT / ITL | E2E | +|---|---:|---:|---:|---:| +| Incumbent | 6862.2667 / 758.8100 | 1009.7167 ms | 34.3400 ms | 5371.0000 ms | +| Geometric scratch | 6863.8867 / 758.9867 | 965.0133 ms | 34.6833 ms | 5369.7067 ms | + +The apparent -4.427% TTFT coincides with a **+0.9998% TPOT regression**, so +the candidate fails the Pareto gate. Same-tool traces identify wait migration, +not saved compute: `cudaFree` falls 506 calls / 2960.361 ms to 476 / 27.705 ms +(exactly 30 argmax frees), while the same 567 `cudaStreamSynchronize` calls rise +18,027.497 -> 20,878.260 ms. Raw A/B is +`/tmp/qwen35-ab-argmax-geometric-06db3bbb/`; candidate profiler hashes are +`1021b045...239e7` and `808a7516...b7be4`. Product and tests are removed; the +falsified allocation hypothesis remains documented. + +## 2026-08-09 — BF16 GDN vector writeback rejected and removed + +Implementation `29de225c8` plus repair `3ca49e926` passed fresh mutation +review (16 cases / 2,568 assertions), operator CUDA rebuild and focused +`test_gdn_decode_fused` + `test_ops_gdn`. All four profiled legs were exact. + +| Arm | all fused decode | y800 decode | total / output tok/s | TTFT | TPOT | E2E | +|---|---:|---:|---:|---:|---:|---:| +| REG-a | 251.210058 ms | 139.127859 ms | 6760.82 / 747.59 | 1027.95 | 34.81 | 5449.30 | +| VEC-a | 250.472657 ms | 138.702659 ms | 6802.32 / 752.18 | 1015.76 | 34.66 | 5417.91 | +| VEC-b | 251.214368 ms | 139.138796 ms | 6804.21 / 752.39 | 1015.60 | 34.65 | 5416.35 | +| REG-b | 251.404693 ms | 139.111428 ms | 6807.51 / 752.76 | 1013.82 | 34.65 | 5413.73 | + +Counterbalanced means improve y800 only **0.1430%** and all fused decode only +**0.1846%**; VEC-b loses to both controls on y800. It therefore misses the +required 1% and raw-leg gates. `nsys` records unchanged gx8/gy800/bx128, +56 registers/thread, 9,728 dynamic shared bytes and zero local bytes. Evidence +is `/tmp/qwen35-gdn-vecstore-{rega,veca,vecb,regb}-3ca49e926.*`. Product and +tests are removed; REGSTATE remains the local opt-in. diff --git a/.agents/kernel-matrix.md b/.agents/kernel-matrix.md index 290f26474..773809fcb 100644 --- a/.agents/kernel-matrix.md +++ b/.agents/kernel-matrix.md @@ -174,18 +174,17 @@ release gates are unchanged. [Spec and evidence](specs/sm120-qwen35-conv-chunkin **2026-08-09 `KERNEL-SSM-MAMBA` sm_120 campaign anchor.** [#206](https://github.com/mudler/vllm.cpp/issues/206) tracks the RTX 5070 Ti -Qwen3.5-4B Pareto campaign. The historical sealed measurement reported a -throughput pass (**6784.743 vs 6643.593 tok/s**) while TTFT -(**1018.040 vs 937.584 ms**), TPOT/ITL (**34.740 vs 33.906 ms**) and VRAM -(**13053.3 vs 12820 MiB**) remain open. The reviewed combined default-off K=4 -causal-conv plus 16-token post-conv A/B is accepted locally: **+0.7175%** total -throughput, **-1.340%** TTFT, **-0.556%** TPOT and **-0.716%** E2E, with six -identical token hashes. The cross-engine throughput and TTFT ratios are now -**VOID/PENDING** because local timed string tokenization while pinned vLLM -pre-tokenized before its timer; TPOT remains comparable and open. The committed -next gate pre-tokenizes local prompts before `t0`, retains -`VT_BENCH_PRETOKENIZE=0` same-binary rollback, and reruns both engines. Lifecycle -stays `INVENTORIED`; generic Mamba and 27B/35B coverage are unchanged. +Qwen3.5-4B Pareto campaign. Reviewed atomic pretoken admission removed the +frontend/batching confound; the corrected exact three-repetition comparison is +**6831.71 vs 6643.40 tok/s (1.0283x)**, while TTFT, TPOT/ITL and E2E remain +**1.0853x / 1.0165x / 1.0288x slower** and VRAM is **+118.7 MiB**. The exact +default-OFF GDN decode BV16+swizzle+REGSTATE stack improves local throughput, +TTFT, TPOT and E2E, and slack-only memset adds a smaller further local win. +Two deeper candidates are closed: geometric argmax scratch merely moved wait +from `cudaFree` to stream synchronization and regressed TPOT ~1%; BF16 vector +writeback improved y800 only 0.143% with one losing raw leg. Both products/tests +were removed while their specs and same-tool traces remain. Lifecycle stays +`INVENTORIED`; generic Mamba and 27B/35B coverage are unchanged. [Campaign spec](specs/sm120-qwen35-pareto-2026-08-09.md). ## Count invariants diff --git a/.agents/roadmap_v1.md b/.agents/roadmap_v1.md index 55a2bf0c6..5be9f3844 100644 --- a/.agents/roadmap_v1.md +++ b/.agents/roadmap_v1.md @@ -101,7 +101,7 @@ models we already ship + benchmark. Full seam map + M0–M5 W-plan: | MEM | `ROAD-V1-MEM` | **Memory budgeting: auto-size to the declared workload by default, optional total-footprint cap, pre-flight error instead of an OOM (user-directed 2026-08-06, [#83](https://github.com/mudler/vllm.cpp/issues/83))** — the user-facing wart that every operator hits before they hit any perf question: vLLM makes you compute your own VRAM budget, express it as a PERCENT, and nail it or OOM | [coverage view §2](feature-matrix.md#2-kv-cache--memory), [porting inventory](porting-inventory.md) | — **M1+M2 LANDED 2026-08-08** ([`specs/kv-sizing.md`](specs/kv-sizing.md)): the absolute `--kv-cache-memory` knob sizes the pool via a group-aware `KVBytesPerBlock` divisor, `--num-blocks` is the override, both mirrored on the C ABI at v16; `ResolveNumBlocks` precedence `num_blocks > bytes > 256`, CPU-gated (`KVBytesPerBlock` 5/5 + `test_capi` v16). M3 (the `gpu_memory_utilization` profile run) stays dgx-gated. (M0 design grounded in vLLM `config/cache.py` + `gpu_worker.py:497-599`; GB10 unified-pool caveat carried) | `M1+M2 DONE` | **Source-verified 2026-08-06 (records-only, NO code).** WE ARE CURRENTLY BEHIND vLLM ON THIS AXIS, not ahead: there is NO memory profiling at all and the KV pool is a RAW BLOCK COUNT the user types by hand — `EngineParams::num_blocks = 256` (`include/vllm/entrypoints/model_loader.h:58`, beside `block_size = 32` `:57` / `max_model_len` `:59` / `max_num_seqs = 8` `:60`), exposed verbatim as `--num-blocks N` (`examples/server/main.cpp:100,203-204,370`), carried on the C ABI as `vllm_model_params.num_blocks` at the same 256 default (`src/capi/vllm_c.cpp:429,485`), landing as `BlockPool(num_gpu_blocks, ...)` which asserts `> 0` and otherwise TRUSTS it (`include/vllm/v1/core/block_pool.h:96,223`; `src/vllm/v1/core/block_pool.cpp:51`). So a user must convert "40 GB free, 32k context, concurrency 8" into a block count themselves — strictly worse ergonomics than a percentage. UPSTREAM HAS THREE KNOBS, all `config/cache.py`, all T0, all rowed at [porting-inventory.md](porting-inventory.md) §T0: `gpu_memory_utilization` (fraction of TOTAL, default 0.9), `kv_cache_memory_bytes` (absolute KV pool) and `num_gpu_blocks_override` (exact block pin), sized off a profile run as `total x utilization - non-torch - peak activation`. **Mirroring that is necessary but does NOT solve the complaint**, for three source-grounded reasons: (1) the fraction is of TOTAL not FREE, so on any shared card the right fraction is a function of what someone else already holds — exactly the arithmetic the user is being asked to do; (2) WEIGHTS LOAD BEFORE THE KNOB ENGAGES (utilization sizes the KV pool AFTER the model is resident), so an oversized model OOMs during load and never reaches the check — which is the failure operators actually hit; (3) 0.9 is taken whether or not it is needed (a 4B model on an 80 GB card reserves 72 GB it will never touch and blocks everything else on the device). THE DESIGN (user-ratified 2026-08-06) is three modes: **Mode 1 default = SIZE TO THE DECLARED WORKLOAD** — per-allocation-class accounting BEFORE allocating anything (weights from checkpoint metadata / safetensors header / GGUF manifest, known before reading a byte of tensor data; CUDA context measured at creation; peak activation from a profile run at `max_num_batched_tokens`; KV for `max_model_len x max_num_seqs` at the resolved `block_size`/KV dtype; CUDA-graph capture-set footprint) → allocate exactly that and LEAVE THE REMAINDER OF THE DEVICE FREE. This is the SURPASS over vLLM, which takes its 90% regardless of whether the workload needs 8 GiB or 80. **Mode 2 = a cap on the TOTAL ENGINE FOOTPRINT** (weights + activations + KV + graph pools + context), NOT on the KV pool alone — load-bearing, because a KV-only cap cannot prevent the weight-load OOM; three spellings of the same cap: `--memory-limit 40GiB` (absolute, the primary form), `--gpu-memory-utilization 0.85` (vLLM's exact flag name and fraction semantics so existing vLLM launch lines port unchanged, per [[mirror-vllm-always-no-asking]]) and `--num-gpu-blocks-override N` (upstream's reproducibility escape hatch — this is where today's `--num-blocks` GOES, demoted from primary knob to explicit override), with precedence spelled out and TESTED, not left to argument order. **Mode 3 = REFUSE BEFORE ALLOCATING** with the full per-class breakdown and remedies COMPUTED FROM THE ACTUAL BUDGET (`--max-model-len N` / `--max-num-seqs M` / `--kv-cache-dtype fp8` / smaller quant, each with the GiB it recovers) — "you are 43.9 GiB over and here are the three levers that close it" is the difference between a usable error and a stack trace. **UNIFIED-MEMORY HAZARD (not hypothetical):** on GB10 the ~119 GiB pool is UNIFIED, a fraction-of-total setting reserves HOST RAM too, and `gpu_memory_utilization=0.85` has HARD-REBOOTED our DGX three separate times ([[gb10-unified-memory-oom-reboots-box]]) — hence absolute bytes is the PRIMARY form with the percentage kept only for vLLM compatibility, and hence the accounting must know whether the pool is unified, which makes free/total + an is-unified predicate a PLATFORM-SEAM question belonging behind `ROAD-V1-C1`'s abstraction (note `Platform::needs_weight_staging()` was deliberately NOT `is_unified_memory()` because the latter FLIPS GB10 — the distinction matters here) rather than a CUDA-specific branch, since discrete and unified devices need different safety margins. CORRECTNESS: pool size changes preemption/scheduling TIMING but not emitted tokens, so the SACRED token-exact gates are unaffected — and M2's gate makes that explicit by re-running them with no block flag at all. **Next gate = M0 the `specs/kv-sizing.md` spike (accounting model + precedence rules + upstream `file:line`); then M1 a `MemoryBudget` computing required bytes per class WITHOUT allocating plus the Platform free/total + is-unified seam, unit-gated predicted-vs-actual weight bytes; M2 auto-sizing as the default with `--num-blocks` demoted to `--num-gpu-blocks-override`, gated by every existing model gate running with NO block flag and staying token-exact; M3 the three caps + precedence through the server flags and the C ABI, gated by our KV pool matching vLLM's own at a matched `--gpu-memory-utilization`; M4 the pre-flight refusal, gated by a deliberately over-subscribed config exiting cleanly on GB10 (non-zero exit, no OOM, NO BOX REBOOT) covering the unified-pool path specifically; M5 (optional) a runtime guard failing the REQUEST rather than the engine.** Docs (README, [STATUS](../docs/STATUS.md)) update in the SAME change as whichever milestone shifts externally-visible behaviour, per [[keep-readme-current]] | | 1 | `ROAD-V1-C1` | **Extensibility-first (USER PRIORITY 2026-07-18):** drop-in kernel ABI + the MISSING Platform seam + model self-registration — make new GPUs/archs/models ADDITIVE (plan: [extensibility-platform-seam-2026-07-18.md](specs/extensibility-platform-seam-2026-07-18.md)) | [`BACKEND-ABI-VT`](backend-matrix.md), [kernel matrix](kernel-matrix.md) | exhaustive kernel/dependency inventory and [raw-pointer adapter ABI](specs/dropin-kernel-abi.md) accepted; additive W0 implemented and CPU 94/94. `CLAIM-BACKEND-ABI-W0-GPU-1` repaired the GCC13/doctest blocker without runtime changes; exact sm_121a all-target build, focused CUDA/ABI sanitizer, and both gate-model tests pass at `1141b79`. Cross-arch/trace/A-B and scalar-forwarder/backend-shim debts remain explicit | `PARTIAL` | **★ NEW ORDER-1 HEAD (user-directed 2026-07-19): the PORTABLE AUTOMATIC OP-FUSION FRAMEWORK (`KERNEL-FUSION-FRAMEWORK`, spike [portable-fusion-framework.md](specs/portable-fusion-framework.md), `SPIKE`).** The extensibility cornerstone: fusions DECLARED ONCE (backend-agnostic `constexpr FusedRecipe` catalog above `vt::`, transcribing vLLM's finite pattern-pass set `passes/fusion/*` @ `pass_manager.py:138-200`, mirroring the `CustomOp` `forward_native`/`forward_cuda` seam `custom_op.py:103`) and REALIZED PER-BACKEND through the existing `vt::` op table (Tier-0 composite = the CPU oracle inherited free by any backend; Tier-1 interpreter = one kernel port per backend lights up every recipe). Makes a new vLLM fusion PR a ONE-declaration port, a new GPU an additive catalog realization, a new model an additive pattern declaration — the PR-#4 remedy, composed with the Platform/attn-registry/model-registry seams below. The TDR Phase-0 skeleton is already LANDED (`fused_recipe.h`/`recipes.h` one recipe + `OpId::kFusedChain` Tier-0/1 on CPU+CUDA + byte-exact `test_ops_fused_chain.cpp`). **W0 ADOPTED 2026-07-19 (`CLAIM-FUSION-FRAMEWORK-W0`):** the seam is now used in production at ONE real site — the 35B `RunLayerPaged` post-attention layernorm routes its plain add+residual+gemma-RMSNorm through `vt::FusedChain(kFusedAddRmsNorm)` (`VT_FUSED_CHAIN_ADOPT` default-ON / `=0` rollback), behaviour-preserving + byte-identical to the prior hand-call (DGX: clean CUDA `-Werror` 0 warn, byte-exact composite==interp==golden incl. H=2048, 35B 315/315 + 27B 235/235 token-exact BOTH arms, memcheck 0 errors) — proving the declare-once/realize-per-backend seam end-to-end; the current 3-opcode/4-role POD sufficed byte-identically for the plain add+rmsnorm class, so W0 needed NO generalization. **W1 GENERALIZED the POD 2026-07-20 (`CLAIM-FUSION-FRAMEWORK-W1`, `1115648`):** full activation/norm/quant/rope opcode set + indexed operand table; all 5 quant-fused chains declared byte-exact; Tier-0 composite = ONE device-agnostic walker (kills CPU/CUDA oracle drift); infrastructure only, no call site changed (DGX: `-Werror` 0-warn, byte-exact CPU 196 + CUDA 361, memcheck 0, 27B 235/235 + 35B 315/315 both arms). **W2 MIGRATED the hand-fusions 2026-07-20 (`CLAIM-FUSION-FRAMEWORK-W2`):** the framework now OWNS the fusion dispatch — each recipe binds (new backend-agnostic `FusedRecipe.fast_op`) to its EXISTING single-launch fused kernel, so `FusedChain(recipe)` dispatches to the SAME fast kernel the model called directly pre-migration (byte-identical + perf-neutral by construction; composite is the graceful fallback + oracle). SIX call sites migrated behind `VT_FUSED_CHAIN_ADOPT` (`kSiluMulFp4Quant`/`kSigmoidGateFp4Quant`/`kRmsNormGatedQuantFp8`×2/`kRmsNormQuantFp8`/`kAttnQkNormRopeGate`×2). DGX: `-Werror` 0-warn, byte-exact CPU 228 + CUDA 420, memcheck 0, 27B 235/235 + 35B 315/315 BOTH arms. **W3 MECHANICAL-SYNC PROOF LANDED 2026-07-20 (`CLAIM-FUSION-FRAMEWORK-W3`):** ported a NEW, previously-unported vLLM fusion pass — `SiluMulFp8StaticQuantPattern` (`act_quant_fusion.py:81` → `_C.silu_and_mul_quant`, the static-per-tensor-FP8 sibling of `kSiluMulFp4Quant`) — as ONE `constexpr FusedRecipe kSiluMulQuantFp8` + its byte-exact test, touching EXACTLY 2 shared files (`recipes.h` + `test_ops_fused_chain.cpp`), NO kernel/dispatch/composite-walker/model-site edit and NO new primitive (composite = existing `vt::MoeSiluMul` + `vt::QuantFp8Static`; `fast_op=kNoFastOp`). The PR-#4 additivity test made concrete: a whole new fusion pattern = one declaration. DGX: `-Werror` 0-warn, byte-exact CUDA 432, memcheck 0, no token regression (recipe declared-only) 27B 235/235 + 35B 315/315. **W4 BACKEND-ADDITIVITY PROOF LANDED 2026-07-20 (`CLAIM-FUSION-FRAMEWORK-W4`) — the W-series proof milestone is DONE:** made the additivity claim EXECUTABLE — new test `test_fused_chain_additivity.cpp` treats the EXISTING CPU backend AS the 'second backend' relative to CUDA (no mock `DeviceType` — that would edit the core enum + every switch, ironically non-additive) and, in ONE generic loop over the WHOLE catalog (all 7 recipes), asserts each runs byte-exact on the CPU backend via the Tier-0 composite — 4 CPU-full end-to-end + 3 fp8-prefix (byte-exact prefix + the FULL composite asserted to THROW on CPU, documenting the CUDA-only static-fp8 backend-negotiated tail, §3b/§6). Additivity evidence: catalog `recipes.h` grew 1→6→7 while the composite walker stayed ONE per-OPCODE function (12 `FOp::` cases) + the CPU/CUDA `kFusedChain` registration ONE line each + `cpu_ops.cpp` never `#include`s `recipes.h` — W3's whole new recipe `kSiluMulQuantFp8` is in ZERO backend TUs, inherited free. CPU `-Werror` 0-warn, `test_fused_chain_additivity` 17/17 + `test_ops_fused_chain` 228/228; engine byte-identical (no `src/`/`include/` change) ⇒ 27B 235/235 + 35B 315/315 structurally unchanged; memcheck N/A. Honest deferred (named, non-blocking the ORDER-1 milestone): Tier-1 perf interpreter for the quant chains (composite-only today), a REAL Metal/Vulkan catalog realization (M4 HW-blocked), and per-recipe fast single-launch kernels. Honest payoff: perf ceiling ~3.5%/step compute-bound on 35B (NOT a perf lever — tasks #61/#62; W0 is perf-neutral by construction); primary value = extensibility + mechanical upstream-sync + CPU/CUDA oracle-drift elimination. Incremental W0 adopt-one **(DONE)** → W1 generalize POD **(DONE)** → W2 migrate hand-fusions **(DONE)** → W3 mechanical-sync proof **(DONE)** → W4 backend-additivity proof **(DONE)** → Wn honest re-measure (optional, off the extensibility critical path). **W-SERIES ORDER-1 PROOF MILESTONE DONE 2026-07-20 (`CLAIM-FUSION-FRAMEWORK-W4`).** **PRIOR extensibility items (all LANDED, the seams this composes with):** **#1 extensibility item — extract the Platform seam — LANDED 2026-07-18** (`BACKEND-PLATFORM` `ACTIVE`, `CLAIM-BACKEND-PLATFORM-1`): `include/vllm/platforms/interface.h` + `src/vllm/platforms/{platform,cpu,cuda}.cpp` mirror `vllm/platforms/interface.py:134-229` 1:1; `CurrentPlatform()` self-registered per `DeviceType`; the 7 memory-model/residency `device.type == kCUDA` sites (of PR #4's ~37) now route through it → new-GPU memory model is ONE additive `platforms/.cpp`. Behavior-preserving (clean CPU build + `test_platform` + full CPU CTest green; DGX 235/235 + 315/315 pending). **Item 2 residency-as-Platform-capability LANDED 2026-07-19** (`CLAIM-BACKEND-PLATFORM-2`): the host-free / load-stream / DevicePool-cap decisions in `qwen3_5.cpp` now READ `GetPlatform(.device.type).residency_policy()` (per-device) via the pure `ShouldReleaseHostWeights`/`ShouldInterleaveLoadStream` helpers + `device_pool_cap_bytes`, not an inline `device.type`/env gate; `CudaPlatform.release_host_weights_after_upload` flipped false→true (now CONSUMED ⇒ reproduces today's GB10 host-free-after-Marlin + ~4 GiB load peak EXACTLY); `MarlinMoeEnabled()` stays the orthogonal kernel-path gate. **A new (discrete) GPU sets `residency_policy()` values ⇒ ZERO model edit.** Behavior-preserving (clean CPU build + `test_platform` consumption cases 7/43 + full CPU CTest + tools 164/164 green; **DGX-CONFIRMED @ `62fc0e0`: clean CUDA `-Werror`, 27B 235/235 + 35B 315/315 token-exact, 35B VmHWM ≈ 4.0 GiB load-stream win preserved, memcheck 0 errors**). Then item 3 drop-in ABI family migration. **Item 4 attn-backend registry LANDED 2026-07-19** (`CLAIM-ATTN-REGISTRY-1`, `BACKEND-ATTN-REGISTRY`): NEW `include/vllm/v1/attention/registry.{h,cpp}` `(DeviceType,name)` registry + `SelectAttentionBackendName` selector (mirror `registry.py` self-registration + `cuda.py:361-470` `get_attn_backend_cls`/`_get_backend_priorities`); `Platform::get_attn_backend_priority()` filled (was the item-1 STUB) → capability-ordered name lists on `CudaPlatform` (major-10 vs else) + `CpuPlatform`; FLASH_ATTN/GDN self-register. **Adding a backend's attention = 1 self-registering TU + 1 priority slot, ZERO selector/model/runner edit.** Behavior-preserving — the walk returns FLASH_ATTN on CUDA+CPU (the same FA2 attention runs); clean CPU build + `test_attn_backend_registry` (8/25) + full CPU CTest, tools 164/164, checkers green; **DGX-CONFIRMED @ `2c732e7`: 27B 235/235 + 35B 315/315 token-exact (FA2 sm_121a), memcheck 0/315**. **Item 5 model self-registration LANDED 2026-07-19** (`CLAIM-MODEL-SELFREG-1`, `MODEL-FACTORY-registry`): the fixed `kRegistrations` array → `REGISTER_VLLM_MODEL(...)` static-`Registrar` idiom (`model_registry.h:167-189`) + Qwen dense/MoE arch entry points split into per-variant TUs (`qwen3_5_dense.cpp`/`qwen3_5_moe.cpp`) over shared `qwen3_5_common.{h,cpp}`, so **adding a model = 1 new TU + 1 REGISTER line, ZERO shared-array edit**; behavior-preserving (clean CPU build + `test_model_registry` extension + full CPU CTest, tools 164/164, checkers green; DGX 27B/35B token-exact pending). Deep `qwen3_5.cpp` machinery factoring deferred. Metal/MLX bring-up proves the seams (needs M4). **★ THE ARCH HALF OF THIS ITEM IS NOW PROVEN BY MEASUREMENT, NOT ARGUED (2026-07-22, `CLAIM-CUDA-SM120-BRINGUP`, [spec §W8](specs/cuda-arch-additivity.md)):** consumer-Blackwell `sm_120a` was brought up as a BUILD-supported target and required **ZERO kernel, model, runner, sampler or feature-table edits** — the additive seams (per-arch FEATURE TABLE, capability-keyed tactic registry keyed on `major == 12`, Platform auto-probe, `pageable && integrated` residency classification) already covered it, so the entire diff is build configuration, a configure-tier test and records. That is the PR-#4 additivity test passed on a real second architecture. It is deliberately NOT a runtime-support claim: no sm_120 board exists here. **★ THE MODEL/QUANT HALF NOW ADVANCES TOO (2026-07-23, `CLAIM-BACKEND-SEAM-S4-1`): the `model_executor/layers/` `LinearMethod`/`QuantizationConfig` seam the [accelerator-seam audit](specs/accelerator-seam-audit.md) §9 named ABSENT now EXISTS in part.** `S4` landed byte-identical: the dense model's projections route through a `method.Apply()` chosen ONCE from the checkpoint (retiring the per-call `IsNvfp4()` tensor-name probe), and 18 shared-layer `device==kCUDA` availability gates became `vt::OpRegistered` op-table queries — the policy(scheme)/implementation(kernel) split the audit's binding rule requires. **DSR 86 → 67**; all six SACRED gates byte-identical on dgx (27B/35B/Coder/dense/OPT/DeepSeek-V2); the fragile 27B-W4A4/fp8-recipe gates are correctly deferred to `S6` behind `S5`'s reference tier. **★ `S6` ASSESSED 2026-07-23 (`CLAIM-BACKEND-SEAM-S6-1`) → NO-OP / BLOCKED, DSR stays 67 (§11):** the deferred fp4/fp8 gates convert ZERO sites byte-identically — every one bottoms out at a **dual-registered** (CPU+CUDA) bespoke op (none CUDA-only, unlike S4's convertible gates), so `OpRegistered(op,dev)` is TRUE on `kCPU` ⇒ the class-A swap is bit-changing on the CPU reference/emulation path (two numerics per device); S5's reference tier does not change this (those CPU kernels are present natively, never a miss). No `src/`/`include/`/test byte changed, no baseline moved. The genuine byte-identical unlock is re-scoped to **`S3`** (Platform capability fields mirroring `supports_fp8`/`cutlass_fp4_supported` — the audit's own class-D fix) and **`S7`** (layer extraction); the plan's `~37` S6 target assumed the class-A `OpRegistered` swap was byte-identical, which holds only for CUDA-only ops (all taken by S4). **★ `S3` LANDED 2026-07-23 (`CLAIM-BACKEND-SEAM-S3-1`) — the byte-identical unlock S6 re-scoped:** mirrors vLLM's `Platform` capability surface (`supports_fp8`/`cutlass_fp4_supported`/`opaque_attention_op`/`is_integrated_gpu`/`support_static_graph_mode`/`is_device_capability_family`, base false in `interface.h`; CudaPlatform answers GB10 values in `cuda.cpp`/`platform.cpp`) and converts **12** deferred `qwen3_5.cpp` gates onto it (7 fp4-act `cutlass_fp4_supported`, 3 fp8-fused `supports_fp8`, 2 decode-graph `support_static_graph_mode`) — byte-identical because a capability answers the base false off CUDA, exactly what `device==kCUDA` did (where S6's `OpRegistered` was TRUE on `kCPU`), and it DECOUPLES (a future accelerator answers for itself). **DSR 67 → 55** (`kcuda` 25→13), baseline lowered same commit, ratchet + 24-case suite green; all seven SACRED gates byte-identical on dgx (27B 235/235 · 35B 315/315 · Coder · dense-32B · OPT · DeepSeek-V2 · Llama), `test_platform` CUDA-leg proves each predicate == former `device==kCUDA` on GB10, memcheck 0 errors, clean CUDA+CPU `-Werror`. Residency/stream/FA2-dtype/merged-layout sites LEFT for `S7`. **★ `S7` LANDED 2026-07-23 (`CLAIM-BACKEND-SEAM-S7-1`) — the seam campaign's TERMINAL runtime-decoupling state, the closest this extensibility work comes to a finish line:** ALL 23 remaining runtime `kCUDA`/`is_cuda()` sites in the shared model layer hoisted onto capabilities — new `Platform::needs_weight_staging()` (the CUDA device-resident staging policy, NOT `is_unified_memory()` which would FLIP GB10; covers residency/merged-GDN/packed-decode/direct-load), `Platform::supports_fa2_attention()` (FA2 dtype), `Backend::SupportsAuxStream()` (MoE aux-stream), reusing S3's `supports_fp8`/`cutlass_fp4_supported`/`support_static_graph_mode`/`is_integrated_gpu` (runner combine/scatter) and `vt::OpRegistered(kMoeGroupedGemmBf16)` (a CUDA-only op). Each returns the former `device==kCUDA` value on GB10 → byte-identical. **DSR 55 → 32 — the IRREDUCIBLE build-gate floor:** the shared model layer holds ZERO runtime device tests; the 32 residual are all `#ifdef VT_*` compile-time gates for kernels that only build on one GPU family (a kernel that only compiles on one arch is legitimately irreducible), so the audit's `<10` is NOT reachable and this is the honest answer to "how additive can the shared layer get" — every runtime device coupling is gone. baseline lowered same commit, ratchet + 24-case suite green; all seven SACRED gates byte-identical on dgx (27B 235/235 · 35B 315/315 · Coder 6/6 · dense-32B 16/16 · OPT 6/6 · DeepSeek-V2 8/8 · Llama 16/16), new `test_platform`/`test_backend`/`test_cuda_backend` cases green, memcheck 0, clean CUDA+CPU `-Werror`. The `layers/`-library physical relocation (shrinking `qwen3_5.cpp` toward `qwen3_next.py`'s 802-line shape) is a follow-on refactor; the device coupling it was to remove is already gone | | 2 | `ROAD-V1-C2` | Model families: Llama/Qwen3/Mistral, MoE, Qwen3-Next | [model matrix](model-matrix.md) | current pin has 353 static IDs; v0.25.0 adds three sync-target rows (MOSS-Transcribe-Diarize, Laguna DFlash, Bailing hybrid MTP), yielding 356 after pin advance. **FIRST ADDITIVE-MODEL BRING-UP W0-W4 LANDED 2026-07-20 — CORRECTNESS COMPLETE (0.6B + 4B gates PASS 16/16; SPEED pending)** ([first-additive-model-qwen3-dense.md](specs/first-additive-model-qwen3-dense.md), `MODEL-TEXT-qwen3-qwen3-for-causal-lm` `ACTIVE`(correctness COMPLETE, speed pending), runner generalization `ENG-RUNNER-MODELSHAPE`, `CLAIM-MODEL-QWEN3-DENSE`) **MLA CAMPAIGN SPIKED 2026-07-21** ([mla-deepseek-campaign](specs/mla-deepseek-campaign.md), `CLAIM-MLA-DEEPSEEK`): five rows `INVENTORIED` -> `SPIKE` (DeepSeek V2 / V3+V3.2 / v1-MHA, Kimi-Linear, MiniMax-M2). **KIMI-LINEAR-48B W0 DEDICATED SPIKE 2026-08-05** ([kimi-linear.md](specs/kimi-linear.md), `CLAIM-KIMI-LINEAR-W0`): full dedicated W0 spike for `MODEL-TEXT-kimi-linear-*` (stays `SPIKE` — actively claimed) — the ONE Kimi text model that FITS one GB10 (91.5 GiB, 0.77x pool) with a real e2e SACRED gate; HEAVY reuse (MLA + sigmoid/`noaux_tc` MoE + GDN family + KDA host refs landed), NET-NEW = the KDA device kernel + NoPE-MLA branch + hybrid schedule/loader; W1 implementation can start. Answers the Tier-3 "MLA = new attention, new campaign" item in [breadth-sweep-plan](specs/breadth-sweep-plan.md) §B.3. Key determinations: GB10/sm_121 selects **`TRITON_MLA`** for dense MLA decode and **`FLASH_ATTN`** for MLA prefill (`vllm/platforms/cuda.py:129-133`, `mla/prefill/selector.py:74-77`), so the sm90/sm100-only MLA kernels are out of reach and out of scope; the cross-cutting cost is the **compressed-latent KV cache** (one 576-wide latent per token, `num_kv_heads=1`, no separate V), which our allocator and `vt::ReshapeAndCache`/`vt::PagedAttention` cannot express; and **only DeepSeek-V2-Lite (~29.3 GiB bf16) fits GB10** — V3/V3.2, Kimi-K2.5, MiniMax-M2/M3 are HW-BLOCKED e2e, Kimi-Linear-48B is HW-MARGINAL. W0-W10 plan recorded; nothing implemented. **GLM + DSA + LATEST-DEEPSEEK SPIKED 2026-07-21** ([glm-dsa-latest-deepseek](specs/glm-dsa-latest-deepseek.md), `CLAIM-GLM-DSA-LATEST-DEEPSEEK`): seven rows `INVENTORIED` -> `SPIKE` (ChatGLM, Glm, Glm4, Glm4Moe, Glm4MoeLite, GlmMoeDsa, DeepSeek-V4). Answers the user's "also glm, and deepseek latest versions" priority. Headline: **`Glm4MoeLiteForCausalLM` / `zai-org/GLM-4.7-Flash` (31.2B, 58.2 GiB bf16) FITS GB10 and is a SECOND MLA gate vehicle that closes BOTH coverage gaps the MLA campaign named as unit-gated-only** (it has `q_lora_rank=768` and `noaux_tc`/`e_score_correction_bias`, which DeepSeek-V2-Lite lacks). **DSA is DOUBLY BLOCKED on GB10:** the sparse XOR filter eliminates `TRITON_MLA` for sparse models leaving `FLASHINFER_MLA_SPARSE_SM120` as the sole candidate, and that path is non-functional on flashinfer 0.6.12 (XQA backend is dense-only, discards `sparse_mla_top_k`); GLM-5 is 1404 GiB and V3.2 642 GiB regardless. DeepSeek-V4 is a NEW architecture (Sinkhorn-normalized Manifold Hyper-Connections, CSA/HCA compressor with recurrent state, hash-routed MoE) and HW-BLOCKED at 148.7 GiB — but its TOKENIZER risk is LOW (standard HF fast BPE; only the chat template needs porting, with upstream golden fixtures). Glm4/Glm need two primitives we have NONE of: partial rotary factor and sandwich norms. Nothing implemented. **NVFP4A16 (W4A16)** on the already-done dense `Qwen3ForCausalLM` (`RedHatAI/Qwen3-32B-NVFP4A16`, 64L) — the QUANT-SCHEME additivity experiment, serving user priorities #2 (models) and #4 (quants) at once. KERNEL LAYER FULLY ADDITIVE (ZERO new kernel code: vLLM FORCES Marlin for `use_a16`, OBSERVED `Using MarlinNvFp4LinearKernel`, and that is the GEMM we already vendored for the 35B). **CORRECTNESS CLOSED 2026-07-21 (W4b):** the strict gate's 4/6 was diagnosed by the ratified TEACHER-FORCING isolation — all 29 divergent positions gap <= 0.0625 nats with 28/29 EXACTLY 0.0, one root flip an EXACT bf16 tie at which vLLM's teacher-forced argmax is OURS and vLLM contradicts its own greedy. **NOT a W4A16 defect: it is the pre-existing dense-forward bf16 near-tie drift, recorded against `MODEL-TEXT-qwen3-qwen3-for-causal-lm`.** Gate closes **6/6** under the ratified near-tie-robust bar with the nats evidence committed. SPEED still pending ⇒ `ACTIVE`, not `DONE`. **GEMMA FAMILY SPIKED 2026-07-24** ([sweep-gemma](specs/sweep-gemma.md), `CLAIM-SWEEP-GEMMA`): four rows `INVENTORIED` → `SPIKE` (Gemma 1/2/3/4). Answers the user's "and then we do gemma" ("gemma 4") next-target. **The newest registered Gemma is Gemma 4** (real, public, but ALL checkpoints multimodal-wrapped `Gemma4*ForConditionalGeneration`, ≥12B, 0.25.0 oracle-support unverified, needs a PLE/YOCO/MoE/k_eq_v stack) — it leads the characterization but is gate-BLOCKED as a first vehicle. **The recent-first gate vehicle that FITS + is oracle-certain is Gemma 3** (`Gemma3ForCausalLM` on `google/gemma-3-1b-it`). Headline: Gemma reduces MOSTLY to landed infra — gemma-RMSNorm, sandwich norms (glm4 `b568d20`), SentencePiece (names "Gemma"), sliding-window (FA-2 + SlidingWindow/ChunkedLocalAttention specs), the `kAttnQkNormRopeGate` QK-norm+rope recipe, tied embeddings are ALL REUSE; the one genuinely-new compute kernel is GeGLU (`gelu_pytorch_tanh`+mul; we have only SiLU), plus the final logit soft-cap + qpas/embed-scale scalars + dual-rope routing. Per-version delta: Gemma-2 has an attn logit soft-cap, Gemma-3 removed it for QK-norm. **GEMMA-3 W0-W2 LANDED 2026-07-24 — CORRECTNESS COMPLETE, the FIRST Gemma family** (`MODEL-TEXT-gemma3-gemma3-for-causal-lm` `ACTIVE`, speed pending): `Gemma3ForCausalLM` on `google/gemma-3-1b-it`. W1 = two additive default-inert vt ops `kGeluAndMul` (GeGLU `gelu_pytorch_tanh`) + `kMulScalar` (bf16 embed-scale), CUDA+CPU, unit 12/12. W2 = `gemma3.{h,cpp}`/`gemma3_weights.cpp`/`gemma3_registry.cpp` reusing the GLM-4 sandwich-norm layout + `dense_attn_block.h` + FA-only KV: GemmaRMSNorm `(1+w)`, per-head Gemma q/k norm, dual per-layer RoPE theta, qpas scale, per-layer sliding window, GeGLU MLP, `sqrt(hidden)` embed-scale, tied lm_head. **SACRED gate STRICT token-exact 48/48** greedy vs vLLM 0.25.0 (K=5 ALL-DETERMINISTIC → STRICT; BOS-verified; tokenizer-free like Mistral's `LOAD-SENTENCEPIECE` path). Loader 340 tensors, registry 23/23, clean `-Werror` 0 warn. **GEMMA-2 + GEMMA-1 W3-W6 LANDED 2026-07-24 — CORRECTNESS COMPLETE** (`MODEL-TEXT-gemma2-gemma2-for-causal-lm` + `MODEL-TEXT-gemma-gemma-for-causal-lm` `ACTIVE`, speed pending): W3 = the logit soft-cap primitives (`vt::SoftCap` final cap + `PagedAttentionArgs.logits_soft_cap` attention cap threaded into the native/CPU/FA-2 attention, default-inert). W4 `Gemma2ForCausalLM` (gemma-2-2b-it) = the inverse of Gemma-3 (BOTH soft-caps, NO QK-norm, single rope) — **near-tie-band SACRED 48/48** (44/48 strict + 4/48 at 0.0-nat vLLM-own ties, 0 forward-divergent; soft-cap PROVEN applied by a cap-on≠cap-off A/B). W5 `GemmaForCausalLM` (gemma-2b) = the original Gemma (two fused norms, head_dim scale) — **STRICT 48/48**. W6 = Gemma-4 honesty pass (HW/DEP-BLOCKED, not registered). Regressions byte-identical (Gemma-3 48/48, Qwen3-dense 184/184, OPT 63/63, Llama 92/92, Mistral 92/92) + DeepSeek-V2 asserts-on 223/223; compute-sanitizer 0; clean `-Werror` 0 warn. Gemma-4 stays `BLOCKED`. | `PARTIAL` | **ACTIVE: the first additive-model bring-up = Qwen3 dense (`Qwen3ForCausalLM`) on `Qwen3-0.6B` BF16 — W0+W1 landed 2026-07-20.** W0 (config+registry stub: new TU `qwen3_dense.cpp`+`qwen3.h`, one `REGISTER_VLLM_MODEL`, full-attention-only KV spec, forward stub) + W1 (the RUNNER GENERALIZATION `ENG-RUNNER-MODELSHAPE`) are DONE and gated: dgx CUDA `-Werror` 0-warn, **27B 235/235 + 35B 315/315 token-exact UNCHANGED** (behaviour-preserving), new CPU runner tests RED(SIGSEGV)→GREEN, registry resolves `Qwen3ForCausalLM`, ASan/UBSan clean. The runner is now MODEL-SHAPE-AGNOSTIC (a full-attention-only KV config allocates+steps without the hybrid GDN path) → every future dense/non-hybrid arch adds new-files-only. Qwen3-0.6B is the only standard-dense arch with a checkpoint + runnable vLLM 0.25.0 oracle on dgx TODAY (no Llama/Mistral checkpoint present → Llama-first needs a download, sequenced as W-next for genuine cross-family additivity). **W2 loader + W3 forward LANDED 2026-07-20:** dense forward `qwen3.cpp` (`Qwen3DenseModel::Forward/ForwardDevice`) composed from vt:: ops + the fusion catalog (2 new byte-exact recipes: `kFusedAddRmsNormStd`, `kAttnQkNormRope`); bf16 attention numerics mirror vLLM. The first pure-dense bf16 model forced out + FIXED 2 genuine latent bugs: tokenizer `kQwen2Classic` (classic Qwen2/Qwen3 pre-tokenizer was hard-rejected) and `cuda_paged_attn.cu` WMMA prefill mistokenizing at head_dim≠256 (now gated to the validated d=256). **W4 CORRECTNESS COMPLETE 2026-07-20 — near-tie-robust gate PASSES on 0.6B AND a bigger 4B.** The 2026-07-20 razor's "vLLM greedy non-deterministic" premise was a BATCHING artifact: per-prompt (batch=1, the gate regime) vLLM 0.25.0 greedy is DETERMINISTIC (0.6B 0-multi/K=10, 4B 0-multi/K=5). Forward correctness is PROVEN by teacher-forcing vLLM on OUR exact prefix (`scripts/qwen3-neartie-gap.py`): at all-but-2 positions vLLM's own argmax given our prefix IS our token (gap 0.0000, bit-identical logprobs — our forward matches vLLM's prefill logits); residual flips are bf16 near-ties (0.6B ≤0.125 nats, 4B ≤0.25) where vLLM's own prefill argmax disagrees with its decode. Gate = our token within 0.5 nats of vLLM's teacher-forced argmax (strict where equal): **Qwen3-0.6B 16/16** (strict 12 + near-tie 4) and the **bigger-model complete-correctness proof Qwen3-4B (36L, GQA 32/8, hidden 2560, same forward code) 16/16** (strict 10 + near-tie 6). Regression 27B 235/235 + 35B 315/315 UNCHANGED, CUDA `-Werror` 0-warn, memcheck 0. Correctness-complete. **SPEED — d128 FA2 PREFILL + DECODE DEFAULT-ON 2026-07-20 (`Qwen3-4B` vs vLLM 0.25.0 production/graphed, in1024/out128) — big gap-close, still below vLLM, `MODEL-TEXT-qwen3-qwen3-for-causal-lm` stays `ACTIVE` NOT `DONE`:** implemented the dominant prefill lever (a d128 FlashAttention-2 varlen prefill — generalized the vendored FA2 launcher to head_dim 128, `VT_FA2_PREFILL_QWEN3` default-ON) and flipped the FA2 varlen d128 decode default ON (near-tie gate re-passes 16/16 on 0.6B + 4B). Total tput now 0.90× (c1)/0.62× (c8) (up from 0.80×/0.48×), c1 decode at parity (TPOT 1.04×, ITL P99 0.98× win); prefill A/B = +7%/+41% total, −55%/−48% TTFT. STILL failing TTFT median 5.85×/10.2× + total <1×: the full prefill STEP (not the attention kernel, now vLLM's FA2 family) is ~6× vLLM = non-attention glue (GEMM/MLP fusion) + host-side launch overhead (un-graphed prefill); plus c8 decode batch efficiency (TPOT 1.38×). Dominant residual lever = portable prefill-step fusion + graphed prefill (secondary = c8 split-KV decode occupancy). **RoPE cos/sin cache flipped DEFAULT-ON 2026-07-20** (`VT_QWEN3_ROPE_CACHE`): the opt-in blocker (an alleged FA2-split-KV-combine run-to-run nondeterminism) was GROUNDED + DISPROVEN — the paged engine is byte-deterministic run-to-run and the combine never launches on the gate (`num_splits==1`); goldens regenerated on the canonical `$HOME/cutlass-4.5.0` build (the flashinfer cutlass copy tips the 27B tok6 razor to 234/235; cutlass-4.5.0 = 235/235), gate 16/16 both, 27B 235/235 + 35B 315/315 unchanged. RoPE-ON closes total tput 0.90×→0.97× (c1) / 0.62×→0.82× (c8) and c1 TTFT ratio 5.85×→2.27×, still `ACTIVE`. **SPEED RE-BOUND 2026-07-21 (same-session, matching-recipe) — TTFT residual RESOLVED, cutlass claim CORRECTED:** the "TTFT 2.27×/5.85×" + "c8 ITL 4.3×" were BAD-DENOMINATOR/num-prompts artifacts — a fresh same-session vLLM capture gives c1 TTFT ~152 ms & c8 ITL P99 ~130 ms, and OURS WINS TTFT at both concurrencies (c1 0.90×, c8 0.38×). **c1 = effective every-axis parity** (tput 0.98× / TPOT 1.01× / TTFT+ITL wins); **c8 residual = decode** (tput 0.93× / TPOT 1.10× / ITL P99 1.12×), which nsys shows is 93% GPU-busy/compute-bound (small-M=8 `cutlass_80_wmma` projections). The **qkv-merge** (new GQA `QkvSplit` op mirroring vLLM `QKVParallelLinear`) was implemented + MEASURED NEUTRAL (doesn't cut decode FLOPs) ⇒ default-OFF. **CUTLASS CLAIM CORRECTED: 27B `test_qwen27_paged_engine` = 235/235 on the FLASHINFER cutlass build** (the "flashinfer ⇒ 234/235" was a build artifact). Stays `ACTIVE`; named residual = c8 decode-GEMM efficiency (a decode-fusion sub-campaign). **SWEEP MODEL #1 — Qwen3-Coder-30B-A3B (`Qwen3MoeForCausalLM`) W0+W1 LANDED 2026-07-21** ([sweep-qwen3-coder-30b.md](specs/sweep-qwen3-coder-30b.md), `MODEL-TEXT-qwen3-moe-qwen3-moe-for-causal-lm` `ACTIVE`, `CLAIM-MODEL-QWEN3-CODER`): the first full-attention BF16 MoE, composed from the done dense attention + the done 35B MoE experts (ZERO runner change). W0 = registry stub (`qwen3_moe_registry.cpp` + `qwen3_moe.h`, one `REGISTER_VLLM_MODEL`, full-attn-only KV, `is_dense_model=false`, W2/W3 throwing stubs). W1 = three behaviour-preserving refactors making the two done pieces reusable: (#1) dense `AttnBlock` + glue extracted to `dense_attn_block.h` (Qwen3-dense byte-identical), (#2) bf16 `MoeBlock` exposed cross-TU via `RunMoeBlock` (`qwen3_5_moe_block.h`; 35B untouched), (#3) no-shared-expert guard (inert for the 35B). Gated: dgx CUDA `-Werror` 0-warn; Qwen3-dense 0.6B+4B 16/16 + 27B 235/235 + 35B 315/315 UNCHANGED; registry resolves `Qwen3MoeForCausalLM`; memcheck 0. W2 bf16 loader → W3 forward → W4 near-tie token-exact → W5 fast bf16 grouped-MoE GEMM remain. Then Llama dense (download), Mistral, MoE families **SWEEP MODEL — GLM-4-9B-0414 (`Glm4ForCausalLM`) G2 LANDED 2026-07-24 — CORRECTNESS COMPLETE** ([glm-dsa-latest-deepseek](specs/glm-dsa-latest-deepseek.md), `MODEL-TEXT-glm4-glm4-for-causal-lm` now `READY` per the [live-state audit](specs/live-state-audit-2026-08-06.md), `CLAIM-GLM-DSA-LATEST-DEEPSEEK` amended) **GLM-4.7-Flash (`Glm4MoeLiteForCausalLM`, 31.2B MLA+MoE) G1 LANDED 2026-07-24 — SACRED gate 8/8, CORRECTNESS COMPLETE** (`MODEL-TEXT-glm4-moe-lite-glm4-moe-lite-for-causal-lm` `ACTIVE`, speed pending): reuses the DeepSeek-V2 MLA stack; first e2e coverage of the q_lora branch + noaux_tc sigmoid router, closing the MLA campaign's two C2 coverage gaps: the FIRST GLM-family model. SACRED gate 16/16 vs vLLM 0.25.0 (STRICT 13/16 + near-tie 3/16, max gap 0 nats; vLLM K=5 self-deterministic ⇒ STRICT bar), speed PENDING. The two "new primitives" the spike named reduced to EXISTING infra: partial + interleaved `RopeFromCache` (`is_neox_style=false`, the DeepSeek-V2 decoupled-rope path) over `rotary_dim=64`; standalone `vt::RmsNorm` sandwich norms. Biased qkv, no QK-norm, GQA 32/2, untied lm_head. New files + one REGISTER, reusing the shared dense glue. **SWEEP MODEL — Llama-3.2 (`LlamaForCausalLM`) W0-W4 LANDED 2026-07-23 — CORRECTNESS COMPLETE** ([sweep-llama-3.2](specs/sweep-llama-3.2.md), `MODEL-TEXT-llama-llama-for-causal-lm` `ACTIVE`, `CLAIM-MODEL-LLAMA-3.2`): the roadmap's explicit "Llama-first" increment and the first mainstream non-Qwen/non-OPT dense arch. `LlamaForCausalLM` (Llama-3.2-1B) = the Qwen3-dense forward with exactly two ADDITIVE deltas — NO qk-norm (shared `AttnBlock` skips it when q_norm/k_norm empty) + llama3 rope-scaling (4 default-0 `RopeArgs` fields + a `Llama3ScaleFreq` kernel helper, no-op elsewhere; formula verified 2e-7 rel vs vLLM) — reusing the shared dense forward VERBATIM (`LlamaModel == Qwen3DenseModel`). 3 new Llama files, ZERO edit to runner/scheduler/platforms/attn-registry/`hf_config`/any qwen3-opt model. vLLM 0.25.0 greedy MEASURED DETERMINISTIC (K=6, 0 multi-valued cells) ⇒ STRICT token-exact bar, PASS **16/16 (12 strict + 4 near-tie band, max gap 0.0000 nats, 0 divergent)** — at all 13 divergent positions vLLM's own teacher-forced argmax given our prefix IS our token. A correctness-fatal tokenizer bug (Llama's `Sequence` post_processor wrapping `TemplateProcessing` ⇒ BOS 128000 never prepended, silently 1/16) was isolated by a CUDA prefill-argmax diagnostic (forward proven 4/4 correct given vLLM's exact tokens) and fixed byte-preservingly (Qwen/OPT/DeepSeek unaffected — all ByteLevel/top-level-TemplateProcessing). Regressions 27B 235/235 · 35B 315/315 · Coder 6/6 · Qwen3-dense 16/16 · OPT 6/6 · DeepSeek-V2 8/8 UNCHANGED; `-Werror` 0-warn; memcheck 0; DSR 67. SPEED pending (head_dim 64 → generic paged path; Llama-3.2-3B head_dim-128 is the FA2-toggle W-next). **MLA CAMPAIGN W0+W1 LANDED 2026-07-21** (`CLAIM-MLA-DEEPSEEK`; rows STAY `SPIKE` — W0/W1 make no model supported). **W0 grounded every fact the spike flagged as an unverified source read; ALL CONFIRMED, none contradicted:** DeepSeek-V2-Lite fetched to dgx (30 GB, 4 shards) and loading in the vLLM 0.25.0 oracle; the oracle's own DEBUG startup on sm_121 prints `Using TRITON_MLA attention backend out of potential backends: ['TRITON_MLA']` and `Using FLASH_ATTN MLA prefill backend` — so the dense-MLA decode + MLA-prefill targets are OBSERVED, not inferred, and the sm90/sm100-only MLA kernel class stays out of scope; the real `config.json` confirms every §5.1 number (`kv_lora_rank=512`, `qk_nope=128`, `qk_rope=64` -> the **576-wide latent**, `v_head_dim=128`, `q_lora_rank=null`, `n_group=topk_group=1`, softmax/greedy, 64+2 experts, 27 layers) plus `is_neox_style=False` and the mscale2 scale correction; and BOTH recorded coverage gaps (no `fused_qkv_a_proj` branch, no `e_score_correction_bias`) are confirmed real, so those pieces stay unit-gated only. **W1 = the behaviour-preserving spec-driven KV allocation, ZERO MLA math:** the attention cache is now sized `num_blocks * spec->page_size_bytes()` and viewed from the spec's own `block_size`/`num_kv_heads`/`head_size`/`dtype` instead of the hardcoded `2 * block * Hkv * Dh` with shape reconstructed from the HF config (`runner.cpp`), plus `MLAAttentionSpec` with upstream's factor-1 single-tensor page formula (`kv_cache_interface.py:397-398`) registered against the ORDINARY `FullAttentionManager` (`single_type_kv_cache_manager.py:1539`) — the spike's key finding, which is why block manager/prefix caching/eviction need no change. Gated: dgx clean CUDA `-Werror` 0 warnings/0 errors; **27B 235/235 + 35B 315/315 + Qwen3-Coder 6/6 + Qwen3-dense 16/16 ALL UNCHANGED** (behaviour-preserving proven, not assumed); `test_runner` 15/15, `test_kv_cache_interface` 21/21 (4 new MLA-spec cases), `test_llm_engine` 5/5; the new path is proven EXERCISED (not merely compiled) by `fa_page_size_bytes()` + a `page_size_padded` case no HF-config formula can produce. **W2 + W3 LANDED 2026-07-21** (base `a05437f`; rows STAY `SPIKE` — still no MLA attention math, no MLA model, no forward). **W2 = the MLA branch of `_get_backend_priorities` the pre-W2 comment deferred, ported as DATA:** the whole of `cuda.py:84-176` (BOTH branches — MLA sm_100 including the `:96-115` adaptive sparse tail, MLA sm_12x, MLA `else`, and the two pre-existing non-MLA arms) is now a TABLE in the new header `include/vllm/platforms/cuda_attn_priority.h`, one row per upstream arch arm keyed on `(use_mla, major)`, so a future arch is a ROW rather than a code path; putting it in a header (not the CUDA-only TU) also let the CPU test tier assert the REAL table and DELETED the hand-copied `FakeCudaPlatform` duplicate. On sm_121 a `use_mla=true` request now RESOLVES to `TRITON_MLA`, matching the W0 oracle observation. **The sparse/DSA seam is left OPEN and unit-proven:** GB10's row keeps both upstream entries and the sparse one loses to a real FILTER — `AttentionBackend::is_mla()`/`is_sparse()` checked against the request (`backend.py:307-360 validate_configuration`) — so a future DSA backend is selected purely by declaring `is_sparse() == true`, with ZERO edit to the table or the selector. `TritonMLABackend` lands the NAME plus upstream's 3-D `get_kv_cache_shape` (no K/V axis; `num_kv_heads != 1` REFUSED), `get_impl_cls()` deliberately still `nullptr`. MLA prefill priority ported too (GB10 -> `[FLASH_ATTN]` alone). **W3 = the two new `vt::` ops, both CPU-reference-gated.** `vt::ConcatAndCacheMla` mirrors `csrc/libtorch_stable/cache_kernels.cu:401-442` — and per the whole-chain rule this was VERIFIED, not assumed, to be vLLM's OWN csrc kernel (`_custom_ops.py:2532` -> `torch.ops._C_cache_ops`), with no flashinfer/cutlass variant in the dense-bf16 path; it concatenates the latent + rope part into ONE 576-wide entry, the write `ReshapeAndCache`'s K/V-pair signature cannot express, stride-driven so a per-layer cache slice and the two column halves of `kv_a_proj_with_mqa` both work copy-free. The **grouped-topk (`noaux_tc`) router** extension — flagged in `coordination.md` as SHARED with `CLAIM-GLM-DSA-LATEST-DEEPSEEK` and "must not be implemented twice" — is landed HERE and that claim now consumes it: additive `MoeRouterTopKArgs` fields + an optional `e_score_correction_bias`, with `num_expert_group == 0` still dispatching the ORIGINAL kernel so the 27B/35B/Coder/dense routers are byte-identical BY CONSTRUCTION. **Stated plainly: the `noaux_tc` correctness evidence is UNIT-ONLY.** V2-Lite has `n_group=topk_group=1` and no bias, so the e2e vehicle exercises none of it; the gate is `tests/vt/test_ops_moe_router_grouped.cpp` at DeepSeek-V3's REAL dimensions (256 experts, n_group=8, topk_group=4, sigmoid, scaling 2.5, WITH the bias) against an INDEPENDENT sort-based transcription of the upstream formula. **W4 LANDED 2026-07-22** (base `ed2c342`; rows STAY `SPIKE` — W4 adds a kernel and fills a `nullptr`, it makes no model supported). **`vt::MlaDecodeAttention` — the MQA decode over the compressed latent (QK 576 / V 512, `num_kv_heads=1`), a structure port of the two-stage split-KV pair W0 OBSERVED EXECUTING:** `MlaDecodeStage1` <- `_fwd_grouped_kernel_stage1` (`triton_decode_attention.py:278-458`, the `IS_MLA` branch whose `v = tl.trans(k)` at `:424-431` is the whole MLA trick — V is the leading 512 columns of the SAME latent row already loaded as K, so one shared-memory tile serves as both), `MlaDecodeStage2` <- `_fwd_kernel_stage2` (`:575-639`), `ComputeNumKvSplits` <- `_compute_num_kv_splits` (`triton_mla.py:40-47`), split workspace <- `_reserve_attn_logits_workspace` (`:57-78`) realized as the house grow-only per-stream scratch. **Honest reuse verdict:** our FA-2 split+combine machinery fit at the ALGORITHM level (the split schedule, the LSE merge algebra, the fixed-ascending no-atomicAdd determinism rule) and NOT at the code level — the vendored FA-2 launcher takes separate 4-D k/v caches and is instantiated for symmetric head_dim {128,256}, which cannot express a 3-D single-buffer cache with QK 576 / V 512; that is recorded in the TU header rather than forced. **Evidence is unit-level and deliberately strong** (there is no e2e model until W7): [`tests/vt/test_ops_mla_attn.cpp`](../tests/vt/test_ops_mla_attn.cpp), a port of `tests/kernels/attention/test_mla_decode_cpu.py` whose `ref_mla` becomes an INDEPENDENT TWO-PASS oracle (a different algorithm from the streaming online-softmax both impls use) plus its NaN-padding out-of-bounds detector, run at the REAL V2-Lite geometry (576/512/64, block 16, mscale^2 scale) over ragged / multi-block / single-block / every split boundary (`num_kv_splits` ∈ {1..512} incl. splits > seq_len) / 128-head V3 / non-BLOCK_H head counts / a 288-256 block-32 geometry / bf16 + f32, with run-to-run BIT-exactness. dgx sm_121: 11/11 cases, 2,303,193 assertions; `compute-sanitizer` memcheck **0 errors**, racecheck **0 hazards**, synccheck **0 errors**; clean CUDA build **0 warnings / 0 errors**; **27B 235/235 + 35B 315/315 + Coder 6/6 + Qwen3-dense 16/16 + OPT 6/6 ALL UNCHANGED**. `TritonMLABackend::get_impl_cls()` is no longer `nullptr` — it returns a real `TritonMLAImpl` whose `forward_mqa` is the 1:1 counterpart of `triton_mla.py:189-260`; PREFILL is W5 and `forward()` refuses a prefill-shaped batch BY NAME rather than producing wrong numbers. NO speed number — decode perf is W9. **W5 LANDED 2026-07-22** (base `5395203`; rows STAY `SPIKE`). **MLA PREFILL + the workspace-bounded CHUNKED-CONTEXT loop.** Three new ops — `vt::MlaPrefillAttention` (<- `mla/prefill/flash_attn.py:153-248`, the ONLY MLA prefill backend reachable on sm_121 and the one W0 OBSERVED the oracle logging), `vt::GatherMlaCache` (<- `cache_kernels.cu:992-1064`) and `vt::MergeAttnStates` (<- `merge_attn_states.cu:18-192`, both `-inf` edge cases verbatim) — plus the loop itself in the new `mla_chunked_context.h` (<- `mla_attention.py:1422-1451,1667-1745,2094-2199,2344-2425`), which is what keeps a long-context prefill inside a bounded workspace instead of materializing a 3 GB up-projected context. **The vendored FA-2 launcher WAS generalized, and W4's prediction that it would be tractable held for a reason worth recording: upstream does not ask FA-2 for asymmetric head dims either.** `requires_v_padding` is TRUE on GB10, so upstream ZERO-PADS V from 128 to 192 and slices the output back — the kernel stays a plain SYMMETRIC head_dim-192 instantiation. The whole change is two new explicit instantiations of an UNCHANGED generic template, one new launcher entry for the contiguous-varlen mode, and the pad/slice pair; the paged launcher every non-MLA prefill calls is textually untouched (211 insertions / **0 deletions**), and 27B 235/235 + 35B 315/315 + Coder 138/138 + Qwen3-dense 664/664 + OPT 36/36 are all UNCHANGED. Evidence is UNIT-ONLY (there is still no model): 4/4 cases / **2,377,052 assertions** and 5/5 / **306,037 assertions** on dgx sm_121 at the real QK 192 / V 128 geometry, against an INDEPENDENT double-precision two-pass oracle and — for the loop — a SINGLE-SHOT whole-sequence oracle that never chunks, over exact / +1 / -1 chunk boundaries, zero-context and zero-key-in-chunk requests, ADVERSARIAL reverse-interleaved block tables, NaN-poisoned outputs and run-to-run bit-exactness; memcheck **0**, racecheck **0 hazards**, synccheck **0**. A genuine upstream FA-2 quirk was found on the way and worked around rather than papered over: its EMPTY-K early exit ignores the unpadded-LSE flag, which a zero-key chunk request would turn into an out-of-bounds LSE write. **W6 LANDED 2026-07-22** (base `2846467`; rows STAY `SPIKE` — W6 adds an attention LAYER, not a model). **The MLA attention BLOCK + LOAD-TIME WEIGHT ABSORPTION — the piece that finally COMPOSES W3's cache write, W4's MQA decode and W5's MHA prefill into one layer:** the projections with BOTH `q_lora_rank` branches (`fused_qkv_a_proj` -> `q_a_layernorm` -> `q_b_proj`, or the direct `q_proj`), the two RMSNorms (the rope part deliberately NOT normed), the DECOUPLED RoPE (`is_neox_style=False`, only the trailing 64-dim slice rotates) with its YaRN cos/sin cache and the SEPARATE mscale^2 softmax-scale correction, the `kv_b_proj -> W_UK/W_UV` split, the prefill-MHA / decode-MQA dispatch with decode tokens packed FIRST, and the `kv_b_proj` up-projection callback W5 left open. **The spike's most useful prediction held: absorption needed NO new attention kernel** — it is a LOAD-TIME weight transform plus TWO batched GEMMs, so the entire new-kernel surface is two general primitives, `vt::BatchedMatmul` (<- `torch.bmm` at `mla_attention.py:789,1034`; on CUDA torch resolves that to cuBLAS `gemmStridedBatchedEx`, and ours is the cuBLASLt strided-batched form of the same GEMM) and `vt::ConcatMlaNopeRope` (<- `concat_mla_q`, generalized so one op also serves `_concat_k_nope_k_pe`). **The absorbed-vs-unabsorbed equivalence — the heart of W6 — is PROVEN NUMERICALLY, three independent ways, rather than argued:** an INDEPENDENT double-precision block oracle computes the attention BOTH ways and agrees to < 1e-11 (the identity itself); our absorbed decode reproduces the UNABSORBED oracle to < 2e-4 in f32; and — the strongest — the SAME batch is driven once through our ABSORBED MQA decode kernel (QK 576 / V 512, one KV head, K/V never materialized) and once through our UNABSORBED materialized-MHA prefill path (QK 192 / V 128 plus the chunked-context loop), agreeing to < 3e-4 (CPU f32) / < 4e-2 (CUDA bf16) with nothing but the weights shared between them. Evidence on dgx sm_121: `test_mla_attention_block.cpp` 10/10 cases / 2,372,644 assertions and `test_ops_mla_absorb.cpp` 9/9 / 1,644,807 (CUDA cases proven to EXECUTE; NaN-poisoned outputs; run-to-run BIT-exact), porting `tests/kernels/test_concat_mla_q.py` in both arms. memcheck 0, racecheck 0 hazards, synccheck 0 (the last needing `--num-cuda-barriers 65536` — the default table overflows on a binary driving this many kernel families and the tool then reports a bogus launch failure, a trap worth knowing). Clean CUDA build 0 warn / 0 err; **27B 235/235 + 35B 315/315 + Coder 138/138 + Qwen3-dense 664/664 + OPT 36/36 ALL UNCHANGED**. **Coverage stated plainly: the `q_lora` query branch has NO e2e coverage and cannot get any on GB10** — DeepSeek-V2-Lite has `q_lora_rank=null`, so it is unit-gated at DeepSeek-V3's real dimensions only; GLM-4.7-Flash (`q_lora_rank=768`, 58.2 GiB, fits) is what would close it. **W7 LANDED 2026-07-22** (base `ce43c51`; the row STILL stays `SPIKE`). **THE DEEPSEEK-V2 MODEL — registry + config parse + loader + forward: the first MLA model in the tree, and the first one that runs a real MLA checkpoint end to end.** Four new files plus ONE shared-code edit (a two-line additive condition in `runner.cpp` recognising a `kMlaAttention` KV group as the model's attention group — upstream maps MLA onto the ordinary `FullAttentionManager`, so block tables/prefix caching/eviction are untouched). **LOADER GATE PASSED on the real 4-shard DeepSeek-V2-Lite: 5291/5291 checkpoint tensors accounted for, ZERO unmapped and ZERO leftover** (4/4 cases / 37,331 assertions), every shape asserted including the LOAD-TIME `kv_b_proj -> W_UK_T [16,128,512]` / `W_UV [16,512,128]` absorption split — the same transform, at the same lifecycle point, as upstream's `process_weights_after_loading`. **V2-Lite takes the DIRECT `q_proj` query branch** (`q_lora_rank: null`), asserted with the fused branch EMPTY on every layer. **FORWARD GATE PASSED and obviously right, not merely finite: the real checkpoint prefill of `The capital of France is` -> argmax ` Paris`** (top-5 ` Paris`/` the`/` a`/` one`/` also`, run-to-run bit-exact) — the direct analogue of the Qwen3-Coder W3 sanity case. **BATCH-ORDERING GATE:** the ordering invariant W6 measured 0.86 relative error from is now VALIDATED, not assumed — `BuildMlaBatchSplit` throws (naming the request and citing the upstream line) if a decode follows a prefill or a with-context prefill follows a context-free one. **SHARED EXPERTS — new for this family and UNGATED unlike Qwen3.6's sigmoid-gated one — gated two ways:** a MoE layer with every routed expert zeroed is BIT-IDENTICAL to a dense layer holding the same MLP, and turning the shared expert off CHANGES the logits. **The CUDA path is EXERCISED, not merely compiled:** a case at the real MLA head geometry drives the CUDA MLA kernels and the CUDA-only grouped bf16 MoE GEMM, bit-exact on device and within 0.0061 worst relative logit error of the CPU reference path. 11/11 forward cases; memcheck/racecheck/synccheck all **0**; clean CUDA build **0 warn / 0 err**; **regression set UNCHANGED**. **Only `DeepseekV2ForCausalLM` is REGISTERED** — `DeepseekForCausalLM` (plain MHA), V3 (fp8/671B) and V3.2 (DSA indexer) are REFUSED BY NAME in the config parse rather than falsely claimed. A pre-existing tree-wide hazard was found on the way and recorded: the shared `DevicePool` is a process-wide singleton keyed only on a byte size class, so a single process driving BOTH a CPU and a CUDA forward hands the second backend the first's recycled pointers. **NEXT: W8 — the SACRED token-exact gate on DeepSeek-V2-Lite** (wire the paged engine to produce the MLA batch order the model already validates, capture oracle goldens, run the STRICT form W0 determined). A loading, forwarding model is NOT a supported model, so no model row moves until that gate passes. **W8 LANDED 2026-07-22 — THE SACRED CORRECTNESS GATE PASSES 8/8, and `MODEL-TEXT-deepseek-v2-deepseek-v2-for-causal-lm` moves `SPIKE` -> `ACTIVE` (correctness COMPLETE, speed PENDING). NOT `DONE` — that additionally requires vLLM-speed parity on every axis, which is W9 and has NO number yet; the other four campaign rows stay `SPIKE`.** An 8-prompt battery is driven through the FULL paged `LLMEngine` and compared to the pinned vLLM 0.25.0 oracle: **8/8 PASS — STRICT token-exact 5/8, near-tie band 3/8, 92/128 tokens strictly exact, max teacher-forced gap 0.25 nats, 0 forward-divergent** (223 assertions). **The bar was ARRIVED AT by measurement, not chosen:** vLLM is DETERMINISTIC on this model at batch=1 (W0's K=5 8/8, re-confirmed by W8's own capture at T=16 with 0 multi-valued cells), so the STRICT form ran FIRST and came out 5/8; the ratified TEACHER-FORCING diagnostic then showed **36 divergent positions with 35 at gap EXACTLY 0.0000 nats** — vLLM's own argmax GIVEN OUR PREFIX is our token, so they are the downstream tail of one earlier flip — **exactly ONE root flip with any gap at all (prompt[3] tok 9, 0.2500 nats, inside the ratified 0.5-nat band and equal to the landed Qwen3-dense 4B gate's worst)**, and **ZERO tokens outside vLLM's top-20**, with the per-position nats COMMITTED as goldens and anything beyond the band still FAILING. **W8's first job — the scheduler/runner wiring — turned out to need NO new code, for a non-accidental reason:** `runner.cpp:671` already reorders with `decode_threshold = 1`, exactly MLA's `reorder_batch_threshold` (`mla_attention.py:1420`), and its `decode -> short_extend -> long_extend -> pure_prefill` ordering satisfies BOTH MLA invariants (decodes form a batch prefix; with-context prefills lead the prefill tail). W8 PROVES that end to end rather than duplicating it, with new DIAGNOSTIC `MlaBatchSplitStats` counters and a non-vacuity bar: the battery is admitted CONCURRENTLY with staggered arrival, producing **7 genuinely MIXED decode+prefill steps at up to 8 concurrent requests** with `BuildMlaBatchSplit` (which throws naming the request) never firing, plus a prefix-cache-driven **with-context prefill**, and a phase-0 check that the engine really allocated the compact MLA cache (`fa_page_size_bytes = 36864`, no factor 2). **THE REAL BLOCKER WAS THE TOKENIZER, NOT THE MODEL:** the first run REFUSED to load (`unsupported normalizer "Sequence"`), and behind it sat a whole NEW pre-tokenizer family — DeepSeek's is a HF `Sequence` PIPELINE of SEVEN stages (five `Split(Isolated)` over ENUMERATED codepoint ranges, then `Digits(individual_digits=true)`, then `ByteLevel(use_regex=false)`), whose stage ORDER is load-bearing because stage 2's punctuation class spans 0x3A-0x7E and CONTAINS A-Z/a-z. Landed as `SplitPattern::kDeepSeek` with the five patterns compared VERBATIM at load, and MEASURED token-for-token against the REAL HF `tokenizers` library over a stage-stress corpus (**6/6 cases / 2461 assertions**). **The TOKENIZATION goldens earned their keep by REFUTING a fix that was already written:** `tokenizer_config.json` declares `add_bos_token: true`, which reads as exactly the OPT missing-BOS bug — but vLLM's resolved tokenizer (`TokenizersBackend`) adds NO BOS, our loader already matched bit-for-bit, and the "fix" would have BROKEN a passing gate; it was reverted and the measured behaviour PINNED by a guard case ([[ground-premises-before-dispatching]]). Regression set UNCHANGED (27B 235/235, 35B 315/315, Coder 6/6, Qwen3-dense 16/16, OPT 6/6, plus every tokenizer test — W8 touches SHARED tokenizer code, so that was proved, not assumed); clean CUDA rebuild 0 warn/0 err; local CPU suite 151/151; memcheck/racecheck/synccheck 0. Batch invariance is REPORTED (6/8) and deliberately NOT a bar, because the ORACLE itself changed on 3/8 of this battery under batched generation (W0). One W9 input recorded: the oracle must run `moe_backend='triton'` — vLLM's auto-selected FlashInfer CUTLASS unquantized MoE REBOOTED dgx three times on GB10's unified memory. **W9 SPEED CLOSE LANDED 2026-07-22 — the track has its FIRST binding speed number, and it is an ATTRIBUTED MISS: `MODEL-TEXT-deepseek-v2-deepseek-v2-for-causal-lm` STAYS `ACTIVE` (correctness COMPLETE, speed SHORT), NOT `DONE`** ([grid](../docs/BENCHMARKS.md), [spike §W9](specs/mla-deepseek-campaign.md)). Denominator SETTLED with evidence — CUTLASS MoE has now rebooted dgx **five times** (two more at W9, the second on a pristine box with a 0 GiB page cache and every mitigation applied, both deaths at the identical post-`torch.compile` phase), so `--moe-backend triton` IS vLLM's best STABLE GRAPHED configuration here and is the bar; the substitution does not flatter us, we lose to it. `nsys` (both sides, `--cuda-graph-trace=node`) overrode the plan: the lever was not the planned MLA fusion recipes but `MlaDecodeStage1` sitting at **44.7% of all GPU time and ~180x off its own memory-bound floor** on a **2-CTA grid at batch 1**; applying upstream's own occupancy target made it **18.3x faster** (837 -> 45.8 us) for **+69.5%/+53.3%/+32.0%/+19.5%** end-to-end at c1/c2/c4/c8, while the planned decode-graph sibling is worth only ~+2% (this decode is GPU-bound). Grid vs vLLM: output throughput **0.87/0.95/0.86/0.88** (was 0.50 at c1), TTFT **1.06/1.14/0.96/0.88** (we WIN at c4/c8), TPOT **1.11/0.97/1.16/1.17**. SACRED gate **8/8 UNCHANGED** with both levers default-ON; a real latent CUDA-graph use-after-free in the MLA metadata upload was found and fixed (its whole class now guarded); regression set UNCHANGED; clean rebuild 0 warn/0 err; sanitizers 0. **NEXT LEVER, NAMED: route the batch-1 dense projections off cuBLAS `gemvx` (31.8% of our GPU time) onto a tensor-core GEMM — vLLM splits the same work `gemvx` 12.7% + `nvjet_sm121_tst_mma_*` 6.6%.** **W10 BLOCKED-ROW HONESTY PASS LANDED 2026-07-22 — the campaign's W-plan is COMPLETE; records only (no code, no build, no GPU work, nothing downloaded, no number claimed).** Rows set to their final honest state: `MODEL-TEXT-deepseek-v2-deepseek-v3-for-causal-lm` (V3 + V3.2, and Kimi-K2/K2.5's text backbone by config composition) and `MODEL-TEXT-minimax-m2-mini-max-m2-for-causal-lm` move `SPIKE` -> `BLOCKED`, joined cross-claim by `MODEL-TEXT-deepseek-v2-glm-moe-dsa-for-causal-lm` (GLM-5) under `CLAIM-GLM-DSA-LATEST-DEEPSEEK`; each is HW-BLOCKED on 119 GiB (~642 GiB fp8 / ~428 GiB / 1404 GiB) and the two DSA models are additionally DEP-BLOCKED — for a SPARSE model the XOR filter eliminates `TRITON_MLA`, leaving `FLASHINFER_MLA_SPARSE_SM120` alone, whose sm12x dispatch goes to flashinfer's DENSE-ONLY XQA backend that discards `sparse_mla_top_k` (upstream's own test monkeypatches the probe and asserts nothing numerical). `MODEL-TEXT-deepseek-v2-deepseek-for-causal-lm` stays `SPIKE` with the record repaired to say it is plain MHA and needs NO MLA; Kimi-Linear stays `SPIKE` (MLA half unlocked, KDA a separate kernel campaign, HW-MARGINAL). Each blocked row states what CAN still be gated (config resolution, weight-map on a slice, unit parity at real dimensions) versus what CANNOT (anything e2e). **Two PERMANENT coverage gaps now stated in the rows:** the `noaux_tc` grouped router and the `q_lora` query branch have NO e2e coverage and are unit-gated only, because V2-Lite is `n_group=topk_group=1`/softmax with no `e_score_correction_bias` and `q_lora_rank=null`. **NAMED NEXT VEHICLE: GLM-4.7-Flash** (`MODEL-TEXT-glm4-moe-lite-glm4-moe-lite-for-causal-lm`, 31.2B / 58.2 GiB, FITS GB10) — the only reachable checkpoint that closes BOTH gaps. **BLOCK NOT CLOSEABLE, nothing archived:** the DeepSeek-V2 row is `ACTIVE`, not `DONE`, so the plan/spec stay LIVE; the single open item is the named `gemvx` -> tensor-core dispatch lever. **MISTRAL FIFTH FAMILY W0-W3 LANDED 2026-07-23** ([sweep-mistral](specs/sweep-mistral.md), `MODEL-TEXT-mistral-mistral-for-causal-lm` `ACTIVE`, `CLAIM-MODEL-MISTRAL`): the closest-to-Llama dense arch (vLLM `mistral.py` = "Mistral adaptation of the LLaMA architecture") — plain rope θ1e6 (no rope_scaling) + qk-norm-optional + untied lm_head + null sliding_window, all PRE-EXISTING ⇒ NO new primitive, 3 new files + additive CMake/registry-test rows only, ZERO shared-code edit. **MODEL forward gate 30/30 greedy tokens vs vLLM 0.25.0** (tokenizer-free: fed vLLM's exact prompt tokens through our CUDA prefill; 29 STRICT token-exact + 1 near-tie, 0 forward-divergent; vLLM greedy det 4/5 K=3). W2 loader real-weights 1541 assertions. **REAL FINDING:** Mistral's SentencePiece/Metaspace tokenizer is unsupported by our ByteLevel-BPE tokenizer → the FULL paged-engine SACRED gate is BLOCKED, the pre-inventoried `LOAD-SENTENCEPIECE` row (SentencePiece tokenizer family). `-Werror` 0-warn, DSR 32, regressions UNCHANGED (Llama paged 16/16, Qwen3-dense forward 1031, registry 299; MoE/GDN gates unaffected by construction). SPEED + full paged gate both PENDING (row `ACTIVE`, not `DONE`). **OLMo-2 SPIKED 2026-07-24** ([sweep-olmo2](specs/sweep-olmo2.md), `CLAIM-SWEEP-OLMO2`): one row `INVENTORIED` → `SPIKE` (`MODEL-TEXT-olmo2-olmo2-for-causal-lm`, covering `Olmo2ForCausalLM` + its `Olmo3ForCausalLM` alias). Answers the breadth-sweep §B.3 Tier-2 rank-8 "GLM4 / Olmo2-3" item (GLM-4 + Gemma landed; OLMo-2 next). **HEADLINE: OLMo-2 is the cleanest dense bring-up yet — ZERO new compute kernels.** The two distinctive facts both reduce to WIRING over landed ops: (1) the **pure post-norm (`norm_after`) placement** is a strict SUBSET of the GLM-4/Gemma sandwich (keeps ONLY the standalone-output-norm op `glm4.cpp:174-178` — the exact primitive flagged — DROPS the pre-norms, plain residual add); (2) the **QK-norm is FULL-WIDTH not per-head** → reuses `vt::RmsNorm` at a `[T,q_size]`/`[T,kv_size]` shape but CANNOT use the fused per-head `kAttnQkNormRopeGate`. Everything else REUSES (plain RMSNorm, SiLU SwiGLU, NeoX rope, GQA paged glue, Gemma-3 sliding-window for Olmo-3, tied embeddings, packed loader, ByteLevel BPE). Gate vehicle `allenai/OLMo-2-0425-1B` (1.485B, ~2.77 GiB, fits GB10 tight ~30 GiB free); Olmo-3 rides the same row (0.25.0 oracle-support UNVERIFIED). OLMo-1 (non-parametric LayerNorm), OLMoE/FlexOlmo (MoE), OlmoHybrid (SSM) stay `INVENTORIED`. Nothing implemented. | -| 2a | `ROAD-V1-C2-LOCAL-BF16` | Close local discrete-Blackwell Qwen3.5 plain-BF16 production parity | [`MODEL-MM-qwen3-5-qwen3-5-for-conditional-generation`](model-matrix.md), [`LOAD-SAFETENSORS-DIRECT-DENSE`](engine-matrix.md), [`SERVE-CLI-BENCH`](engine-matrix.md), [`KERNEL-SSM-MAMBA`](kernel-matrix.md) | Exact `(sequence, 8-token chunk)` GDN conv dispatch is byte-identical and default ON. Rebased-main same-binary reprofile: **718.704→233.955 ms (3.072x)** and whole run **+2.272%**; sealed vLLM remains 145.421 ms (**1.609x residual**). The reviewed combined default-off K4 causal-conv + 16-token post-conv local A/B is exact and improves total throughput **0.7175%**, TTFT **1.340%**, TPOT **0.556%** and E2E **0.716%**. The #206 pretoken frontend is CPU-GREEN, but its first real A/B failed output identity (98/128 requests; 15,507/16,384 positions) because one-at-a-time AsyncLLM publication changed the execution wave; all apparent timings are VOID. Atomic all-or-zero AsyncLLM wave admission matching vLLM's add-all-before-step behavior is now implemented and CPU-GREEN (12/12·433 async gate); fresh mutation review and real retry remain pending. Cross-engine throughput and TTFT remain **VOID/PENDING**; TPOT and VRAM remain OPEN | `GATING` | Fresh mutation review; operator CPU gate; then restart counterbalanced `STRING→TOKENS→TOKENS→STRING` GPU A/B and pinned-vLLM rerun. Latency, VRAM and 27B/35B correctness remain gates; no 4B-to-gate-model extrapolation. Evidence: [campaign spec](specs/sm120-qwen35-pareto-2026-08-09.md), [production baseline](../docs/bench-evidence/qwen35-4b-sm120-main-20260807.md) | +| 2a | `ROAD-V1-C2-LOCAL-BF16` | Close local discrete-Blackwell Qwen3.5 plain-BF16 production parity | [`MODEL-MM-qwen3-5-qwen3-5-for-conditional-generation`](model-matrix.md), [`LOAD-SAFETENSORS-DIRECT-DENSE`](engine-matrix.md), [`SERVE-CLI-BENCH`](engine-matrix.md), [`KERNEL-SSM-MAMBA`](kernel-matrix.md) | Exact `(sequence, 8-token chunk)` GDN conv dispatch is byte-identical and default ON. Rebased-main same-binary reprofile: **718.704→233.955 ms (3.072x)** and whole run **+2.272%**. The reviewed combined default-off K4 causal-conv + 16-token post-conv A/B is exact and improves total throughput **0.7175%**, TTFT **1.340%**, TPOT **0.556%** and E2E **0.716%**. Atomic all-or-zero AsyncLLM wave admission now matches vLLM's add-all-before-step behavior and fixes the pretokenized exactness confound. Corrected three-repetition cross-engine result: throughput **6831.71 vs 6643.40 tok/s (1.0283x PASS)**; host PSS passes; TTFT **1016.51 vs 936.59 ms (1.0853x slower)**, TPOT **34.477 vs 33.917 ms (1.0165x slower)**, E2E **5395.13 vs 5244.07 ms (1.0288x slower)** and VRAM **+118.7 MiB** remain open. Exact BV16+swizzle+REGSTATE and slack-memset local opt-ins improve all enclosing axes. Geometric argmax scratch (wait migration, +1% TPOT) and BF16 vector writeback (only 0.143% y800; unstable) are rejected and removed | `GATING` | Preserve the throughput stack; use same-tool interval profiling to split remaining TTFT intake/prefill and then target the largest attributable interval. VRAM and 27B/35B correctness remain gates; no 4B-to-gate-model extrapolation. Evidence: [campaign spec](specs/sm120-qwen35-pareto-2026-08-09.md), [production baseline](../docs/bench-evidence/qwen35-4b-sm120-main-20260807.md) | | 3 | `ROAD-V1-C3` | MTP k=1 + GDN speculative path, then DFlash, DSpark and heterogeneous-vocabulary TLI | [engine matrix](engine-matrix.md), [coverage view §8](feature-matrix.md#8-speculative-decoding) | MTP and DFlash specs exist. **M-mtp-0 CLOSED 2026-07-24** - the standalone MTP draft head is oracle-parity-proven on BOTH gate checkpoints (op-level parity vs a dumped k=1 vLLM oracle, not a token-generation SACRED gate). **I2 SCHEDULER-HALF LANDED 2026-07-24** ([mtp-spec-decode §2.7](specs/mtp-spec-decode.md)): host-side spec-decode scheduler/engine plumbing + the FROZEN spec-metadata ABI that I3 (rejection sampler) and I5 (verify/propose runner) build against - `SpeculativeConfig`, `DraftTokenIds`, `Request::spec_token_ids`/`NumTokensWithSpec`, the first population of `scheduled_spec_decode_tokens`, `Scheduler::update_draft_token_ids`, the `take_draft_token_ids` seam, `EngineCore::post_step`, `InputBatch::num_accepted_tokens`/`update_req_spec_token_ids`; DEFAULT-OFF and INERT (no `SpeculativeConfig` => `num_lookahead_tokens == 0` => byte-identical engine). `SPEC-MTP` **STAYS `GATING`** because M-mtp-1..4 (greedy rejection, GDN spec slots, k>1, CUDA graphs) are still open, so spec decode remains user-invisible. DSpark is user-promoted scope with DeepSeek-V4/Qwen3 draft models, reduced-vocabulary handling and full-CUDA-graph behavior inventoried under `SPEC-DSPARK`; tokenizer-agnostic target<->draft mapping is separately inventoried as `SPEC-TLI`. Their dedicated spikes are not written **I3 GREEDY REJECTION SAMPLER LANDED 2026-07-24** (`SPEC-REJECTION` `READY` -> `ACTIVE`): per-request logits expansion to `1 + k_i` rows plus the greedy accept rule (accept a draft iff it equals the target argmax; on the first mismatch emit the target argmax and stop; bonus token when all k accept), CUDA==CPU bit-exact at vocab 248320. **I4 GDN-HALF LANDED 2026-07-24** (`SPEC-GDN-SEGMENTS` -> `ACTIVE`): the GDN spec metadata split + decode->prefill reclassification, the `T>1`/`IS_SPEC` recurrence with per-timestep snapshots, the conv sliding window honouring `num_accepted`, k+1 slot allocation - bit-exact rollback. **I5a GDN LAYER ROUTING + runner spec-metadata upload LANDED 2026-07-24** (`CLAIM-SPEC-MTP-I5A`): `GdnBlockPaged` now routes a pure-spec batch through `vt::GdnSpecDecode`/`vt::CausalConv1dSpecUpdate` and the runner uploads I4's six spec device tensors - first of the scoped M-mtp-1 sub-increments (I5a GDN wiring -> I5b prepare_prefill -> I5c MTP paged propose -> I5d config+runner-loop+27B token gate, spec §5), DEFAULT-OFF INERT, bit-exact vs the I4 ops, no e2e loop yet. **I5b `prepare_prefill_inputs` LANDED 2026-07-24** (`CLAIM-SPEC-MTP-I5B`): the drafter prefill input-prep - shift each request's `input_ids` left one within its query span, splice the just-sampled next token, `query_len -= num_rejected`, emit last-token index / query_start_loc / seq_lens into the `SpecPrefillInputs` struct; a HOST routine (no new CUDA kernel; mirrors our DEVICE-NEUTRAL `prepare_inputs`/`combine_sampled_and_draft_tokens` family), unit-gated 7 cases / 27 assertions RED-first, DEFAULT-OFF INERT, additive by construction. **I5c MTP PAGED PROPOSE + DRAFT KV LAYER LANDED 2026-07-24** (`CLAIM-SPEC-MTP-I5C`): `Qwen3_5MTPModel::ForwardPaged` runs the head + one full_attention decoder layer over the head's OWN paged draft KV layer (ReshapeAndCache + PagedAttention over the target's block table / slot mapping); `MakeQwen3_5KVCacheSpec(num_spec>0)` adds that draft KV layer (`fa_draft` FullAttentionSpec group, index num_hidden_layers); `ForwardDeviceTap` exposes the `[T,H]` post-final-norm hidden tap (INERT); and `MtpProposePrefill` is the callable k=1 propose (I5b shift-splice -> one paged forward -> argmax at last_token_indices, early-exit). CORE PROOF: the paged forward reproduces I1's standalone head logits/argmax on BOTH gate checkpoints; a two-step drive proves the draft-KV write/read (RED control diverges). DEFAULT-OFF INERT (no spec config -> draft KV layer not allocated, tap nullptr, target forward byte-identical); NOT wired into the runner step loop. **I5d-pre REGISTRY/FORWARD-SEAM ENABLING REFACTOR LANDED 2026-07-25** (`CLAIM-SPEC-MTP-I5D-PRE`): a scoping pass found the model seam is fully TYPE-ERASED, so the runner cannot reach the concrete target weights / hidden-state tap / loaded MTP weights the I5d loop needs. Four ADDITIVE, inert-when-spec-off access paths + one latent-bug fix - the `hidden_tap` out-field on the type-erased `ModelForwardInput` (routes to the existing `ForwardDeviceTap`), a `LoadedModel::BuildMtpDraft` virtual (typed path to the draft, null for non-MTP), MTP weight loading + shard retention in `FromModelDir` behind `EngineParams::speculative_config`, and the `GPUModelRunner` ctor widened with optional draft/draft-KV/`SpeculativeConfig`; PLUS the latent `initialize_kv_cache` fix (select the FIRST non-eagle full-attn group as the target so a third `fa_draft` group can't displace it; byte-identical at num_spec==0). DEFAULT-OFF INERT, unit-gated RED-first, spec-off SACRED gates byte-identical. **I5d CONFIG RUNTIME + VERIFY/PROPOSE RUNNER LOOP LANDED as a spec-off-byte-identical PARTIAL 2026-07-25** (`CLAIM-SPEC-MTP-I5D`): `--speculative-config` JSON parse -> `EngineParams` -> `LoadedEngine` resolution (widened KV `MakeQwen3_5KVCacheSpec(num_spec>0)`, `BuildMtpDraft`, forced sync scheduling, `MakeScheduler(spec)`, `EngineCore(check_for_draft=true)`) + the full runner loop (draft splice, hidden-tap capture, GDN builder spec-overload feed, k+1 GDN state-slot remap + widened conv cache + draft-KV alloc, `MtpProposePrefill`, `take_draft_token_ids`, acceptance telemetry). CUDA `-Werror` 0 warnings, cutlass-ON; spec-OFF SACRED byte-identical (27B 235/235, 35B 315/315, Coder 138/138 + spec unit tests ALL PASS). **The three-way 27B token gate is NOT yet passing**: the spec-ON engine RUNS the loop end to end and MEASURES the blocker (`test_qwen27_spec_decode`) - it throws on the FIRST prefill step at `gdn_state_gather: working/cache row shapes must match` (`src/vt/ops.cpp:1773`) because I4's spec conv rollback needs the conv row widened to `(K-1)+num_spec` while the non-spec GDN conv ops assume `(K-1)`. `SPEC-MTP` STAYS `GATING`. **I5e LANDED 2026-07-25 (`CLAIM-SPEC-MTP-I5E`): the non-spec GDN conv ops made widened-cache-aware (mirror vLLM `state_len=KERNEL_WIDTH-1` + physical `stride_conv_state_tok`, leading `(K-1)` sub-window, byte-identical at `num_spec==0`) AND the async input-combine forced off under spec (it overwrote the verify batch's draft position with the committed token -> 0 acceptance, RCA'd on the real 27B). **THE THREE-WAY 27B SINGLE-REQUEST GREEDY GATE PASSES**: our-spec-ON == vLLM `--speculative-config mtp` greedy == our-spec-OFF token-for-token, **acceptance 16/16 drafts accepted** (~16 target steps saved); spec-OFF SACRED byte-identical (27B 235/235, 35B 315/315, Coder 138/138), compute-sanitizer 0 on the spec step. `SPEC-MTP` LEAVES `GATING` (single-request greedy correctness PROVEN); NOT `DONE` - the MIXED `GdnBlockPaged` concurrency split/merge + the throughput A/B vs vLLM same-config are I6. **I6 LANDED 2026-07-25 (`CLAIM-SPEC-MTP-I6`, `benchmark_binding=true`): the §5 c1 THROUGHPUT GATE — ours spec-ON AT/ABOVE vLLM spec-ON on EVERY measured axis at c1** (TPOT 66.2/62.95 vs 69.1/65.3 ms prose/code, ours ~1.04x faster; output tput +4.6%/+3.9%; ITL/TTFT lower; acceptance ours 0.85/0.92 vs vLLM 0.838, within noise; spec helps both ~1.5-1.6x TPOT; ours ~4% faster spec-OFF too), via an additive example-only `--speculative-config` bench flag (NO engine code touched). STAYS `ACTIVE`: the c>1 mixed-batch path is still refused + owes a c>1 A/B, and no server-facing spec flag yet. | `ACTIVE` | M-mtp-0, I2 scheduler-half, I3 rejection sampler, I4 GDN spec slots, I5a GDN layer routing, I5b prepare_prefill, I5c MTP paged propose + draft KV, I5d-pre the registry/forward enabling seam, and I5d config runtime + verify/propose runner loop (spec-off byte-identical) are landed; next (before `SPEC-MTP` leaves `GATING`) is closing the measured I5d gate blocker - make the non-spec GDN conv ops widened-cache-aware (mirror vLLM `causal_conv1d` `state_len=width-1+(seqlen-1)`) + the MIXED `GdnBlockPaged` split/merge - then the passing M-mtp-1 27B k=1 greedy three-way token gate + acceptance, then M-mtp-2 35B, then DFlash, the DSpark spike/gates and TLI. **DFlash D0-redo + D1 LANDED 2026-07-26 (`CLAIM-DFLASH-D0D1`, [dflash-spec-decode §0](specs/dflash-spec-decode.md)): `SPEC-DFLASH` UNBLOCKED + `ACTIVE` on the advanced pin `555967922`/vLLM 0.26.0.dev0.** The prior 0.25.0 ORACLE-BLOCKED verdict is SUPERSEDED — under `VLLM_USE_V2_MODEL_RUNNER=1` (vllm#40898 resolved) the mixed-SWA/full z-lab 27B draft CONSTRUCTS and the drafter is ALIVE (acceptance 2.21/8.80/4.75/4.57 > 1, `num_spec=16`, flashinfer-native fp8-KV; goldens committed). Gate FORM measured STRICT MODE-MATCHED (vLLM-ON run-deterministic K>=3 but != vLLM-OFF at k=16 near-ties — NOT the MTP three-way identity). **D1 `DF-AUX-TAPS` DONE:** the single hidden tap is generalized to the multi-tap `[T,H×taps]` (`ForwardDeviceMultiTap` capturing `(hidden+res)` at `target_layer_ids`), config-gated byte-identical off; unit gate 598 assertions (RED-first), CUDA 697/697 + sanitizer 0, 27B MTP e2e 9/9 + 27B SACRED 235/235 byte-identical (inertness). **D2-D5 LANDED 2026-07-26 (`CLAIM-DFLASH-D2`/`D3`/`D4D5`/`D5`):** the drafter model + the project's first non-causal in-block attention (D2, GPU parity vs the real vLLM draft), context-KV precompute + `prepare_dflash_inputs` (D3, GPU numeric-parity 61/61), the non-autoregressive whole-block propose brick + `dflash` config-select (D4), and the RUNNER-LOOP INTEGRATION + 27B e2e (D5): the full verify/propose loop is wired (separate z-lab draft load + target-shared bf16 embed/lm_head, aux-tap capture, per-request combined-feature context accumulation honoring num_rejected, `propose_drafts_dflash`) and RUNS end to end - `test_qwen27_dflash_spec_decode` 2/4 STRICT token-exact vs the vLLM-DFlash-ON golden + acceptance ~ vLLM on ALL 4 (19/39/29/25 vs 17/39/30/25). The 2 divergences are SINGLE bf16 near-tie flips (ratified near-tie ROOT = the D3 inline context-KV recompute envelope), NOT a wiring bug; inertness SACRED 235/235 + MTP 9/9 byte-identical; CUDA `-Werror` clean, no new kernel. NOT a clean strict-4/4 pass - STRICT 4/4 token-identity + the speed A/B = D6 (persistent paged draft-KV bit-matching vLLM's fused projections + the uniform-1+k FULL CG). Capture tool + goldens: `scripts/spec/d{0,2,3}_dflash_*.py`, `tests/parity/goldens/dflash_27b{,_draft,_kvprep}/`. **D6-D9 SPEED CAMPAIGN 2026-07-27:** D6 c1 A/B + STRICT-4/4 bf16-irreducibility RCA; D7 device-resident within-step forward (bit-identical); D8 acceptance RCA + FINAL golden A/B (ours 0.69× vLLM). **D9 (`CLAIM-DFLASH-D9`) PERSISTENT PAGED DRAFT-KV LANDED (bit-identical, +22.7%): `AppendContextKVHost` + `ForwardBlockLogitsWithPrecomputedKV` replace the O(context²) per-step recompute with an append-only per-request store; ours-ON 20.99→25.75 tok/s = 0.917× vLLM-ON (28.09, was 0.69×); e2e 27/27 SAME tokens, SACRED 235/235 + MTP 9/9 byte-identical, CUDA `-Werror` clean, no new kernel. D8's "bf16 acceptance ceiling" REFUTED — same-trajectory per-step acceptance == vLLM (ratio 1.00) and ours realized acceptance (3.68/step) > vLLM (3.31); the SOLE residual (~8%) is the FULL uniform-(1+k) CUDA graph (eager-vs-graphed), a closeable increment. SPEC-DFLASH stays `ACTIVE` (speed not yet ≥ vLLM).** **SPEC-MTP → `DONE` 2026-07-26 (`CLAIM-SPEC-MTP-DONE`, records-only, closing commit I7 `72f9fb1`):** the user ratified the c>1 near-tie+SPEED criterion, closing both I6-owed items (mixed-batch concurrency + server/CLI/C-ABI `--speculative-config`); MTP k=1 is COMPLETE + gated. **M-mtp-2 CLOSED 2026-07-26 (`CLAIM-SPEC-MTP-M-MTP-2`): the 35B `Qwen3_5MoeMTP` full e2e three-way token gate PASSES** — our spec-ON == our spec-OFF == vLLM 0.25.0 `--speculative-config mtp` greedy == vLLM spec-OFF, 16/16 vs the `greedy_ids` anchor (STRICT, c1), acceptance 16/16 both sides; c1 spec-ON 1.19x TPOT / +16.3% output-tput vs spec-OFF (acceptance 0.908) — MoE speedup transfers; spec-OFF byte-identical (test+docs-only). MTP is now `DONE` on BOTH gate models (`MODEL-SPEC-qwen3-5-mtp-qwen3-5-moe-mtp` `GATING`→`DONE`). **DFlash D11+D12 2026-07-27 — the FULL uniform-(1+k) CUDA graph is being built in three parts:** D11 (`CLAIM-DFLASH-D11`) landed Part A (the device-store primitive, CPU-gated); **D12 (`CLAIM-DFLASH-D12`) landed A-wire (the D11 device store is now the PRODUCTION path; GPU-gated e2e 27/27 all-exact acceptance 19/39/29/25 + SACRED 235/235 + MTP 9/9 byte-identical) + Part B (`vt::DFlashPagedBlockAttention`, the capture-safe paged kernel; `test_ops_dflash_paged_block_attn` 795648/795648 CPU==CUDA + cross-check vs materialized `DFlashBlockAttention` + compute-sanitizer 0; NO function-local host `cu_seqlens` upload = capture-UAF fixed).** Speed UNCHANGED 0.917× (A-wire eager + Part B not yet wired). The SOLE remaining piece is Part C (static-shape capture + device mask-scatter + `BeginCapture`/replay + the ≥vLLM c1 A/B); if ours-ON-graphed ≥ vLLM-ON → SPEC-DFLASH DONE → C3 complete. C3 stays `ACTIVE` (DFlash Part C + DSpark/TLI remain) | | 4 | `ROAD-V1-C4` | Quantization: llama.cpp breadth/speed, NVFP4/FP8/MX, MLX native | [quantization matrix](quantization-matrix.md) | coverage spike merged; `QUANT-GGUF-CPU-THREADPOOL` W1-W3 implemented and correctness-gated, still `GATING` (its reproduction now exists — same-binary 1-vs-20-thread A/B is prefill 12.47x / decode 8.05x / RSS 1.000x, so **decode misses the >=10x bar**). **GGUF COMPUTE-IN-QUANT IS NOW LIVE AND DEFAULT-ON (2026-07-22, `CLAIM-QUANT-GGUF-CIQ-G4-1`):** [compute-in-quant GEMM](specs/gguf-compute-in-quant-gemm.md) **G1-G4** — block dtypes + traits, the Q8_0/Q8_K activation quantizers, the six generic `vec_dot`, `kMatmulBTQuant`, and now the ROUTING (`vt::MatmulBT` dispatches a block-dtype weight to it) — plus [keep-quant loader](specs/gguf-keep-quant-loader.md) **L1-L4**, whose master switch defaults ON wherever that op is registered for the running device (CPU today; a CUDA build still expands). Six encodings (Q4_0, Q8_0, Q3_K, Q4_K, Q5_K, Q6_K) now carry `C` = `Y`. **Correctness held exactly**: the 35B GGUF gate is 16/16 token-exact vs the same-file llama.cpp oracle with the quant path on, and the bench model's output tokens are byte-identical across the pre-G4, post-G4 and `VT_CPU_REF=1` arms — no golden regenerated. **Binding CPU A/B** (idle dgx aarch64, one flock, same binary, 3 reps): decode **3.45x**, prefill **4.16x**, peak RSS **1.16x less**; vs llama.cpp we went from 11.7x / 34.1x / 2.66x behind to **3.38x / 8.20x / 2.29x**. The projected 9-17x did **not** hold, for a measured reason: 60 % of that file's weight bytes are `f16`, which no block encoding covers. Keep-quant loader **L4** is therefore MEASURED-and-NOT-MET on RSS; other leaf specs open. **THAT #1 LEVER IS NOW LANDED (2026-07-22, `CLAIM-KERNEL-CPU-ELEM-GEMM-1`, new row [`KERNEL-GEMM-CPU-ELEM`](kernel-matrix.md)):** [the elementwise CPU GEMM](specs/cpu-elementwise-gemm.md) **E1-E4** — per-dtype specialization out of the K loop, 16 independent accumulators instead of one, AArch64 NEON + x86-64 SSE2/F16C tiers behind a runtime probe, and M-blocking — all **BYTE-IDENTICAL** to the historical kernel (`memcmp` gate, exhaustive 65,536-pattern widening check, same token md5), so nothing was regenerated. Binding same-binary A/B: prefill **3.41x**, decode **3.11x**; op-level bf16 18-24 -> 69-351 GFLOP/s. **vs llama.cpp: decode AT PARITY (1.03x), prefill 2.34x behind, RSS 2.29x worse. THEN loader L5 LANDED (2026-07-23, `CLAIM-QUANT-GGUF-KEEPQ-L5-1`):** [keep-quant loader](specs/gguf-keep-quant-loader.md) **L5** — mmap in-place residency (borrow kept blocks out of the read-only mapping, refcounted), tied-head sharing (one bf16 vocab matrix for embed+lm_head), and a read-once page release (port of llama.cpp `unmap_fragment`) — took **peak RSS 6.401 -> 3.884 GiB (2.29x -> 1.39x llama.cpp)** with decode UNCHANGED and output tokens byte-identical (md5 `d235db12f2cd304007530286a1755c95`). The remaining ~1.09 GiB over llama.cpp is the f16 expansion (no block encoding covers f16). | `PARTIAL` | **THE OWED FRESH PROFILE IS DONE (2026-07-23) and it re-ranks the plan.** A `vt::GetOp` hook (100% of wall time, reverted before binding) on the CURRENT binary: prefill is **no longer GEMM-bound** — kMatmulBTQuant 37%, **kGdnPrefill 25%**, kMatmul 12%, kMatmulBT 10%, **kPagedAttention 10%**; the two non-GEMM kernels (GDN linear-attention recurrence + paged attention) run **SINGLE-THREADED** on the CPU and are now the top prefill levers. Decode is memory-bound matmul at parity, no kernel work owed. **RE-RANK: G5/G6/G7 all only speed the already-fast quant GEMM and rank BELOW the two serial non-GEMM kernels; the new #1 CPU lever is threading kGdnPrefill + kPagedAttention.** **THAT #1 LEVER IS NOW LANDED (2026-07-23, `CLAIM-CPU-THREAD-GDN-PAGED-1`, [two-kernel threading](specs/cpu-thread-gdn-paged-2026-07-23.md)):** kGdnPrefill chunks over the (sequence, value-head) axis and kPagedAttention over query-token rows, both via the existing `ParallelForRows`, both **byte-identical** (qwen35 output-token md5 `d235db12f2cd304007530286a1755c95` unchanged at threads 1/4/20 + `VT_CPU_REF=1`, determinism battery extended, CPU ctest 158/158). **Binding dgx aarch64 (idle): prefill 1.382× same-binary (73.0→100.9 t/s), 2.43×→1.76× behind llama.cpp pp128; decode at parity; op-scaling 1→20 GdnPrefill 7.08× / PagedAttention 8.96×; fresh profile shows the two kernels 35%→8.6% of prefill and re-ranks the NEW bottleneck to the GEMMs (kMatmulBTQuant 50% + kMatmul 16% + kMatmulBT 14% = 80%) ⇒ next CPU lever is the SIMD/repack GEMM tiers (G5/G6/G7).** **THE FIRST SUCH TIER IS NOW LANDED (2026-07-23, `CLAIM-QUANT-GGUF-CIQ-G6-1`, [compute-in-quant GEMM](specs/gguf-compute-in-quant-gemm.md) G6):** the Arm **i8mm mmla `nrc==2`** `vec_dot` 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 (decode M=1 → portable, unchanged); runtime `HWCAP2_I8MM` probe + `VT_CPU_QUANT_MMLA` defeat + per-file `+i8mm`. **BYTE-IDENTICAL** where the math allows (q8_0/q4_0 bit-exact to the scalar tier, q4_K/q6_K within NMSE ≤ 5e-4), bit-identical across threads 1/2/4/20, e2e token md5 `d235db12f2cd304007530286a1755c95` byte-identical (mmla on/off/`VT_CPU_REF=1`), 35B GGUF gate 16/16 vs llama.cpp on both files. **Op-level portable→i8mm: q4_K 7–8.4×, q6_K 3.8–4.5×, q8_0 ~1.2×**; e2e prefill same-binary **1.084×** on the q8_0-dominant bench file (1.56×→**1.44× behind** llama.cpp pp128, Amdahl-bounded — the big k-quant win lands on the APEX 35B files). Fresh bottleneck: the elementwise f16/f32 GEMM (~30%, unchanged) is now co-dominant on this mixed file. CUDA `-Werror` 0-warn, regression set UNCHANGED. `docs/BENCHMARKS.md` ACCEPTED. RSS deficit closed to 1.39x by L5; the last RSS lever is an f16 keep-as-is compute path, not this loader. **THEN the GDN split-projection orientation LANDED (2026-07-23, `CLAIM-CPU-GDN-ORIENT-1`, [GDN projection orientation](specs/cpu-gdn-proj-orientation-2026-07-23.md)):** a fresh op-dispatch profile of the current binary (warm prefill, `vt::GetOp` hook + per-GEMM shape histogram, reverted before binding) found the four GDN input projections (`in_proj_qkv/z/b/a`, 72 GEMMs, **17.9%** of prefill: `kMatmulBTQuant 50.7% / kMatmul 17.9% / kMatmulBT 14.9%`) were the LAST weight family `LoadGdnGguf` still transposed into [K,N] (nk=false → the N-striding `kMatmul`, no M-blocking) after G4's `expand_nk` gave every other expanded weight the file's own [N,K] order. New `GgufLoadPolicy::gdn_expand_nk` + `MakeGdnProj` keep them [N,K] nk=true → the M-blocked `kMatmulBT`; **BYTE-IDENTICAL** (same sequential f32 K-reduction, only the weight offset differs — token md5 `d235db12f2cd304007530286a1755c95` unchanged across on/`VT_GGUF_GDN_NK=0`/`VT_CPU_REF=1` and threads 1/4/20), `test_qwen36_gguf_engine` 2/2·28/28·16/16 on APEX. **Binding same-binary prefill 1.090× / decode 1.09× (44.1→40.4 ms TPOT = 1.01× llama tg32, at parity), 1.44×→1.32× behind llama.cpp pp128, RSS unchanged.** Fresh post-change profile: `kMatmul` **17.9%→0% (72→0 calls, ELIMINATED)**, absorbed into `kMatmulBT` (14.9%→27.7%); **next CPU prefill lever = the quant GEMM (kMatmulBTQuant, now 55%): G7 repack-at-load.** **G7 LANDED 2026-07-23 (`CLAIM-QUANT-GGUF-CIQ-G7-1`, [compute-in-quant GEMM](specs/gguf-compute-in-quant-gemm.md) G7):** q8_0 repacked once at load into the i8mm `block_q8_0x4` interleave (ported llama.cpp `repack.cpp` `q8_0_4x8`), `kMatmulBTQuant` dispatches a pre-shuffled gemm/gemv with no per-block register shuffles. **BIT-IDENTICAL** (byte-permute weight + non-fused `vmlaq_f32`, 305-assertion memcmp across decode/prefill/bf16-out/threads, token md5 `d235db12f2cd304007530286a1755c95` unchanged on/`VT_CPU_QUANT_REPACK=0`/`VT_CPU_REF=1`; a `ResidentWeight`/`MakeTensor` flag-drop that produced all-zero tokens was caught by the E2E gate and fixed). Op-level q8_0 **3.7–5.9×** (518→2401 / 583→3456 / 514→1902 GFLOP/s); **E2E prefill 1.92× same-binary (1096→572 ms), 223.8 t/s vs llama.cpp pp128 177.3 = 1.26× — AT/BEYOND PARITY** (was ~1.5× behind), decode at parity, RSS 3.884 GiB unchanged. Fresh profile: q8_0 GEMM 55%→~21%; **the CPU prefill-lever search is CLOSED — the sole remaining gap to llama.cpp is peak RSS (1.39×), not prefill.** CUDA-inert (gated off any non-CPU-quant device), CUDA `-Werror` 0-warn, regression set UNCHANGED. **THE RSS GAP IS NOW CORRECTLY ATTRIBUTED (2026-07-23, `CLAIM-QUANT-GGUF-KEEPF16-L6-1`, [keep-quant loader](specs/gguf-keep-quant-loader.md) L6): it is NOT the f16 expansion.** L6 implemented keep-f16 residency (keep the file's 56 F16 weights + tied head resident as F16 and compute on them, mirroring llama.cpp `ggml_vec_dot_f16`) and MEASURED it **RSS-NEUTRAL** (3.884 → 3.832 GiB, −52 MB) and prefill-regressive (TTFT 577 → ~1000 ms, from 1.25× ahead of llama.cpp to 0.72× behind) — because L5's page-release had ALREADY dropped the f16 file pages, so keep-f16 only swaps an anonymous bf16 buffer for equal-size file-backed f16 pages. smaps attribution proves our weight residency is at llama.cpp parity (file-backed 2.63 ≈ 2.68 GiB); **the residual ~1.08 GiB is the engine's ANONYMOUS activation/KV workspace, not weights — the real, separate CPU RSS lever.** keep-f16 ships DEFAULT OFF (`VT_GGUF_KEEP_F16=1` opt-in), tokens byte-identical (md5 `d235db1…`), `test_gguf_keep_quant` 35/35 (x86+aarch64), regressions UNCHANGED (27B 235/235, 35B 315/315, Coder 138, dense 184, OPT 63, DeepSeek 223, Llama 92, GGUF engine 28/28). **NEXT CPU RSS lever: profile + shrink the engine's activation/KV working set, NOT the weight loader** | | 5 | `ROAD-V1-C5` | Sliding window, local attention, YaRN/long context | [engine matrix](engine-matrix.md), [coverage view §§2,11](feature-matrix.md#2-kv-cache--memory), [joint spike](specs/sliding-local-yarn-long-context.md) | **CUDA GPU CLOSURE 2026-07-27 (`CLAIM-ROADMAP-C5`, dgx GB10 sm_121a, clean build of `489f7771`, oracle vLLM 0.26.0.dev0):** the shared scaled-RoPE + local-mask CUDA path COMPILES `-Werror`-clean and RUNS on GB10; the C5 feature-positive correctness gates that were the stated `GATING` blocker now PASS — SWA (Gemma-2/Gemma-3 48/48), LongRoPE (Phi-4-mini 16/16, RED-first), llama3-rope (Llama-3.2-1B 16/16), dynamic-NTK (InternLM2 16/16); both RoPE 0.26-oracle recaptures BIT-IDENTICAL to goldens (zero drift). Leaves `ATTN-SLIDING-WINDOW`/`ATTN-ROPE-{LLAMA3,LONGROPE,DYNAMIC-NTK}`/`ATTN-YARN` → `ACTIVE` | `PARTIAL` | (RI) **Honest residual (vehicle-blocked, not skipped):** YaRN model e2e (no cached Nomic/gpt-oss consumer) + chunked-local model e2e (no Llama4 row) are REACHABLE-BLOCKED — operator/formula stay GPU/G3-gated; long-context positive-mask (prompt > W) SWA model e2e + the KV memory-optimization G8; and the roadmap-wide every-axis SPEED tail (all C5 leaves correctness-complete, speed-pending, mirroring their model consumers). Not row-DONE until speed + the blocked vehicles close | diff --git a/.agents/specs/sm120-qwen35-argmax-scratch-growth-2026-08-09.md b/.agents/specs/sm120-qwen35-argmax-scratch-growth-2026-08-09.md index cc8910d13..799ea148b 100644 --- a/.agents/specs/sm120-qwen35-argmax-scratch-growth-2026-08-09.md +++ b/.agents/specs/sm120-qwen35-argmax-scratch-growth-2026-08-09.md @@ -248,5 +248,29 @@ largest same-tool interval; never declare a ceiling. ## Outcome -Pending fresh RED-first implementation, mutation review, operator verification, -and same-binary GPU evidence. No default or performance claim is made. +**REJECTED and removed, 2026-08-09.** Fresh RED-first implementation and +mutation review passed at `06db3bbb3454f85e20f75b96eeb299afba7036ae` (9 +portable cases / 105 assertions plus the CUDA sampling gate). The immutable +same-binary series +`INC-a -> GEO-a -> GEO-b -> INC-b -> INC-c -> GEO-c` was token-exact in all +six legs (SHA-256 `be20ffbceb61f0264ca21d972bfc5fc51e855e64f2b945de71669cae666aa702`). +Raw evidence is `/tmp/qwen35-ab-argmax-geometric-06db3bbb/`. + +| Arm mean (3 reps) | total tok/s | output tok/s | mean TTFT | mean TPOT / ITL | mean E2E | +|---|---:|---:|---:|---:|---:| +| Incumbent | 6862.2667 | 758.8100 | 1009.7167 ms | 34.3400 ms | 5371.0000 ms | +| Geometric | 6863.8867 | 758.9867 | 965.0133 ms | 34.6833 ms | 5369.7067 ms | +| Change | +0.0236% | +0.0233% | -4.427% | **+0.9998%** | -0.0241% | + +The candidate therefore fails the no-TPOT-regression Pareto gate. Same-tool +`nsys --trace=cuda,nvtx --cuda-graph-trace=node` explains why the apparent +TTFT win is not an end-to-end compute saving: incumbent `cudaFree` was +506 calls / 2960.361 ms (max 200.004 ms), candidate was 476 calls / 27.705 ms +(max 0.688 ms), exactly removing the 30 expected argmax frees; meanwhile +`cudaStreamSynchronize` rose from 18,027.497 to 20,878.260 ms over the same +567 calls. The wait moved to the later synchronization boundary. Candidate +trace hashes are `1021b045...239e7` (`.nsys-rep`) and +`808a7516...b7be4` (SQLite); the incumbent SQLite hash is +`bab04901...be178`. No default or performance credit is taken, and the +selector/product/tests are removed while this falsified hypothesis remains as +evidence. diff --git a/.agents/specs/sm120-qwen35-gdn-decode-bf16-vector-writeback-2026-08-09.md b/.agents/specs/sm120-qwen35-gdn-decode-bf16-vector-writeback-2026-08-09.md index 7e6b643dd..6d79c5a14 100644 --- a/.agents/specs/sm120-qwen35-gdn-decode-bf16-vector-writeback-2026-08-09.md +++ b/.agents/specs/sm120-qwen35-gdn-decode-bf16-vector-writeback-2026-08-09.md @@ -1,6 +1,6 @@ # sm_120 Qwen3.5 GDN decode BF16 vector-writeback discriminator -**Lifecycle:** `SPIKED`; implementation and measurement pending +**Lifecycle:** `REJECTED/REMOVED`; exact but below the timing gate **Owner rows:** `KERNEL-SSM-MAMBA`, `ROAD-V1-C2-LOCAL-BF16` @@ -130,3 +130,38 @@ unavailable 27B/35B acceptance. Add only portable parser/eligibility/mapping/pack-reference tests and capture their red failure. Do not edit CUDA recurrence code until that evidence exists. + +## Outcome — rejected and removed + +Reviewed implementation `29de225c8` plus repair `3ca49e926` passed the fresh +mutation re-review. The repaired portable suite ran 16 cases / 2,568 assertions; +independent mutations killed every eligibility predicate, BF16 RN conversion, +either pack, a hardwired-false dispatch, the CUDA route and the vector store. +The operator rebuilt the CUDA translation unit and passed +`test_gdn_decode_fused` plus `test_ops_gdn`. The immutable four-leg run was +token-exact throughout (SHA-256 +`be20ffbceb61f0264ca21d972bfc5fc51e855e64f2b945de71669cae666aa702`). + +| Arm | all fused decode, 1,656 calls | grid-y 800, 864 calls | total / output tok/s | TTFT | TPOT / ITL | E2E | +|---|---:|---:|---:|---:|---:|---:| +| REG-a | 251.210058 ms | 139.127859 ms | 6760.82 / 747.59 | 1027.95 ms | 34.81 ms | 5449.30 ms | +| VEC-a | 250.472657 ms | 138.702659 ms | 6802.32 / 752.18 | 1015.76 ms | 34.66 ms | 5417.91 ms | +| VEC-b | 251.214368 ms | 139.138796 ms | 6804.21 / 752.39 | 1015.60 ms | 34.65 ms | 5416.35 ms | +| REG-b | 251.404693 ms | 139.111428 ms | 6807.51 / 752.76 | 1013.82 ms | 34.65 ms | 5413.73 ms | + +The counterbalanced means improve y800 only **0.1430%** and all fused decode +only **0.1846%**. VEC-b is slower than both controls on y800 and slower than +REG-a on all fused decode, so the candidate fails both the required 1.00% gain +and the every-vector-leg stability gate. The enclosing differences are run +order drift: the final REG-b also beats both vector legs on every reported +wall axis. `nsys` confirms unchanged gx8/gy800/bx128, 56 registers/thread, +9,728 dynamic shared bytes and zero local bytes in every arm; static inspection +shows two 16-byte stores, but the old PTX artifact predates the repaired pack +seam and receives no current-head codegen credit. + +Raw roots are +`/tmp/qwen35-gdn-vecstore-{rega,veca,vecb,regb}-3ca49e926.*`; the GPU was idle +before the locked, cgroup-contained series (RTX 5070 Ti, driver 595.71.05, +P8, 40 C, 0% utilization). The experiment is therefore rejected and its +selector/product/tests are removed. REGSTATE remains the accepted opt-in; no +default, release, pinned-vLLM, 27B or 35B claim changes. diff --git a/docs/BENCHMARKS.md b/docs/BENCHMARKS.md index c847ddb97..f814a1de7 100644 --- a/docs/BENCHMARKS.md +++ b/docs/BENCHMARKS.md @@ -41,7 +41,7 @@ The binding comparison. vLLM runs its **production graphed config**, never | Qwen3.6-27B | NVFP4 (`nvidia` @`0893e160`, ModelOpt `modelopt_mixed`) | 0.25.0 | 0/4 | **BEHIND, uniformly 0.85x** on decode throughput (was 0.72x before the FP8 tower fix); greedy continuation IDENTICAL to vLLM. A different model from the `unsloth` row (NVFP4 MLP + FP8 W8A8 GDN/attn tower) | | Qwen3.6-35B-A3B | NVFP4 `modelopt_mixed` | 0.25.0 | 2/18 | 3-rep grid 2026-08-05 @`1ea26427`: 0.93-1.03x (c4 wins), c16 0.93x. Both c16 levers A/B'd NEG: drain event -1.9%, mirror 0.999x. ★ probe found a prod async batch-1 greedy DEGENERATION bug the mirror fixes | | DeepSeek-V2-Lite | bf16 MLA | 0.25.0 | 4/25 | Attributed miss, row stays `ACTIVE` | -| Qwen3.5-4B | bf16 direct-load | 0.26.0.dev0 | host PSS; frontend exactness failed | #206 pretoken A/B FAILED 98/128; timings VOID. Atomic-wave retry pending (CPU-green). Opt-ins: +0.718% tput/-0.556% TPOT; tput/TTFT VOID, TPOT/VRAM OPEN ([data](bench-evidence/qwen35-4b-sm120-main-20260807.md)) | +| Qwen3.5-4B | bf16 direct-load | 0.26.0.dev0 | **total/output throughput 1.0283x; host PSS lower** | Corrected exact comparison: throughput 1.0283x. OPEN: TTFT/TPOT/E2E 1.0853/1.0165/1.0288x slower; VRAM +118.7 MiB. Local GDN opt-ins improve every timing axis ([data](bench-evidence/qwen35-4b-sm120-main-20260807.md)) | ### GDN prefill causal-convolution by GPU diff --git a/docs/STATUS.md b/docs/STATUS.md index e9d0e5286..7f722e4a9 100644 --- a/docs/STATUS.md +++ b/docs/STATUS.md @@ -78,7 +78,7 @@ token-for-token correctness against the pinned oracle. | Qwen3.6-27B (NVFP4) text generation | Correctness-complete; speed is CHECKPOINT-dependent | Token-exact GB10 on both. `unsloth` @`890bdef7` beats vLLM 0.25.0 every c (1.007-1.045x), 115/124; `nvidia` @`0893e160` (ModelOpt FP8 tower) is **0.85x BEHIND**, decode ~100% GPU-busy | | Qwen3.6-35B-A3B (NVFP4, GDN MoE) | Correctness-complete; binding grid @`a0fa12c7` FLAT 0.935x-0.979x over c1-c32 (CoV <0.81%); the prior 0.87x c2 / 0.92x c8 "weak cells" were harness mismatch, not code; memory PSS 3.81x, GPU 1.40x | Token-exact SYNC+ASYNC; `VT_ASYNC_DEVICE_MIRROR` ON fixes async batch-1 token-0 degeneration; `VT_ASYNC_EXECUTOR` Option A NEUTRAL → OFF | | Qwen3 / Qwen2 dense (BF16) | Correctness-complete, speed-pending. Async-serving P0 FIXED (`ROW-SERVE-ASYNC-DENSE-MIRROR`): classic-dense `Qwen3ForCausalLM` now honors the async device token-ids mirror; CPU-only -Werror test-guard fixes x2 | Near-tie-robust token-exact vs vLLM (Qwen3-0.6B, Qwen3-4B); c1 effective parity, c8 decode residual. **Async device-mirror (`ROW-SERVE-ASYNC-DENSE-MIRROR`, `f9c969ae`): the #31 fix ported to the classic dense family, dgx-VERIFIED.** The shared dense `EmbedInto` (qwen3.cpp) raced the async combine's device input-ids write against a stale host upload → token-0 degeneration on the depth-2 AsyncLLM serving path (quant-independent). `EmbedInto` now consumes the device override published by `ForwardQwen3ForCausalLM`'s `DeviceTokenIdsScope` (27B-dense template); gate `test_qwen3_dense_async_serving` RED on `VT_ASYNC_DEVICE_MIRROR=0`, GREEN default, byte-identical mirror-off. dgx GB10: async gate RED→GREEN 0.6B+4B, SACRED 0.6B+4B 184/184 unchanged (byte-neutral sync path), memcheck 0 errors; Yi30/Qwen3-8B-MXFP4 default-config e2e coherent + 3/4 token-exact (p2 = oracle-ratified near-tie, gap 0.0000), closing the QUANT-CT-MXFP4 async-default residual. RESIDUAL: sibling InternLM2/Mistral/Llama scope one-liner; W4 bench RAN; FA2 GQA-swap default-ON, c2-c8 <1.0x. `FLASH-PTXAS` #82: codegen at PARITY (no ptxas lever); gap=engine context. **D1 (2026-07-31, `CLAIM-D1-BF16-MERGED-QKV`): the bf16 merged-QKV path (`Qwen3QkvMergeEnabled`/`VT_QWEN3_QKV_MERGE`) is now default-ON** — one `vt::MatmulBT` over the merged `[qdim+2kdim,H]` owner + a contiguous `vt::QkvSplit` (OLMo-2 exemplar), replacing three per-shard GEMMs. Bit-exact GEMM math (A/B unit `test_ops_qkv_merge` byte-identical, RED-first); the wider-N cuBLASLt K-reduction flips the 0.6B genuine bf16 near-tie so the SACRED 0.6B golden was regenerated (all tokens within the near-tie band, max 0.125 nats), while Qwen3-4B is byte-neutral (0 diffs, stays STRICT). Re-gated 0.6B 16/16 + 4B 16/16; consistency/launch-count fold (measured NEUTRAL on 4B decode), no new throughput owed | -| Qwen3.5-4B BF16 direct-load on discrete CUDA | Correctness-complete; atomic frontend CPU-green; TPOT/VRAM open | #206 pretoken A/B exactness FAIL 98/128; timings VOID. Atomic: async 13/13·435, core 14/14·111, bench 11/11·80; re-review/retry pending. TPOT 1.025x; VRAM +233 MiB ([data](bench-evidence/qwen35-4b-sm120-main-20260807.md)) | +| Qwen3.5-4B BF16 direct-load on discrete CUDA | Correct; throughput/host PSS ahead; latency/VRAM open | Atomic pretoken exact. Ratios: tput 1.0283x; TTFT/TPOT/E2E 1.0853/1.0165/1.0288x slower; VRAM +118.7 MiB. GDN local stack retained ([data](bench-evidence/qwen35-4b-sm120-main-20260807.md)) | | Qwen3-Coder-30B-A3B MoE (BF16) | Correctness-complete, speed-pending | Near-tie-robust token-exact 6/6; 11 of 16 binding grid cells at or above vLLM. **D1 (2026-07-31): inherits the default-ON bf16 merged-QKV via the shared dense `AttnBlock` — byte-neutral (0 token diffs, golden UNCHANGED); re-gated 6/6** | | Llama-3.x dense (BF16) | Correctness-complete, speed-pending | Near-tie-robust token-exact 16/16 (Llama-3.2-1B); llama3 RoPE scaling | | Mistral dense (BF16) | Correctness-complete, speed-pending | Paged-engine token-exact 16/16 (Mistral-7B-v0.3) | diff --git a/include/vt/cuda/argmax_scratch.h b/include/vt/cuda/argmax_scratch.h deleted file mode 100644 index 5f178a311..000000000 --- a/include/vt/cuda/argmax_scratch.h +++ /dev/null @@ -1,211 +0,0 @@ -// vllm.cpp original -- portable ownership policy for CUDA argmax scratch. -#ifndef VT_CUDA_ARGMAX_SCRATCH_H_ -#define VT_CUDA_ARGMAX_SCRATCH_H_ - -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include - -namespace vt::cuda { - -inline bool ArgmaxGeometricScratchEnabled(const char* value) noexcept { - return value != nullptr && value[0] == '1' && value[1] == '\0'; -} - -template -decltype(auto) DispatchArgmaxScratch(const char* value, Incumbent&& incumbent, - Candidate&& candidate) { - if (ArgmaxGeometricScratchEnabled(value)) { - return std::forward(candidate)(); - } - return std::forward(incumbent)(); -} - -inline std::size_t RoundArgmaxScratchCapacity(std::size_t need) { - if (need <= 1) return need; - std::size_t capacity = 1; - while (capacity < need) { - if (capacity > std::numeric_limits::max() / 2) { - throw std::overflow_error("argmax geometric scratch capacity overflow"); - } - capacity *= 2; - } - return capacity; -} - -struct ArgmaxScratchKey { - int device = 0; - std::uintptr_t stream = 0; - - friend bool operator==(const ArgmaxScratchKey&, const ArgmaxScratchKey&) = default; -}; - -struct ArgmaxScratchView { - float* values = nullptr; - std::int64_t* indices = nullptr; - std::size_t capacity = 0; -}; - -struct ArgmaxScratchDiagnostics { - void* values = nullptr; - void* indices = nullptr; - std::size_t capacity = 0; - std::size_t active_bytes = 0; - std::size_t retired_bytes = 0; - std::size_t growths = 0; -}; - -class ArgmaxScratchOwners { - public: - // The caller's submit callback launches both reduction kernels. It runs while - // the per-(device, stream) owner mutex is held, preventing another host thread - // from replacing or overwriting the pair between those submissions. - template - void Submit(ArgmaxScratchKey key, std::size_t need, CaptureQuery&& is_capturing, - Allocate&& allocate, CleanupPartial&& cleanup_partial, Retire&& retire, - Launch&& submit) { - auto grow = [&](Owner* owner, bool capture_already_checked) { - if (need <= owner->capacity) return; - if (!capture_already_checked && std::forward(is_capturing)()) { - ThrowCaptureMiss(); - } - const std::size_t new_capacity = RoundArgmaxScratchCapacity(need); - if (new_capacity > std::numeric_limits::max() / sizeof(std::int64_t)) { - throw std::overflow_error("argmax geometric scratch byte-size overflow"); - } - - void* new_values = nullptr; - void* new_indices = nullptr; - new_values = std::forward(allocate)(new_capacity * sizeof(float)); - if (new_values == nullptr) { - throw std::runtime_error("argmax geometric scratch value allocation returned null"); - } - try { - new_indices = std::forward(allocate)(new_capacity * sizeof(std::int64_t)); - if (new_indices == nullptr) { - throw std::runtime_error("argmax geometric scratch index allocation returned null"); - } - } catch (...) { - std::forward(cleanup_partial)(new_values); - throw; - } - - void* old_values = owner->values; - void* old_indices = owner->indices; - const std::size_t old_capacity = owner->capacity; - owner->values = new_values; - owner->indices = new_indices; - owner->capacity = new_capacity; - ++owner->growths; - if (old_values != nullptr) { - owner->retired_bytes += old_capacity * sizeof(float); - std::forward(retire)(old_values); - } - if (old_indices != nullptr) { - owner->retired_bytes += old_capacity * sizeof(std::int64_t); - std::forward(retire)(old_indices); - } - }; - - std::unique_lock owners_lock(owners_mu_); - auto it = owners_.find(key); - if (it == owners_.end()) { - if (std::forward(is_capturing)()) ThrowCaptureMiss(); - auto new_owner = std::make_unique(); - Owner* owner = new_owner.get(); - std::unique_lock submit_lock(owner->submit_mu); - // Keep the map lock until the first pair is complete. A failed first - // allocation therefore cannot publish an empty owner or race a second - // creator for the same key. - grow(owner, true); - owners_.emplace(key, std::move(new_owner)); - owners_lock.unlock(); - std::forward(submit)(ArgmaxScratchView{static_cast(owner->values), - static_cast(owner->indices), - owner->capacity}); - return; - } - - Owner* owner = it->second.get(); - owners_lock.unlock(); - std::lock_guard submit_lock(owner->submit_mu); - grow(owner, false); - std::forward(submit)(ArgmaxScratchView{static_cast(owner->values), - static_cast(owner->indices), - owner->capacity}); - } - - std::optional Diagnostics(ArgmaxScratchKey key) const { - Owner* owner = nullptr; - { - std::lock_guard lock(owners_mu_); - const auto it = owners_.find(key); - if (it == owners_.end()) return std::nullopt; - owner = it->second.get(); - } - std::lock_guard lock(owner->submit_mu); - return ArgmaxScratchDiagnostics{ - owner->values, owner->indices, - owner->capacity, owner->capacity * (sizeof(float) + sizeof(std::int64_t)), - owner->retired_bytes, owner->growths}; - } - - std::size_t OwnerCount() const { - std::lock_guard lock(owners_mu_); - return owners_.size(); - } - - private: - struct Owner { - mutable std::mutex submit_mu; - void* values = nullptr; - void* indices = nullptr; - std::size_t capacity = 0; - std::size_t retired_bytes = 0; - std::size_t growths = 0; - }; - - struct KeyHash { - std::size_t operator()(ArgmaxScratchKey key) const noexcept { - std::size_t hash = static_cast(key.stream); - hash ^= static_cast(static_cast(key.device)) + 0x9e3779b97f4a7c15ULL + - (hash << 6) + (hash >> 2); - return hash; - } - }; - - [[noreturn]] static void ThrowCaptureMiss() { - throw std::runtime_error( - "argmax geometric scratch capacity miss during CUDA graph capture; " - "warm or reserve scratch before capture"); - } - - mutable std::mutex owners_mu_; - std::unordered_map, KeyHash> owners_; -}; - -// This is the single production routing seam between cuda_sample and the pure -// ownership state machine. Runtime is intentionally capability-limited: the -// old published blocks can only be retired, while only a never-published -// partial allocation can use stream-ordered cleanup. -template -void SubmitCudaArgmaxScratch(ArgmaxScratchOwners& owners, const QueueLike& queue, Stream stream, - std::size_t need, Runtime& runtime, Launch&& launch) { - const ArgmaxScratchKey key{queue.device.index, reinterpret_cast(stream)}; - owners.Submit( - key, need, [&] { return runtime.IsCapturing(stream); }, - [&](std::size_t bytes) { return runtime.Allocate(bytes, stream); }, - [&](void* p) { runtime.CleanupPartial(p, stream); }, [&](void* p) { runtime.Retire(p); }, - std::forward(launch)); -} - -} // namespace vt::cuda - -#endif // VT_CUDA_ARGMAX_SCRATCH_H_ diff --git a/scripts/check-public-doc-tables.py b/scripts/check-public-doc-tables.py index 590349d4e..82e0669dd 100755 --- a/scripts/check-public-doc-tables.py +++ b/scripts/check-public-doc-tables.py @@ -485,7 +485,12 @@ def features_errors(text: str) -> list[str]: # (works on the 35B gate model, spec-on output token-identical to spec-off, # 1.15x warm against upstream's 1.41x). Re-pinned byte-tight below after the # merge, so the reduction cannot become untracked growth headroom. - "chars": 243570, + # + # 243519 since 2026-08-10 (measured 243519): issue #206 replaces the stale + # pretoken-failure row with the corrected, exact throughput result and its + # still-open latency/VRAM axes. Re-pinned byte-tight after current-main + # reconciliation so the shorter row cannot become untracked headroom. + "chars": 243519, "h2_sections": 11, "long_paragraphs": 82, "oversized_cells": 44, diff --git a/scripts/env-doc-allowlist.txt b/scripts/env-doc-allowlist.txt index 58ce81cb9..5abe69189 100644 --- a/scripts/env-doc-allowlist.txt +++ b/scripts/env-doc-allowlist.txt @@ -16,7 +16,6 @@ VT_ROCM_GEMM_COMPUTE VT_ROCM_GEMV VT_ROCM_HIPBLASLT VLLM_MM_TOWER_PROFILE -VT_ARGMAX_GEOMETRIC_SCRATCH VT_ARCH_TACTIC_STATS VT_ATTN_DECODE_GQA VT_ATTN_DECODE_OPT @@ -76,7 +75,6 @@ VT_GDN_BF16 VT_GDN_CHUNKED VT_GDN_CHUNKO_TRITON VT_GDN_CHUNK_VEC -VT_GDN_DECODE_BF16_VECSTORE VT_GDN_DECODE_NW VT_GDN_DELTAH_TRITON VT_GDN_DELTAH_VEC diff --git a/src/vt/cuda/cuda_gdn.cu b/src/vt/cuda/cuda_gdn.cu index bcda6fda4..0ad4eff66 100644 --- a/src/vt/cuda/cuda_gdn.cu +++ b/src/vt/cuda/cuda_gdn.cu @@ -270,23 +270,6 @@ __device__ inline void Store(__nv_bfloat16* p, int64_t i, float v) { p[i] = __float2bfloat16(v); // round-to-nearest-even, same as host F32ToBF16 } -// Pack exactly eight scalar BF16 RN conversions into one aligned 16-byte -// transaction. The aggregate is alignment-safe; no potentially unaligned -// scalar object is type-punned. -__device__ inline void StoreBf16Vec8(__nv_bfloat16* p, float v0, float v1, - float v2, float v3, float v4, float v5, - float v6, float v7) { - static_assert(sizeof(__nv_bfloat16) * 8 == sizeof(int4)); - static_assert(alignof(int4) == 16); - const float values[8] = {v0, v1, v2, v3, v4, v5, v6, v7}; - const GdnDecodeBf16Pack8 packed = GdnDecodePackBf16Vec8(values); - const int4 transaction{static_cast(packed.words[0]), - static_cast(packed.words[1]), - static_cast(packed.words[2]), - static_cast(packed.words[3])}; - *reinterpret_cast(p) = transaction; -} - template __device__ inline float RoundToStorage(float v); template <> @@ -2795,17 +2778,13 @@ void GdnPackedDecodeKernelCuda(Queue& q, Tensor& out, // a power of two dividing 32 and the launcher only uses NW>1 when BV==32, so a // block is always a whole number of warps) — the xor butterfly stays in-warp. template + bool REGSTATE> __global__ void GdnDecodeFusedKernel(Tout* out, const Tin* q, const Tin* k, const Tin* v, const float* g, const float* beta, TState* state, const int32_t* state_idx, int64_t hk_n, int64_t dk, int64_t hv_n, int64_t dv, int64_t bv, float scale) { static_assert(!REGSTATE || (SWIZZLED && NW == 8), "register state requires the production swizzled NW8 layout"); - static_assert(!BF16_VECSTORE || - (REGSTATE && SWIZZLED && NW == 8 && - std::is_same::value), - "BF16 vector writeback requires BF16 REGSTATE BV16+SWIZZLE+NW8"); const int64_t i_v = blockIdx.x; // value-dim tile const int64_t i_nh = blockIdx.y; // fused (sequence, v-head) const int64_t i_n = i_nh / hv_n; // sequence == decode token index @@ -2921,35 +2900,19 @@ __global__ void GdnDecodeFusedKernel(Tout* out, const Tin* q, const Tin* k, cons #pragma unroll for (int off = 1; off < NW; off <<= 1) o += __shfl_xor_sync(0xffffffffu, o, off); if (vrow < dv && wk == 0) Store(out, (i_n * hv_n + hv) * dv + vrow, o); - if constexpr (REGSTATE && !BF16_VECSTORE) { + if constexpr (REGSTATE) { for (int64_t j = 0; j < 16; ++j) { r[GdnDecodeRegisterSharedColumn(wk, j)] = rr[j]; } } - if constexpr (BF16_VECSTORE) { - // Eligibility proves Dv=Dk=128 and BV16: every lane owns one complete, - // valid 16-element logical slice. Each destination begins at a 32-byte - // boundary, so both 16-byte stores are naturally aligned. - static_assert(sizeof(TState) == 2); -#pragma unroll - for (int pack = 0; pack < 2; ++pack) { - const int64_t column = - GdnDecodeBf16PackLogicalColumn(wk, pack, 0); - const int base = pack * 8; - StoreBf16Vec8(s_head + static_cast(vi) * 128 + column, - rr[base], rr[base + 1], rr[base + 2], rr[base + 3], - rr[base + 4], rr[base + 5], rr[base + 6], rr[base + 7]); - } - } else { - __syncthreads(); + __syncthreads(); - // Coalesced write-back of the updated slice from f32 registers to the - // configured fp16/bf16/fp32 temporal cache. - for (int64_t e = tid; e < tile; e += blockDim.x) { - const int64_t c = e % dk; - const int64_t sc = SWIZZLED ? (c % 16) * 8 + c / 16 : c; - Store(s_head, e, sbh[(e / dk) * sdk + sc]); - } + // Coalesced write-back of the updated slice from f32 registers to the + // configured fp16/bf16/fp32 temporal cache. + for (int64_t e = tid; e < tile; e += blockDim.x) { + const int64_t c = e % dk; + const int64_t sc = SWIZZLED ? (c % 16) * 8 + c / 16 : c; + Store(s_head, e, sbh[(e / dk) * sdk + sc]); } } @@ -2957,8 +2920,7 @@ template void LaunchGdnDecodeFusedNW(cudaStream_t s, Tensor& out, const Tensor& q_in, const Tensor& k, const Tensor& v, const Tensor& g, const Tensor& beta, Tensor& state, const int32_t* state_idx, int64_t n, const GdnArgs& args, - const GdnDecodeLaunchContract& contract, - const GdnDecodeBf16VecstoreCapability& bf16_vecstore) { + const GdnDecodeLaunchContract& contract) { const int64_t hk_n = q_in.shape[1], dk = q_in.shape[2]; const int64_t hv_n = v.shape[1], dv = v.shape[2]; const int64_t bv = contract.value_tile; @@ -2978,37 +2940,13 @@ void LaunchGdnDecodeFusedNW(cudaStream_t s, Tensor& out, const Tensor& q_in, con args.scale); }, [&](const GdnDecodeLaunchContract&) { - if constexpr (std::is_same::value) { - bf16_vecstore.Dispatch( - [&] { - GdnDecodeFusedKernel - <<(contract.block_threads), - contract.shared_bytes, s>>>( - out.Ptr(), q_in.Ptr(), k.Ptr(), - v.Ptr(), g.Ptr(), beta.Ptr(), - state.Ptr(), state_idx, hk_n, dk, hv_n, dv, - bv, args.scale); - }, - [&] { - GdnDecodeFusedKernel - <<(contract.block_threads), - contract.shared_bytes, s>>>( - out.Ptr(), q_in.Ptr(), k.Ptr(), - v.Ptr(), g.Ptr(), beta.Ptr(), - state.Ptr(), state_idx, hk_n, dk, hv_n, dv, - bv, args.scale); - }); - } else { - GdnDecodeFusedKernel - <<(contract.block_threads), - contract.shared_bytes, s>>>( - out.Ptr(), q_in.Ptr(), k.Ptr(), - v.Ptr(), g.Ptr(), beta.Ptr(), - state.Ptr(), state_idx, hk_n, dk, hv_n, dv, bv, - args.scale); - } + GdnDecodeFusedKernel + <<(contract.block_threads), + contract.shared_bytes, s>>>( + out.Ptr(), q_in.Ptr(), k.Ptr(), + v.Ptr(), g.Ptr(), beta.Ptr(), + state.Ptr(), state_idx, hk_n, dk, hv_n, dv, bv, + args.scale); }); } else { GdnDecodeFusedKernel @@ -3033,28 +2971,23 @@ template void LaunchGdnDecodeFused(cudaStream_t s, Tensor& out, const Tensor& q_in, const Tensor& k, const Tensor& v, const Tensor& g, const Tensor& beta, Tensor& state, const int32_t* state_idx, int64_t n, const GdnArgs& args, - const GdnDecodeLaunchContract& contract, - const GdnDecodeBf16VecstoreCapability& bf16_vecstore) { + const GdnDecodeLaunchContract& contract) { switch (contract.lanes_per_row) { case 2: LaunchGdnDecodeFusedNW(s, out, q_in, k, v, g, beta, state, state_idx, - n, args, contract, - bf16_vecstore); + n, args, contract); break; case 4: LaunchGdnDecodeFusedNW(s, out, q_in, k, v, g, beta, state, state_idx, - n, args, contract, - bf16_vecstore); + n, args, contract); break; case 8: LaunchGdnDecodeFusedNW(s, out, q_in, k, v, g, beta, state, state_idx, - n, args, contract, - bf16_vecstore); + n, args, contract); break; default: LaunchGdnDecodeFusedNW(s, out, q_in, k, v, g, beta, state, state_idx, - n, args, contract, - bf16_vecstore); + n, args, contract); break; } } @@ -3065,19 +2998,16 @@ template void LaunchGdnDecodeFusedS(cudaStream_t s, Tensor& out, const Tensor& q_in, const Tensor& k, const Tensor& v, const Tensor& g, const Tensor& beta, Tensor& state, const int32_t* state_idx, int64_t n, const GdnArgs& args, - const GdnDecodeLaunchContract& contract, - const GdnDecodeBf16VecstoreCapability& bf16_vecstore) { + const GdnDecodeLaunchContract& contract) { if (state.dtype == DType::kBF16) LaunchGdnDecodeFused(s, out, q_in, k, v, g, beta, state, state_idx, - n, args, contract, - bf16_vecstore); + n, args, contract); else if (state.dtype == DType::kF16) LaunchGdnDecodeFused(s, out, q_in, k, v, g, beta, - state, state_idx, n, args, contract, - bf16_vecstore); + state, state_idx, n, args, contract); else LaunchGdnDecodeFused(s, out, q_in, k, v, g, beta, state, state_idx, n, args, - contract, bf16_vecstore); + contract); } // Decode dispatch. state_idx == nullptr: compact [n,Hv,Dv,Dk] state (row==i_n). @@ -3112,26 +3042,20 @@ void GdnDecodeFusedCuda(Queue& q, Tensor& out, const Tensor& q_in, const Tensor& args, "gdn_decode"); return; } - const GdnDecodeBf16VecstoreCapability bf16_vecstore{ - std::getenv("VT_GDN_DECODE_BF16_VECSTORE"), contract, dv, dk, nw, - state.dtype == DType::kBF16, true}; if (q_in.dtype == DType::kF32) { if (out.dtype == DType::kF32) LaunchGdnDecodeFusedS(s, out, q_in, k, v, g, beta, state, state_idx, n, - args, contract, bf16_vecstore); + args, contract); else LaunchGdnDecodeFusedS(s, out, q_in, k, v, g, beta, state, state_idx, - n, args, contract, - bf16_vecstore); + n, args, contract); } else { if (out.dtype == DType::kF32) LaunchGdnDecodeFusedS<__nv_bfloat16, float>(s, out, q_in, k, v, g, beta, state, state_idx, - n, args, contract, - bf16_vecstore); + n, args, contract); else LaunchGdnDecodeFusedS<__nv_bfloat16, __nv_bfloat16>( - s, out, q_in, k, v, g, beta, state, state_idx, n, args, contract, - bf16_vecstore); + s, out, q_in, k, v, g, beta, state, state_idx, n, args, contract); } }; DispatchGdnDecodeValueTile(std::getenv("VT_GDN_DECODE_BV"), diff --git a/src/vt/cuda/cuda_sample.cu b/src/vt/cuda/cuda_sample.cu index 86685e169..652988503 100644 --- a/src/vt/cuda/cuda_sample.cu +++ b/src/vt/cuda/cuda_sample.cu @@ -18,13 +18,10 @@ #include #include -#include #include #include #include "vt/backend.h" -#include "vt/cuda/argmax_scratch.h" -#include "vt/cuda/graph_safe_scratch.h" #include "vt/ops.h" namespace vt::cuda { @@ -158,7 +155,6 @@ __global__ void ArgmaxFinalKernel(int64_t* out, const float* part_val, const int float* g_argmax_val = nullptr; int64_t* g_argmax_idx = nullptr; size_t g_argmax_cap = 0; // capacity in elements -ArgmaxScratchOwners g_argmax_geometric_owners; void EnsureArgmaxScratch(size_t elems) { if (elems <= g_argmax_cap) return; @@ -194,34 +190,6 @@ bool FastArgmaxEnabled() { return on; } -const char* ArgmaxGeometricScratchSelector() { - static const char* selector = [] { - const char* value = std::getenv("VT_ARGMAX_GEOMETRIC_SCRATCH"); - return ArgmaxGeometricScratchEnabled(value) ? "1" : nullptr; - }(); - return selector; -} - -struct CudaArgmaxScratchRuntime { - bool IsCapturing(cudaStream_t stream) const { - cudaStreamCaptureStatus status = cudaStreamCaptureStatusNone; - Check(cudaStreamIsCapturing(stream, &status), "argmax scratch capture query"); - return status != cudaStreamCaptureStatusNone; - } - - void* Allocate(size_t bytes, cudaStream_t stream) const { - void* p = nullptr; - Check(cudaMallocAsync(&p, bytes, stream), "argmax geometric scratch allocation"); - return p; - } - - void CleanupPartial(void* p, cudaStream_t stream) const { - Check(cudaFreeAsync(p, stream), "argmax geometric scratch partial cleanup"); - } - - void Retire(void* p) const { RetireGraphScratch(p); } -}; - void GreedyArgmaxCuda(Queue& q, Tensor& token_ids, const Tensor& logits) { const int64_t n = logits.shape[0], v = logits.shape[1]; if (n == 0 || v == 0) return; @@ -239,36 +207,13 @@ void GreedyArgmaxCuda(Queue& q, Tensor& token_ids, const Tensor& logits) { if (bpr > kBlock) bpr = kBlock; // pass 2 reduces bpr partials with kBlock threads if (bpr < 1) bpr = 1; - DispatchArgmaxScratch( - ArgmaxGeometricScratchSelector(), - [&] { - // Exact rollback: retain the original process-global exact-size - // cudaFree/cudaMalloc allocator and launch sequence byte-for-byte. - EnsureArgmaxScratch(static_cast(n) * bpr); - dim3 grid1(static_cast(bpr), static_cast(n)); - ArgmaxPartialKernel<<>>(g_argmax_val, g_argmax_idx, - logits.Ptr(), v, bpr); - ArgmaxFinalKernel<<(n), kBlock, 0, s>>>( - token_ids.Ptr(), g_argmax_val, g_argmax_idx, bpr); - Check(cudaGetLastError(), "greedy_argmax launch"); - }, - [&] { - if (static_cast(n) > - std::numeric_limits::max() / static_cast(bpr)) { - throw std::overflow_error("argmax scratch element-count overflow"); - } - const size_t need = static_cast(n) * static_cast(bpr); - CudaArgmaxScratchRuntime runtime; - SubmitCudaArgmaxScratch( - g_argmax_geometric_owners, q, s, need, runtime, [&](ArgmaxScratchView scratch) { - dim3 grid1(static_cast(bpr), static_cast(n)); - ArgmaxPartialKernel<<>>(scratch.values, scratch.indices, - logits.Ptr(), v, bpr); - ArgmaxFinalKernel<<(n), kBlock, 0, s>>>( - token_ids.Ptr(), scratch.values, scratch.indices, bpr); - Check(cudaGetLastError(), "greedy_argmax launch"); - }); - }); + EnsureArgmaxScratch(static_cast(n) * bpr); + dim3 grid1(static_cast(bpr), static_cast(n)); + ArgmaxPartialKernel<<>>(g_argmax_val, g_argmax_idx, logits.Ptr(), v, + bpr); + ArgmaxFinalKernel<<(n), kBlock, 0, s>>>(token_ids.Ptr(), + g_argmax_val, g_argmax_idx, bpr); + Check(cudaGetLastError(), "greedy_argmax launch"); } // --- compute_probs / compute_logprobs (block-per-row softmax) --------------- diff --git a/src/vt/cuda/gdn_decode_fused.h b/src/vt/cuda/gdn_decode_fused.h index 9b4eb2c05..c98d5a6c5 100644 --- a/src/vt/cuda/gdn_decode_fused.h +++ b/src/vt/cuda/gdn_decode_fused.h @@ -5,7 +5,6 @@ #include #include -#include namespace vt::cuda { @@ -36,13 +35,6 @@ inline constexpr bool GdnDecodeRegstateFlagIsOn(const char* env_value) { return env_value != nullptr && env_value[0] == '1' && env_value[1] == '\0'; } -// Direct BF16 vector writeback is a strict, independently selectable -// discriminator. It is never implied by REGSTATE and every non-exact spelling -// leaves the incumbent shared-memory writeback intact. -inline constexpr bool GdnDecodeBf16VecstoreFlagIsOn(const char* env_value) { - return env_value != nullptr && env_value[0] == '1' && env_value[1] == '\0'; -} - inline constexpr int64_t GdnDecodeStateStride(bool swizzled, int64_t dk, int lanes_per_row) { return dk + (swizzled ? lanes_per_row : 1); @@ -62,67 +54,6 @@ inline constexpr int64_t GdnDecodeRegisterLogicalColumn(int lane, int slot) { return static_cast(lane) * 16 + slot; } -#if defined(__CUDACC__) -__host__ __device__ -#endif -inline constexpr int64_t GdnDecodeBf16PackLogicalColumn(int lane, int pack, - int element) { - return static_cast(lane) * 16 + pack * 8 + element; -} - -#if defined(__CUDACC__) -__host__ __device__ -#endif -inline constexpr int64_t GdnDecodeBf16PackByteOffset(int lane, int pack) { - return GdnDecodeBf16PackLogicalColumn(lane, pack, 0) * 2; -} - -static_assert(GdnDecodeBf16PackByteOffset(0, 0) == 0); -static_assert(GdnDecodeBf16PackByteOffset(0, 1) == 16); -static_assert(GdnDecodeBf16PackByteOffset(7, 1) == 240); - -// Portable representation of the exact 16-byte transaction emitted by the -// CUDA writeback. Keeping conversion and packing here lets the host tests -// validate the bytes that StoreBf16Vec8 actually consumes. -struct alignas(16) GdnDecodeBf16Pack8 { - uint32_t words[4]; -}; - -#if defined(__CUDACC__) -__host__ __device__ -#endif -inline uint16_t GdnDecodeBf16RnBits(float value) { - uint32_t bits = 0; -#if defined(__CUDA_ARCH__) - bits = __float_as_uint(value); -#else - std::memcpy(&bits, &value, sizeof(bits)); -#endif - const uint32_t exponent = bits & 0x7f800000u; - const uint32_t mantissa = bits & 0x007fffffu; - if (exponent == 0x7f800000u && mantissa != 0) { - return static_cast((bits >> 16) | 0x0040u); - } - const uint32_t rounding_bias = 0x7fffu + ((bits >> 16) & 1u); - return static_cast((bits + rounding_bias) >> 16); -} - -#if defined(__CUDACC__) -__host__ __device__ -#endif -inline GdnDecodeBf16Pack8 GdnDecodePackBf16Vec8(const float* values) { - GdnDecodeBf16Pack8 packed{}; - for (int pair = 0; pair < 4; ++pair) { - const uint32_t lo = GdnDecodeBf16RnBits(values[pair * 2]); - const uint32_t hi = GdnDecodeBf16RnBits(values[pair * 2 + 1]); - packed.words[pair] = lo | (hi << 16); - } - return packed; -} - -static_assert(sizeof(GdnDecodeBf16Pack8) == 16); -static_assert(alignof(GdnDecodeBf16Pack8) == 16); - #if defined(__CUDACC__) __host__ __device__ #endif @@ -142,46 +73,6 @@ struct GdnDecodeLaunchContract { bool regstate; }; -// The vector specialization is intentionally narrower than REGSTATE. Keeping -// the dtype and regular-decode predicates explicit here makes host dispatch and -// portable mutation tests agree on every gate. -inline constexpr bool GdnDecodeBf16VecstoreEligible( - const char* env_value, const GdnDecodeLaunchContract& contract, int64_t dv, - int64_t dk, int requested_nw, bool state_is_bf16, - bool regular_fused_decode) { - return GdnDecodeBf16VecstoreFlagIsOn(env_value) && contract.regstate && - contract.swizzled && - contract.selected_tile == GdnDecodeValueTile::kBv16 && - contract.value_tile == 16 && contract.value_tiles == 8 && - contract.lanes_per_row == 8 && contract.block_threads == 128 && - contract.shared_bytes == 9728 && contract.should_launch && dv == 128 && - dk == 128 && requested_nw == 8 && state_is_bf16 && - regular_fused_decode; -} - -struct GdnDecodeBf16VecstoreCapability { - const char* env_value; - GdnDecodeLaunchContract contract; - int64_t dv; - int64_t dk; - int requested_nw; - bool state_is_bf16; - bool regular_fused_decode; - - constexpr bool Eligible() const { - return GdnDecodeBf16VecstoreEligible( - env_value, contract, dv, dk, requested_nw, state_is_bf16, - regular_fused_decode); - } - - template - inline decltype(auto) Dispatch(IncumbentLaunch&& incumbent_launch, - VectorLaunch&& vector_launch) const { - if (Eligible()) return vector_launch(); - return incumbent_launch(); - } -}; - inline constexpr GdnDecodeLaunchContract GdnDecodeLaunchContractFor( const char* bv_env_value, const char* swizzle_env_value, const char* regstate_env_value, int64_t dv, int64_t dk, @@ -233,14 +124,6 @@ inline decltype(auto) DispatchGdnDecodeStateStorage( return shared_launch(contract); } -template -inline decltype(auto) DispatchGdnDecodeBf16Vecstore( - bool vector_eligible, IncumbentLaunch&& incumbent_launch, - VectorLaunch&& vector_launch) { - if (vector_eligible) return vector_launch(); - return incumbent_launch(); -} - // Shared callback dispatcher used by production and portable tests. The // callbacks receive fully resolved geometry and exactly one is invoked. template diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt index 03e6a97e1..3ed1ec824 100644 --- a/tests/CMakeLists.txt +++ b/tests/CMakeLists.txt @@ -894,7 +894,6 @@ target_include_directories(test_gdn_packed_decode_triton PRIVATE ${CMAKE_SOURCE_ # in cuda_gdn.cu and is covered by same-process byte-exact CUDA comparisons. vllm_cpp_add_test(test_gdn_decode_fused vt/test_gdn_decode_fused.cpp) target_include_directories(test_gdn_decode_fused PRIVATE ${CMAKE_SOURCE_DIR}/src) -target_compile_definitions(test_gdn_decode_fused PRIVATE VLLM_CPP_SOURCE_DIR="${CMAKE_SOURCE_DIR}") # CPU-tier contract for the two NUMERICS-NEUTRAL FP4-quant vectorized-load # fast-path flags (VT_FP4_QUANT_FAST + VT_SILU_FP4_FAST, both default OFF / # '1'-opt-in). Reaches the internal header under src/ (not public API); the @@ -1048,8 +1047,6 @@ target_include_directories(test_ops_conv1d_depthwise PRIVATE ${CMAKE_SOURCE_DIR} vllm_cpp_add_test(test_ops_attn_relpos vt/test_ops_attn_relpos.cpp) target_include_directories(test_ops_attn_relpos PRIVATE ${CMAKE_SOURCE_DIR}/src) vllm_cpp_add_test(test_ops_paged_attn vt/test_ops_paged_attn.cpp) -vllm_cpp_add_test(test_argmax_scratch vt/test_argmax_scratch.cpp) -target_compile_definitions(test_argmax_scratch PRIVATE VLLM_CPP_SOURCE_DIR="${CMAKE_SOURCE_DIR}") vllm_cpp_add_test(test_ops_sample vt/test_ops_sample.cpp) vllm_cpp_add_test(test_ops_penalties vt/test_ops_penalties.cpp) diff --git a/tests/vt/test_argmax_scratch.cpp b/tests/vt/test_argmax_scratch.cpp deleted file mode 100644 index 52e8967e2..000000000 --- a/tests/vt/test_argmax_scratch.cpp +++ /dev/null @@ -1,375 +0,0 @@ -#include - -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include - -#include "vt/cuda/argmax_scratch.h" - -namespace { -using vt::cuda::ArgmaxScratchKey; -using vt::cuda::ArgmaxScratchOwners; -using vt::cuda::ArgmaxScratchView; - -void* FakePtr(std::uintptr_t value) { - return reinterpret_cast(value); -} - -TEST_CASE("argmax geometric scratch selector is strict and dispatches exactly once") { - const char* disabled[] = {nullptr, "", "0", "true", "01", "2", " 1", "10", "1x", "1 "}; - for (const char* value : disabled) { - int incumbent = 0; - int candidate = 0; - const int selected = vt::cuda::DispatchArgmaxScratch( - value, - [&] { - ++incumbent; - return 17; - }, - [&] { - ++candidate; - return 29; - }); - CHECK(selected == 17); - CHECK(incumbent == 1); - CHECK(candidate == 0); - } - int incumbent = 0; - int candidate = 0; - CHECK(vt::cuda::DispatchArgmaxScratch( - "1", - [&] { - ++incumbent; - return 17; - }, - [&] { - ++candidate; - return 29; - }) == 29); - CHECK(incumbent == 0); - CHECK(candidate == 1); -} - -TEST_CASE("argmax geometric scratch grows transactionally by powers of two") { - ArgmaxScratchOwners owners; - const ArgmaxScratchKey key{3, 0x1234}; - std::uintptr_t next = 0x1000; - std::vector allocations; - std::vector cleanups; - std::vector retired; - std::vector seen_caps; - auto allocate = [&](std::size_t bytes) { - allocations.push_back(bytes); - next += 0x100; - return FakePtr(next); - }; - auto submit = [&](ArgmaxScratchView view) { seen_caps.push_back(view.capacity); }; - - for (std::size_t need : {1U, 3U, 5U, 9U, 17U, 32U, 7U, 31U}) { - owners.Submit( - key, need, [] { return false; }, allocate, [&](void* p) { cleanups.push_back(p); }, - [&](void* p) { retired.push_back(p); }, submit); - } - CHECK(seen_caps == std::vector{1, 4, 8, 16, 32, 32, 32, 32}); - CHECK(allocations == std::vector{4, 8, 16, 32, 32, 64, 64, 128, 128, 256}); - CHECK(cleanups.empty()); - CHECK(retired == std::vector{FakePtr(0x1100), FakePtr(0x1200), FakePtr(0x1300), - FakePtr(0x1400), FakePtr(0x1500), FakePtr(0x1600), - FakePtr(0x1700), FakePtr(0x1800)}); - const auto diag = owners.Diagnostics(key); - REQUIRE(diag.has_value()); - CHECK(diag->capacity == 32); - CHECK(diag->active_bytes == 32 * (sizeof(float) + sizeof(std::int64_t))); - CHECK(diag->retired_bytes == (1 + 4 + 8 + 16) * (sizeof(float) + sizeof(std::int64_t))); - CHECK(diag->active_bytes + diag->retired_bytes < 2 * diag->active_bytes); -} - -TEST_CASE("argmax geometric scratch preserves the published pair on allocation failure") { - ArgmaxScratchOwners owners; - const ArgmaxScratchKey key{0, 0x44}; - std::uintptr_t next = 0x2000; - std::vector cleanups; - std::vector retired; - int calls = 0; - auto ok_alloc = [&](std::size_t) { - next += 0x100; - return FakePtr(next); - }; - owners.Submit( - key, 3, [] { return false; }, ok_alloc, [&](void* p) { cleanups.push_back(p); }, - [&](void* p) { retired.push_back(p); }, [&](ArgmaxScratchView) { ++calls; }); - const auto before = *owners.Diagnostics(key); - - int allocation_number = 0; - CHECK_THROWS_AS(owners.Submit( - key, 5, [] { return false; }, - [&](std::size_t) -> void* { - ++allocation_number; - if (allocation_number == 2) - throw std::runtime_error("second allocation failed"); - next += 0x100; - return FakePtr(next); - }, - [&](void* p) { cleanups.push_back(p); }, - [&](void* p) { retired.push_back(p); }, [&](ArgmaxScratchView) { ++calls; }), - std::runtime_error); - const auto after = *owners.Diagnostics(key); - CHECK(after.values == before.values); - CHECK(after.indices == before.indices); - CHECK(after.capacity == before.capacity); - CHECK(cleanups == std::vector{FakePtr(0x2300)}); - CHECK(retired.empty()); - CHECK(calls == 1); -} - -TEST_CASE("argmax geometric scratch rejects a null first allocation transactionally") { - ArgmaxScratchOwners owners; - const ArgmaxScratchKey absent{4, 0x404}; - int launches = 0; - int retirements = 0; - CHECK_THROWS_AS(owners.Submit( - absent, 3, [] { return false; }, [](std::size_t) -> void* { return nullptr; }, - [](void*) { FAIL("null first allocation must not need cleanup"); }, - [&](void*) { ++retirements; }, [&](ArgmaxScratchView) { ++launches; }), - std::runtime_error); - CHECK(owners.OwnerCount() == 0); - CHECK_FALSE(owners.Diagnostics(absent).has_value()); - CHECK(retirements == 0); - CHECK(launches == 0); - - const ArgmaxScratchKey existing{4, 0x405}; - std::uintptr_t next = 0x5000; - owners.Submit( - existing, 3, [] { return false; }, - [&](std::size_t) { - next += 0x100; - return FakePtr(next); - }, - [](void*) {}, [&](void*) { ++retirements; }, [&](ArgmaxScratchView) { ++launches; }); - const auto before = *owners.Diagnostics(existing); - CHECK_THROWS_AS( - owners.Submit( - existing, 5, [] { return false; }, [](std::size_t) -> void* { return nullptr; }, - [](void*) { FAIL("null first allocation must not need cleanup"); }, - [&](void*) { ++retirements; }, [&](ArgmaxScratchView) { ++launches; }), - std::runtime_error); - const auto after = *owners.Diagnostics(existing); - CHECK(after.values == before.values); - CHECK(after.indices == before.indices); - CHECK(after.capacity == before.capacity); - CHECK(after.growths == before.growths); - CHECK(retirements == 0); - CHECK(launches == 1); -} - -TEST_CASE("CUDA argmax adapter binds device stream capture allocation and retirement") { - struct FakeQueue { - struct { - int index; - } device; - } queue{{7}}; - struct CapabilityLimitedRuntime { - bool capturing = false; - std::uintptr_t next = 0x6000; - std::vector capture_streams; - std::vector> allocations; - std::vector> partial_cleanups; - std::vector retirements; - - bool IsCapturing(void* stream) { - capture_streams.push_back(reinterpret_cast(stream)); - return capturing; - } - void* Allocate(std::size_t bytes, void* stream) { - allocations.emplace_back(bytes, reinterpret_cast(stream)); - next += 0x100; - return FakePtr(next); - } - void CleanupPartial(void* p, void* stream) { - partial_cleanups.emplace_back(p, reinterpret_cast(stream)); - } - void Retire(void* p) { retirements.push_back(p); } - } runtime; - - ArgmaxScratchOwners owners; - void* const stream = FakePtr(0x777); - std::vector launches; - vt::cuda::SubmitCudaArgmaxScratch(owners, queue, stream, 3, runtime, - [&](ArgmaxScratchView view) { launches.push_back(view); }); - vt::cuda::SubmitCudaArgmaxScratch(owners, queue, stream, 5, runtime, - [&](ArgmaxScratchView view) { launches.push_back(view); }); - CHECK(owners.OwnerCount() == 1); - CHECK(owners.Diagnostics(ArgmaxScratchKey{7, 0x777})->capacity == 8); - CHECK_FALSE(owners.Diagnostics(ArgmaxScratchKey{0, 0x777}).has_value()); - CHECK_FALSE(owners.Diagnostics(ArgmaxScratchKey{7, 0}).has_value()); - CHECK(runtime.capture_streams == std::vector{0x777, 0x777}); - CHECK(runtime.allocations == std::vector>{ - {16, 0x777}, {32, 0x777}, {32, 0x777}, {64, 0x777}}); - CHECK(runtime.partial_cleanups.empty()); - CHECK(runtime.retirements == std::vector{FakePtr(0x6100), FakePtr(0x6200)}); - REQUIRE(launches.size() == 2); - CHECK(launches[0].capacity == 4); - CHECK(launches[1].capacity == 8); - - CapabilityLimitedRuntime captured; - captured.capturing = true; - CHECK_THROWS_AS(vt::cuda::SubmitCudaArgmaxScratch( - owners, queue, FakePtr(0x778), 1, captured, - [](ArgmaxScratchView) { FAIL("capture miss must not launch"); }), - std::runtime_error); - CHECK(captured.capture_streams == std::vector{0x778}); - CHECK(captured.allocations.empty()); - CHECK(captured.retirements.empty()); - CHECK_FALSE(owners.Diagnostics(ArgmaxScratchKey{7, 0x778}).has_value()); -} - -TEST_CASE("CUDA argmax call site binds the capability adapter to graph-safe runtime operations") { - std::ifstream input(std::string(VLLM_CPP_SOURCE_DIR) + "/src/vt/cuda/cuda_sample.cu"); - REQUIRE(input.good()); - const std::string source((std::istreambuf_iterator(input)), - std::istreambuf_iterator()); - const auto runtime_begin = source.find("struct CudaArgmaxScratchRuntime {"); - const auto runtime_end = source.find("\n};", runtime_begin); - REQUIRE(runtime_begin != std::string::npos); - REQUIRE(runtime_end != std::string::npos); - const std::string runtime = source.substr(runtime_begin, runtime_end - runtime_begin); - CHECK(runtime.find("cudaStreamIsCapturing(stream, &status)") != std::string::npos); - CHECK(runtime.find("return status != cudaStreamCaptureStatusNone") != std::string::npos); - CHECK(runtime.find("cudaMallocAsync(&p, bytes, stream)") != std::string::npos); - CHECK(runtime.find("cudaFreeAsync(p, stream)") != std::string::npos); - CHECK(runtime.find("void Retire(void* p) const { RetireGraphScratch(p); }") != - std::string::npos); - - const auto route = source.find("SubmitCudaArgmaxScratch(", runtime_end); - REQUIRE(route != std::string::npos); - const std::string invocation = source.substr(route, 180); - CHECK(invocation.find("g_argmax_geometric_owners, q, s, need, runtime") != std::string::npos); -} - -TEST_CASE("argmax geometric scratch isolates keys and fails capture misses closed") { - ArgmaxScratchOwners owners; - std::uintptr_t next = 0x3000; - int allocations = 0; - int submits = 0; - auto allocate = [&](std::size_t) { - ++allocations; - next += 0x100; - return FakePtr(next); - }; - auto cleanup = [](void*) {}; - auto retire = [](void*) {}; - auto submit = [&](ArgmaxScratchView) { ++submits; }; - const ArgmaxScratchKey a{0, 0x99}; - const ArgmaxScratchKey b{1, 0x99}; - const ArgmaxScratchKey c{0, 0x100}; - owners.Submit(a, 4, [] { return false; }, allocate, cleanup, retire, submit); - owners.Submit(b, 4, [] { return false; }, allocate, cleanup, retire, submit); - owners.Submit(c, 4, [] { return false; }, allocate, cleanup, retire, submit); - CHECK(owners.OwnerCount() == 3); - CHECK(owners.Diagnostics(a)->values != owners.Diagnostics(b)->values); - CHECK(owners.Diagnostics(a)->values != owners.Diagnostics(c)->values); - - int capture_queries = 0; - owners.Submit( - a, 4, - [&] { - ++capture_queries; - return true; - }, - allocate, cleanup, retire, submit); - CHECK(capture_queries == 0); // a warmed hit is capture-safe - CHECK_THROWS_AS(owners.Submit( - a, 5, - [&] { - ++capture_queries; - return true; - }, - allocate, cleanup, retire, submit), - std::runtime_error); - CHECK(capture_queries == 1); - CHECK(allocations == 6); - CHECK(submits == 4); - CHECK(owners.Diagnostics(a)->capacity == 4); - - CHECK_THROWS_AS( - owners.Submit( - ArgmaxScratchKey{7, 0x777}, 1, [] { return true; }, allocate, cleanup, retire, submit), - std::runtime_error); - CHECK(owners.OwnerCount() == 3); // an absent-owner miss changes no state -} - -TEST_CASE("argmax geometric scratch submit lock spans the complete launch callback") { - ArgmaxScratchOwners owners; - const ArgmaxScratchKey key{0, 0x55}; - std::uintptr_t next = 0x4000; - auto allocate = [&](std::size_t) { - next += 0x100; - return FakePtr(next); - }; - owners.Submit( - key, 4, [] { return false; }, allocate, [](void*) {}, [](void*) {}, [](ArgmaxScratchView) {}); - - std::mutex mu; - std::condition_variable cv; - bool first_between_kernels = false; - bool release_first = false; - std::atomic second_entered{false}; - std::thread first([&] { - owners.Submit( - key, 4, [] { return false; }, allocate, [](void*) {}, [](void*) {}, - [&](ArgmaxScratchView) { - std::unique_lock lock(mu); - first_between_kernels = true; - cv.notify_all(); - cv.wait(lock, [&] { return release_first; }); - }); - }); - { - std::unique_lock lock(mu); - cv.wait(lock, [&] { return first_between_kernels; }); - } - std::thread second([&] { - owners.Submit( - key, 4, [] { return false; }, allocate, [](void*) {}, [](void*) {}, - [&](ArgmaxScratchView) { second_entered = true; }); - }); - std::this_thread::sleep_for(std::chrono::milliseconds(20)); - CHECK_FALSE(second_entered.load()); - { - std::lock_guard lock(mu); - release_first = true; - } - cv.notify_all(); - first.join(); - second.join(); - CHECK(second_entered.load()); -} - -TEST_CASE("argmax geometric scratch rejects capacity and byte overflow") { - CHECK(vt::cuda::RoundArgmaxScratchCapacity(17) == 32); - CHECK(vt::cuda::RoundArgmaxScratchCapacity(32) == 32); - CHECK_THROWS_AS(vt::cuda::RoundArgmaxScratchCapacity(std::numeric_limits::max()), - std::overflow_error); - ArgmaxScratchOwners owners; - int allocations = 0; - CHECK_THROWS_AS(owners.Submit( - ArgmaxScratchKey{0, 1}, std::numeric_limits::max() / 4 + 1, - [] { return false; }, - [&](std::size_t) { - ++allocations; - return FakePtr(1); - }, - [](void*) {}, [](void*) {}, [](ArgmaxScratchView) {}), - std::overflow_error); - CHECK(allocations == 0); -} -} // namespace diff --git a/tests/vt/test_gdn_decode_fused.cpp b/tests/vt/test_gdn_decode_fused.cpp index bf62468fb..468618701 100644 --- a/tests/vt/test_gdn_decode_fused.cpp +++ b/tests/vt/test_gdn_decode_fused.cpp @@ -2,12 +2,7 @@ // value tile; the CUDA recurrence remains in cuda_gdn.cu. #include -#include -#include -#include #include -#include -#include #include "vt/cuda/gdn_decode_fused.h" @@ -17,13 +12,6 @@ using vt::cuda::GdnDecodeLaunchContractFor; using vt::cuda::GdnDecodeRegisterLogicalColumn; using vt::cuda::GdnDecodeRegisterSharedColumn; using vt::cuda::GdnDecodeRegstateFlagIsOn; -using vt::cuda::DispatchGdnDecodeBf16Vecstore; -using vt::cuda::GdnDecodeBf16PackByteOffset; -using vt::cuda::GdnDecodeBf16PackLogicalColumn; -using vt::cuda::GdnDecodeBf16VecstoreEligible; -using vt::cuda::GdnDecodeBf16VecstoreFlagIsOn; -using vt::cuda::GdnDecodeBf16VecstoreCapability; -using vt::cuda::GdnDecodePackBf16Vec8; using vt::cuda::GdnDecodeSharedColumn; using vt::cuda::GdnDecodeStateStride; using vt::cuda::GdnDecodeValueTile; @@ -291,195 +279,3 @@ TEST_CASE("fused GDN decode register slots map bijectively to logical and shared CHECK(shared_seen[c]); } } - -namespace { - -uint16_t ReferenceBf16Rn(float value) { - uint32_t bits = 0; - static_assert(sizeof(bits) == sizeof(value)); - std::memcpy(&bits, &value, sizeof(bits)); - const uint32_t exponent = bits & 0x7f800000u; - const uint32_t mantissa = bits & 0x007fffffu; - if (exponent == 0x7f800000u && mantissa != 0) { - return static_cast((bits >> 16) | 0x0040u); - } - const uint32_t rounding_bias = 0x7fffu + ((bits >> 16) & 1u); - return static_cast((bits + rounding_bias) >> 16); -} - -} // namespace - -TEST_CASE("VT_GDN_DECODE_BF16_VECSTORE selects only exact one") { - CHECK(GdnDecodeBf16VecstoreFlagIsOn("1")); - constexpr const char* invalid[] = { - nullptr, "", "0", "01", "10", "1 ", " 1", "+1", "1x", "true", "yes"}; - for (const char* value : invalid) { - CAPTURE(value == nullptr ? "" : value); - CHECK_FALSE(GdnDecodeBf16VecstoreFlagIsOn(value)); - } -} - -TEST_CASE("GDN BF16 vector writeback requires every production predicate") { - const auto production = - GdnDecodeLaunchContractFor("16", "1", "1", 128, 128, 8); - auto eligible = [&](const char* vecstore, const auto& contract, int64_t dv, - int64_t dk, int nw, bool state_is_bf16, - bool regular_fused_decode) { - return GdnDecodeBf16VecstoreEligible( - vecstore, contract, dv, dk, nw, state_is_bf16, - regular_fused_decode); - }; - CHECK(eligible("1", production, 128, 128, 8, true, true)); - CHECK_FALSE(eligible(nullptr, production, 128, 128, 8, true, true)); - CHECK_FALSE(eligible("1x", production, 128, 128, 8, true, true)); - CHECK_FALSE(eligible("1", GdnDecodeLaunchContractFor("32", "1", "1", 128, 128, 8), - 128, 128, 8, true, true)); - CHECK_FALSE(eligible("1", GdnDecodeLaunchContractFor("16", "0", "1", 128, 128, 8), - 128, 128, 8, true, true)); - CHECK_FALSE(eligible("1", GdnDecodeLaunchContractFor("16", "1", "0", 128, 128, 8), - 128, 128, 8, true, true)); - CHECK_FALSE(eligible("1", production, 127, 128, 8, true, true)); - CHECK_FALSE(eligible("1", production, 128, 127, 8, true, true)); - CHECK_FALSE(eligible("1", production, 128, 128, 4, true, true)); - CHECK_FALSE(eligible("1", production, 128, 128, 8, false, true)); - CHECK_FALSE(eligible("1", production, 128, 128, 8, true, false)); - - // Corrupt exactly one field at a time. Tests made from another factory - // contract can leave several correlated fields false and therefore survive - // deletion of the predicate they claim to cover. - auto corrupted = production; - corrupted.regstate = false; - CHECK_FALSE(eligible("1", corrupted, 128, 128, 8, true, true)); - corrupted = production; - corrupted.swizzled = false; - CHECK_FALSE(eligible("1", corrupted, 128, 128, 8, true, true)); - corrupted = production; - corrupted.selected_tile = GdnDecodeValueTile::kBv32; - CHECK_FALSE(eligible("1", corrupted, 128, 128, 8, true, true)); - corrupted = production; - corrupted.value_tile = 32; - CHECK_FALSE(eligible("1", corrupted, 128, 128, 8, true, true)); - corrupted = production; - corrupted.value_tiles = 4; - CHECK_FALSE(eligible("1", corrupted, 128, 128, 8, true, true)); - corrupted = production; - corrupted.lanes_per_row = 4; - CHECK_FALSE(eligible("1", corrupted, 128, 128, 8, true, true)); - corrupted = production; - corrupted.block_threads = 64; - CHECK_FALSE(eligible("1", corrupted, 128, 128, 8, true, true)); - corrupted = production; - corrupted.shared_bytes = 9727; - CHECK_FALSE(eligible("1", corrupted, 128, 128, 8, true, true)); - corrupted = production; - corrupted.should_launch = false; - CHECK_FALSE(eligible("1", corrupted, 128, 128, 8, true, true)); -} - -TEST_CASE("GDN BF16 vector writeback dispatches exactly one callback") { - auto select = [](bool vector_eligible) { - int incumbent_calls = 0; - int vector_calls = 0; - const int selected = DispatchGdnDecodeBf16Vecstore( - vector_eligible, - [&] { - ++incumbent_calls; - return 0; - }, - [&] { - ++vector_calls; - return 1; - }); - CHECK(incumbent_calls + vector_calls == 1); - return selected; - }; - CHECK(select(false) == 0); - CHECK(select(true) == 1); -} - -TEST_CASE("GDN BF16 vector packs cover the aligned 128-column state row") { - bool seen[128] = {}; - for (int lane = 0; lane < 8; ++lane) { - const int64_t first = GdnDecodeBf16PackLogicalColumn(lane, 0, 0); - const int64_t second = GdnDecodeBf16PackLogicalColumn(lane, 1, 0); - CHECK(GdnDecodeBf16PackByteOffset(lane, 0) % 16 == 0); - CHECK(GdnDecodeBf16PackByteOffset(lane, 1) % 16 == 0); - CHECK(second == first + 8); - for (int pack = 0; pack < 2; ++pack) { - for (int element = 0; element < 8; ++element) { - const int64_t column = - GdnDecodeBf16PackLogicalColumn(lane, pack, element); - REQUIRE(column >= 0); - REQUIRE(column < 128); - CHECK_FALSE(seen[column]); - seen[column] = true; - } - } - } - for (bool visited : seen) CHECK(visited); - - const std::array values = { - 0.0f, -0.0f, 1.0f, -2.0f, 0.333251953125f, 65504.0f, - 1.00390625f, 1.01171875f, -1.00390625f, -1.01171875f, - 1.0e-20f, -1.0e-20f, 3.14159265f, -2.71828183f, 42.5f, -99.75f}; - for (int pack = 0; pack < 2; ++pack) { - const auto packed = GdnDecodePackBf16Vec8(values.data() + pack * 8); - for (int pair = 0; pair < 4; ++pair) { - const uint16_t lo = static_cast(packed.words[pair]); - const uint16_t hi = static_cast(packed.words[pair] >> 16); - CHECK(lo == ReferenceBf16Rn(values[pack * 8 + pair * 2])); - CHECK(hi == ReferenceBf16Rn(values[pack * 8 + pair * 2 + 1])); - } - } -} - -TEST_CASE("GDN BF16 vector capability binds eligibility and launch") { - const auto production = - GdnDecodeLaunchContractFor("16", "1", "1", 128, 128, 8); - auto select = [&](const char* env, bool state_is_bf16) { - int incumbent = 0; - int vector = 0; - const GdnDecodeBf16VecstoreCapability capability{ - env, production, 128, 128, 8, state_is_bf16, true}; - const int result = capability.Dispatch( - [&] { return ++incumbent, 0; }, [&] { return ++vector, 1; }); - CHECK(incumbent + vector == 1); - return result; - }; - CHECK(select("1", true) == 1); - CHECK(select("0", true) == 0); - CHECK(select("1", false) == 0); -} - -TEST_CASE("CUDA GDN call site uses the tested vector capability and pack seam") { - std::ifstream input(std::string(VLLM_CPP_SOURCE_DIR) + - "/src/vt/cuda/cuda_gdn.cu"); - REQUIRE(input.good()); - const std::string source((std::istreambuf_iterator(input)), - std::istreambuf_iterator()); - - const auto store_begin = source.find("__device__ inline void StoreBf16Vec8("); - const auto store_end = source.find("\n}", store_begin); - REQUIRE(store_begin != std::string::npos); - REQUIRE(store_end != std::string::npos); - const std::string store = source.substr(store_begin, store_end - store_begin); - CHECK(store.find("GdnDecodePackBf16Vec8(values)") != std::string::npos); - CHECK(store.find("*reinterpret_cast(p) = transaction") != - std::string::npos); - - const auto kernel = source.find("if constexpr (BF16_VECSTORE) {"); - REQUIRE(kernel != std::string::npos); - const std::string writeback = source.substr(kernel, 900); - CHECK(writeback.find("for (int pack = 0; pack < 2; ++pack)") != - std::string::npos); - CHECK(writeback.find("StoreBf16Vec8(") != std::string::npos); - - const auto route = source.find( - "const GdnDecodeBf16VecstoreCapability bf16_vecstore{"); - REQUIRE(route != std::string::npos); - const std::string invocation = source.substr(route, 260); - CHECK(invocation.find("VT_GDN_DECODE_BF16_VECSTORE") != std::string::npos); - CHECK(invocation.find("state.dtype == DType::kBF16, true") != - std::string::npos); - CHECK(source.find("bf16_vecstore.Dispatch(") != std::string::npos); -} diff --git a/tests/vt/test_ops_gdn.cpp b/tests/vt/test_ops_gdn.cpp index 0b842ae93..f624887c8 100644 --- a/tests/vt/test_ops_gdn.cpp +++ b/tests/vt/test_ops_gdn.cpp @@ -1869,7 +1869,6 @@ struct GdnDecodeTestEnv { const char* swizzle; const char* regstate; const char* nw; - const char* bf16_vecstore = nullptr; }; void SetOrUnsetEnv(const char* name, const char* value) { @@ -1915,7 +1914,6 @@ void RunGdnDecodeExactCase(int64_t dv, DType io_dtype, DType state_dtype, SetOrUnsetEnv("VT_GDN_DECODE_BV", env.bv); SetOrUnsetEnv("VT_GDN_DECODE_SWIZZLE", env.swizzle); SetOrUnsetEnv("VT_GDN_DECODE_REGSTATE", env.regstate); - SetOrUnsetEnv("VT_GDN_DECODE_BF16_VECSTORE", env.bf16_vecstore); DeviceTensor dq(gpu, gq.q, io_dtype, {n, hk, dk}, q.data()); DeviceTensor dkt(gpu, gq.q, io_dtype, {n, hk, dk}, k.data()); DeviceTensor dvt(gpu, gq.q, io_dtype, {n, hv, dv}, v.data()); @@ -1946,7 +1944,6 @@ void RunGdnDecodeExactCase(int64_t dv, DType io_dtype, DType state_dtype, run(incumbent, out_incumbent, state_incumbent); run(candidate, out_candidate, state_candidate); unsetenv("VT_GDN_DECODE_REGSTATE"); - unsetenv("VT_GDN_DECODE_BF16_VECSTORE"); unsetenv("VT_GDN_DECODE_SWIZZLE"); unsetenv("VT_GDN_DECODE_BV"); unsetenv("VT_GDN_DECODE_NW"); @@ -3541,28 +3538,6 @@ TEST_CASE("CUDA fused GDN decode register-state schedule is byte-exact") { nw4_shared, nw4_reg, 6440); } -TEST_CASE("CUDA fused GDN decode BF16 vector writeback is byte-exact") { - if (!HasCuda()) { - MESSAGE("no CUDA backend registered; skipping"); - return; - } - constexpr GdnDecodeTestEnv incumbent{"16", "1", "1", "8", "0"}; - constexpr GdnDecodeTestEnv vectorized{"16", "1", "1", "8", "1"}; - RunGdnDecodeExactCase(128, DType::kBF16, DType::kBF16, false, - incumbent, vectorized, 6450); - RunGdnDecodeExactCase(128, DType::kBF16, DType::kBF16, true, - incumbent, vectorized, 6460); - - // The selector is ineligible for independent fp16/fp32 cache dtypes and a - // partial value dimension; those configurations must remain incumbent. - RunGdnDecodeExactCase(128, DType::kBF16, DType::kF16, true, - incumbent, vectorized, 6470); - RunGdnDecodeExactCase(128, DType::kF32, DType::kF32, true, - incumbent, vectorized, 6480); - RunGdnDecodeExactCase(37, DType::kBF16, DType::kBF16, true, - incumbent, vectorized, 6490); -} - TEST_CASE("CUDA public GdnDecode binds decode layout to exact production geometry") { if (!HasCuda()) { MESSAGE("no CUDA backend registered; skipping"); @@ -3582,10 +3557,8 @@ TEST_CASE("CUDA public GdnDecode binds decode layout to exact production geometr DType::kBF16); const auto g = RandomF32(static_cast(n * hv), 6433, -0.5f, -0.01f); const auto beta = RandomF32(static_cast(n * hv), 6434, 0.1f, 0.9f); - const auto state = - Pack(RandomF32(static_cast(n * hv * dv * dk), 6435, -0.25f, - 0.25f), - DType::kBF16); + const auto state = RandomF32(static_cast(n * hv * dv * dk), 6435, + -0.25f, 0.25f); Backend& gpu = vt::GetBackend(DeviceType::kCUDA); QueueGuard gq(gpu); @@ -3596,7 +3569,7 @@ TEST_CASE("CUDA public GdnDecode binds decode layout to exact production geometr DeviceTensor dvt(gpu, gq.q, DType::kBF16, {n, hv, dv}, v.data()); DeviceTensor dg(gpu, gq.q, DType::kF32, {n, hv}, g.data()); DeviceTensor dbeta(gpu, gq.q, DType::kF32, {n, hv}, beta.data()); - DeviceTensor dst(gpu, gq.q, DType::kBF16, {n, hv, dv, dk}, state.data()); + DeviceTensor dst(gpu, gq.q, DType::kF32, {n, hv, dv, dk}, state.data()); DeviceTensor dout(gpu, gq.q, DType::kBF16, {n, hv, dv}); gpu.Synchronize(gq.q); @@ -3604,40 +3577,33 @@ TEST_CASE("CUDA public GdnDecode binds decode layout to exact production geometr const char* bv_env; const char* swizzle_env; const char* regstate_env; - const char* vecstore_env; bool register_state; - bool vector_writeback; unsigned grid_x; unsigned block_x; unsigned shared_bytes; }; constexpr ExpectedGeometry cases[] = { - {nullptr, "1", "1", "1", false, false, 4, 256, 17536}, - {"16", "0", "1", "1", false, false, 8, 128, 9280}, - {"16", "1", "0", "1", false, false, 8, 128, 9728}, - {"16", "1", "1", "0", true, false, 8, 128, 9728}, - {"16", "1", "1", "1", true, true, 8, 128, 9728}, - {"16", "1", "1", "1x", true, false, 8, 128, 9728}, - {"16", "1", "1x", "1", false, false, 8, 128, 9728}, - {"16", "1x", "1", "1", false, false, 8, 128, 9280}, - {"32", "1", "1", "1", false, false, 4, 256, 17536}, + {nullptr, "1", "1", false, 4, 256, 17536}, + {"16", "0", "1", false, 8, 128, 9280}, + {"16", "1", "0", false, 8, 128, 9728}, + {"16", "1", "1", true, 8, 128, 9728}, + {"16", "1", "1x", false, 8, 128, 9728}, + {"16", "1x", "1", false, 8, 128, 9280}, + {"32", "1", "1", false, 4, 256, 17536}, }; const cudaStream_t stream = static_cast(gq.q.handle); setenv("VT_GDN_DECODE_NW", "8", 1); void* shared_production_function = nullptr; - void* register_production_function = nullptr; for (const auto& expected : cases) { CAPTURE(expected.bv_env == nullptr ? "" : expected.bv_env); CAPTURE(expected.swizzle_env); CAPTURE(expected.regstate_env); - CAPTURE(expected.vecstore_env); if (expected.bv_env == nullptr) unsetenv("VT_GDN_DECODE_BV"); else setenv("VT_GDN_DECODE_BV", expected.bv_env, 1); setenv("VT_GDN_DECODE_SWIZZLE", expected.swizzle_env, 1); setenv("VT_GDN_DECODE_REGSTATE", expected.regstate_env, 1); - setenv("VT_GDN_DECODE_BF16_VECSTORE", expected.vecstore_env, 1); CudaGraphGuard captured; REQUIRE(cudaStreamBeginCapture(stream, cudaStreamCaptureModeThreadLocal) == cudaSuccess); @@ -3667,15 +3633,6 @@ TEST_CASE("CUDA public GdnDecode binds decode layout to exact production geometr REQUIRE(shared_production_function != nullptr); CHECK(params.func != shared_production_function); } - if (expected.register_state && !expected.vector_writeback && - std::string(expected.regstate_env) == "1" && - std::string(expected.vecstore_env) == "0") { - register_production_function = params.func; - } - if (expected.vector_writeback) { - REQUIRE(register_production_function != nullptr); - CHECK(params.func != register_production_function); - } CHECK(params.gridDim.x == expected.grid_x); CHECK(params.gridDim.y == static_cast(n * hv)); CHECK(params.gridDim.z == 1); @@ -3685,7 +3642,6 @@ TEST_CASE("CUDA public GdnDecode binds decode layout to exact production geometr CHECK(params.sharedMemBytes == expected.shared_bytes); } unsetenv("VT_GDN_DECODE_REGSTATE"); - unsetenv("VT_GDN_DECODE_BF16_VECSTORE"); unsetenv("VT_GDN_DECODE_SWIZZLE"); unsetenv("VT_GDN_DECODE_BV"); unsetenv("VT_GDN_DECODE_NW"); From db3336fb8517d5365c088611704c12dc783ef708 Mon Sep 17 00:00:00 2001 From: Richard Palethorpe Date: Sun, 9 Aug 2026 23:55:49 +0100 Subject: [PATCH 15/17] spec(KERNEL-SSM-MAMBA): isolate BF16 cold-prefill TTFT Ground the remaining Qwen3.5 TTFT deficit in the first request wave and specify a strict opt-in pre-serve BF16 warmup discriminator mirroring vLLM's mandatory profile run. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: Codex:gpt-5.6-sol [Codex] --- ...qwen35-bf16-pre-serve-warmup-2026-08-09.md | 96 +++++++++++++++++++ 1 file changed, 96 insertions(+) create mode 100644 .agents/specs/sm120-qwen35-bf16-pre-serve-warmup-2026-08-09.md diff --git a/.agents/specs/sm120-qwen35-bf16-pre-serve-warmup-2026-08-09.md b/.agents/specs/sm120-qwen35-bf16-pre-serve-warmup-2026-08-09.md new file mode 100644 index 000000000..aabf3976b --- /dev/null +++ b/.agents/specs/sm120-qwen35-bf16-pre-serve-warmup-2026-08-09.md @@ -0,0 +1,96 @@ +# sm_120 Qwen3.5 BF16 pre-serve prefill warmup + +**Lifecycle:** `SPIKED`; implementation and GPU A/B pending + +**Owner rows:** `ROAD-V1-C2-LOCAL-BF16`, `KERNEL-SSM-MAMBA` + +**Scope:** Qwen3.5 plain-BF16 CUDA initialization only. No default, release, +server-startup SLA, NVFP4, CPU, 27B/35B or pinned-vLLM performance claim is +authorized by this spike. + +## Ground and falsifiable hypothesis + +Corrected atomic/pretoken comparison `da449a88f` leaves local TTFT 1016.513 ms +versus pinned vLLM 936.589 ms. A current accepted-stack diagnostic at +`6c7e64b90` is exact and reports 6865.95 tok/s, 1008.71 ms TTFT, 34.33 ms TPOT +and 5368.08 ms E2E. Its internal means are 133.406 ms intake, 401.197 ms queue +and 473.969 ms prefill. Pinned vLLM reports 338.960 ms queue and 449.291 ms +prefill; its comparable inferred intake is 154.011 ms. + +The gap is isolated to the first request wave, not steady service: + +| request IDs | local queue + prefill | vLLM queue + prefill | local - vLLM | +|---|---:|---:|---:| +| 0-31 | 2086.086 ms | 1625.343 ms | **+460.743 ms** | +| 32-63 | 478.327 ms | 490.471 ms | -12.144 ms | +| 64-95 | 473.912 ms | 510.115 ms | -36.203 ms | +| 96-127 | 462.336 ms | 527.074 ms | -64.738 ms | + +Thus every post-first wave is already faster; the mean deficit is first-touch +startup cost. This matches the execution-chain difference. Pinned vLLM always +runs `_dummy_run(self.max_num_tokens, is_profile=True)` and synchronizes during +memory profiling (`gpu_worker.py:468-495`, `gpu_model_runner.py:6411-6483`), +then warms declared compile sizes and kernels before serving +(`gpu_worker.py:671-708`). Local `LoadedEngine::WarmupKernels` is invoked before +the async frontend exists, but returns unless the model is NVFP4 W4A4 and three +FP4 flags are on (`src/vllm/entrypoints/model_loader.cpp:862-918`). Plain BF16 +therefore serves its first real prefill cold. + +**Hypothesis:** one pre-serve dummy BF16 prefill at the configured +`max_num_batched_tokens` moves kernel/library/pool first-touch work out of user +TTFT, matching vLLM initialization. It must reduce the first-wave mean by at +least 250 ms and aggregate TTFT by at least 5%, without changing later-wave +queue+prefill, throughput, TPOT/E2E, tokens, steady memory or scheduler state. + +## Exact experimental contract + +Add strict opt-in `VT_BF16_PRE_SERVE_WARMUP=1`; only exact `"1"` selects it. +Eligibility requires CUDA, a loaded non-NVFP4 Qwen3.5 model, positive +`max_num_batched_tokens`, and the synchronous pre-serve phase before +`async_engine()` construction. Unset, `0`, invalid values, CPU, NVFP4 and every +other model retain byte-identical initialization. + +Reuse the existing `WarmupKernels` dummy-request lifecycle rather than adding a +benchmark-only warmup. Choose one valid non-special tokenizer token, construct a +prompt of exactly `max_num_batched_tokens`, request one greedy token with EOS +ignored, drive `engine_core_.step()` until completion, drain the finished marker, +and assert both unfinished and finished scheduler state are empty. Use a distinct +request id. Do not start AsyncLLM, alter benchmark timers, change scheduling, +retain output, or special-case the benchmark executable. The warmup must use the +production model/runner/KV paths and must release its blocks before serving. + +## Red-first, mutation, and correctness gates + +Extract a portable selector/lifecycle seam and first capture RED for the missing +BF16 arm. Tests cover strict parsing, BF16-vs-NVFP4/model/backend eligibility, +exact prompt length, one request, greedy one-token parameters, step-until-done, +the extra finished drain, and empty final scheduler state. Mutations must kill a +permissive selector, BF16 hardwired false, wrong prompt length, skipped final +drain, more than one generated token, and benchmark-only routing. Existing FP4 +warmup tests and defaults remain unchanged. + +Operator CUDA gates: focused warmup/loader tests, Qwen GDN and scheduler tests, +clean `-Werror`, then exact 128x128 Qwen output hash +`be20ffbceb61f0264ca21d972bfc5fc51e855e64f2b945de71669cae666aa702` in +both arms. A second request after warmup must reuse a clean scheduler/block state. + +## Measurement and disposition + +Use one immutable binary, one GPU lock and the standard cgroup. Run +`COLD-a -> WARM-a -> WARM-b -> COLD-b`; only the selector differs. Record load +time separately because moving work before serving is a startup trade, not free +compute. Capture `VT_TTFT_DUMP=1` and total/output throughput, TTFT, TPOT/ITL, +E2E, GPU peak/resident memory and host PSS. All tokens must match before timing. + +Accept as an opt-in only if both WARM raw first-wave queue+prefill means beat +both COLD controls, their mean improves at least 250 ms, aggregate TTFT improves +at least 5%, IDs 32-127 do not regress beyond 1%, and throughput, TPOT, E2E, +steady VRAM/PSS and correctness do not regress. Report added load time and peak +initialization memory. A pass remains default-OFF until 27B/35B correctness and +startup-memory gates justify mirroring vLLM's default. Any miss removes product +and tests while retaining this outcome. The next structural target on failure is +the first-wave chunk scheduling/shape sequence, not another decode microkernel. + +## Outcome + +Pending fresh RED-first implementation, mutation review and operator GPU A/B. From fb0fc14db33135ae6fa69a99b5fd6996a6a17597 Mon Sep 17 00:00:00 2001 From: Richard Palethorpe Date: Mon, 10 Aug 2026 00:03:06 +0100 Subject: [PATCH 16/17] perf(KERNEL-SSM-MAMBA): reject BF16 pre-serve warmup Record the exact cold-versus-warm TTFT discriminator and remove the opt-in after it misses both prespecified magnitude gates. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: Codex:gpt-5.6-sol [Codex] --- .agents/NOW.md | 5 ++-- ...qwen35-bf16-pre-serve-warmup-2026-08-09.md | 23 +++++++++++++++++-- 2 files changed, 24 insertions(+), 4 deletions(-) diff --git a/.agents/NOW.md b/.agents/NOW.md index 3462bf4f4..e6c288471 100644 --- a/.agents/NOW.md +++ b/.agents/NOW.md @@ -21,7 +21,7 @@ Work: exact-chunks on main `1ce0d662b`; sm_120 measured at `3d2581551`. | MiniMax-H3 | **PRUNED ckpts RUN (#241): Q8_0 renders, seam 0.9941** | same-binary A/B | | Kimi-Linear-48B | 122/128 held; grouped router parallelised, e2e NOT ESTABLISHED | ckpt is tiktoken-only: no warm server | | 35B binding grid | @`a0fa12c7`: **flat 0.935-0.979x, NO c2/c8 weak cell** (CoV <0.81%); mem PSS 3.81x | Attribute the flat ~5% mid-band; TTFT c2 0.872x | -| Qwen3.5-4B sm_120 | Exact: tput **1.0283x**; TTFT/TPOT/E2E **1.0853/1.0165/1.0288x slower**; VRAM +118.7 MiB | Profile TTFT split | +| Qwen3.5-4B sm_120 | Exact; tput **1.0283x**. Slower: TTFT/TPOT/E2E **1.0853/1.0165/1.0288x**; VRAM +118.7 MiB | Profile first wave | | RPi5 A76 CPU | **R5 asm GREEN; llama NOT MET**: 0.461x pf, 0.653x dec | W6: BF16 GEMM | | MXFP4 parity | c1 1.020, c2-c8 0.962-0.969. **#82 CLOSED: ptxas-lineage REFUTED** | TERMINAL: at parity | | SERVE-ASYNC-DENSE-MIRROR | **LANDED+VERIFIED** (`f9c969ae`): async mirror, dense Qwen3; SACRED 184/184 | Sibling scope one-liner | @@ -53,7 +53,8 @@ latency/memory on every axis, both gate models, reproduced 2–3x idle. See of roof. Dense-marlin +0.5%; Triton-AOT GDN a WASH. 2. **Spike the Parakeet encoder row** (vLLM: `nano_nemotron_vl.py`; the transducer half is NOT in vLLM: separate call). -3. **Qwen3.5-4B #206:** tput +2.83%; latency/VRAM open. Profile TTFT split. +3. **Qwen3.5-4B #206:** exact, tput +2.83%; latency/VRAM open. Profile first + wave; argmax/writeback closed. 2. **Merge the invocation-parity prevention** (CI guard + AGENTS.md checklist); CUDA build-verify the byte-exact `kGemvHeuristicAlgos` refactor on dgx. 4. **Restore `local-ai-worker`** on dgx at campaign end (`--restart=always`). diff --git a/.agents/specs/sm120-qwen35-bf16-pre-serve-warmup-2026-08-09.md b/.agents/specs/sm120-qwen35-bf16-pre-serve-warmup-2026-08-09.md index aabf3976b..4b1d4826a 100644 --- a/.agents/specs/sm120-qwen35-bf16-pre-serve-warmup-2026-08-09.md +++ b/.agents/specs/sm120-qwen35-bf16-pre-serve-warmup-2026-08-09.md @@ -1,6 +1,6 @@ # sm_120 Qwen3.5 BF16 pre-serve prefill warmup -**Lifecycle:** `SPIKED`; implementation and GPU A/B pending +**Lifecycle:** `REJECTED/REMOVED`; directionally positive but below hard gates **Owner rows:** `ROAD-V1-C2-LOCAL-BF16`, `KERNEL-SSM-MAMBA` @@ -93,4 +93,23 @@ the first-wave chunk scheduling/shape sequence, not another decode microkernel. ## Outcome -Pending fresh RED-first implementation, mutation review and operator GPU A/B. +**REJECTED and removed, 2026-08-10.** Implementation `ef1f404f1` plus review +repair `120b36a0e` passed fresh mutation re-review and the operator's rebuilt +5-case / 42-assertion portable gate. The first exact GPU pair produced identical +token SHA-256 +`be20ffbceb61f0264ca21d972bfc5fc51e855e64f2b945de71669cae666aa702`: + +| Arm | total / output tok/s | TTFT | TPOT | E2E | first-wave queue+prefill | +|---|---:|---:|---:|---:|---:| +| COLD | 6864.21 / 759.03 | 1008.58 ms | 34.34 ms | 5369.49 ms | 2084.065 ms | +| WARM | 6902.89 / 763.30 | 973.79 ms | 34.37 ms | 5339.26 ms | 1941.388 ms | + +Warmup is directionally positive (+0.564% total throughput, -3.45% TTFT, +-0.563% E2E), but the first-wave gain is only **142.676 ms** versus the required +250 ms and aggregate TTFT improves only **3.45%** versus the required 5%. +TPOT also rises 0.03 ms. One leg is sufficient to reject on both hard magnitude +gates; repetitions cannot turn this measured arm into an accepted result without +changing the prespecified thresholds. Raw roots are +`/tmp/qwen35-bf16-warmup-{cold,warm}-120b36a0e.*`. Product, selector and tests +are removed. The remaining TTFT target is the first-wave chunk scheduling/shape +sequence after first-touch warmup, as this spec prescribed. From a0a8a5b6305ecee1a24354a1b430ea15fd3be682 Mon Sep 17 00:00:00 2001 From: Richard Palethorpe Date: Mon, 10 Aug 2026 15:03:25 +0100 Subject: [PATCH 17/17] test(records): sync STATUS ratchet ceiling (#206) FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: Codex:gpt-5.6-sol [Codex] --- tests/scripts/test_check_public_doc_tables.py | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/tests/scripts/test_check_public_doc_tables.py b/tests/scripts/test_check_public_doc_tables.py index 7b7fb0509..0aba4d521 100644 --- a/tests/scripts/test_check_public_doc_tables.py +++ b/tests/scripts/test_check_public_doc_tables.py @@ -475,7 +475,7 @@ def test_the_status_ratchet_only_ever_moves_down(self) -> None: 243584 over 48 commits), so the ceiling never needs to rise. """ ceiling = { - "chars": 243578, + "chars": 243519, "h2_sections": 11, "long_paragraphs": 82, "oversized_cells": 44,