diff --git a/.agents/NOW.md b/.agents/NOW.md index 4c3239905..51eb85483 100644 --- a/.agents/NOW.md +++ b/.agents/NOW.md @@ -8,7 +8,7 @@ characters. ## Live claims -Work: exact-chunks on main `1ce0d662b`; sm_120 measured at `3d2581551`. +Work: exact chunks; opt-in sm_120 post-conv and K4 conv measured. | Claim / track | State | Next command or step | |---|---|---| @@ -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 | Exact 3.072x; post-conv 1.859x; K4 arm 1 exact/+6.44%; 256-channel falsified | Repeat A/B; gate models unavailable | | 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:** post-conv/K4 exact and opt-in; repeat A/B or divide gap. 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`). @@ -81,8 +80,7 @@ SCOPE (`ROAD-V1-D6`). - Mirror vLLM; never ask how a feature should behave. - `nsys` BOTH sides, SAME tool, before any perf claim; cross-tool comparisons never establish invocation parity; whole-run sums mix prefill. -- GPU: park `local-ai-worker`, flock `$HOME/gpu.lock`, single-load - steady-state, never reload per rep, named tmux. +- GPU: flock `/tmp/gpu`; use the 22/25 GiB user-systemd scope; single-load reps. - Never weaken a checker to pass; repair the record. - Work happens in its own worktree on a task branch; the shared checkout stays clean on `main`, never a work surface. Land via `row/*` PR or authorized diff --git a/.agents/benchmark-record.md b/.agents/benchmark-record.md index 48b408ad9..9a232c46d 100644 --- a/.agents/benchmark-record.md +++ b/.agents/benchmark-record.md @@ -17619,3 +17619,117 @@ 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). diff --git a/.agents/completed/state-events/2026-08/STATE-20260808T231000-001.md b/.agents/completed/state-events/2026-08/STATE-20260808T231000-001.md new file mode 100644 index 000000000..02ed30f94 --- /dev/null +++ b/.agents/completed/state-events/2026-08/STATE-20260808T231000-001.md @@ -0,0 +1,22 @@ +# sm_120 post-conv token tile implemented and locally positive + + +After exact chunks was rebased and revalidated, the next same-tool target was +fused GDN post-conv. The current per-V-head split was first falsified: 448.365 +ms versus fast megablock 228.150 ms. Pinned vLLM's 16-token/per-head Triton +schedule is 108.035 ms, selecting token tiling plus Q/K register reuse rather +than another V-only split. + +The new opt-in `GdnPostConvTokenTileKernel` ports that four-warp schedule for +128-wide heads. A first faster reduction changed production tokens and was +rejected. Reproducing the existing 128-lane reduction tree within each warp +restored byte identity while preserving the speedup. Final graph-node A/B is +227.887→122.587 ms (**1.859x**), leaving **1.135x** to vLLM; total/output +throughput improves 0.532%, TTFT 0.850%, TPOT 0.457% and E2E 0.531%. + +Portable flag/grid 6/6·50, CUDA GDN 67/67·4384 and cached Qwen3.5-4B +3/3·1672 are green, including packed BA strides and exact five-output bytes. +The flag remains opt-in: repeated A/B and hardware-unavailable 27B/35B gates +are still required before a default flip. Resume from +`row/KERNEL-SSM-MAMBA-POSTCONV-TILE`; revalidate on current main, then profile +the next whole-run residual and divide it by kernel shape. diff --git a/.agents/completed/state-events/2026-08/STATE-20260808T232000-001.md b/.agents/completed/state-events/2026-08/STATE-20260808T232000-001.md new file mode 100644 index 000000000..2a55ab9eb --- /dev/null +++ b/.agents/completed/state-events/2026-08/STATE-20260808T232000-001.md @@ -0,0 +1,25 @@ +# sm_120 causal-conv K4 specialization locally wins; 256-channel tile falsified + + +Fresh current-head `nsys` selected causal conv after the post-conv residual fell +to 14.476 ms: local causal conv was 234.255 ms versus pinned-vLLM 145.532 ms. +Launch-shape division put 58.1% in the 279/280-program waves. A committed spike +split the structural differences into compile-time K=4 at the same 64-block +feature grid, then two channels/thread at a 32-block grid. + +Fresh implementer/reviewer passes found and repaired two correctness/test issues. +The first serial-stripe arm 2 could overwrite state before its second stripe +loaded initial history; the CUDA matrix caught it. Mutation review then deleted +whole specialized branches without failing byte comparisons; production now +uses the same callback dispatcher as the portable mutation gate. Scoped +re-review is clean. Portable 9/9·88, CUDA GDN 67/67·4631 and Qwen3.5 +paged-forward 4/4·8 pass independently. + +The first production profile series is VOID because `vllm-bench` had not been +relinked; all traces proved the old runtime kernel ran. After explicit relink, +all arm token hashes match. Arm 1 (`K4<1>`) improves causal conv +234.605→219.506 ms (**6.44%**) and every observed enclosing axis slightly; +same-tool vLLM residual is 1.508x. Arm 2 (`K4<2>`) halves blocks but raises +registers to 58 and is 4.05% slower than arm 1, with neutral/negative enclosing +movement, so the 256-channel hypothesis is falsified. Arm 1 stays opt-in: +repeat A/B and unavailable 27B/35B gates remain before any default flip. diff --git a/.agents/kernel-matrix.md b/.agents/kernel-matrix.md index 0a0cf92ec..ad2a37eda 100644 --- a/.agents/kernel-matrix.md +++ b/.agents/kernel-matrix.md @@ -172,6 +172,29 @@ 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-08 `KERNEL-SSM-MAMBA` checkpoint.** The opt-in +`VT_GDN_POSTCONV_TOKEN_TILE` path ports vLLM/FLA's 16-token, per-head, +four-warp post-conv schedule for 128-wide Q/K/V heads. It preserves the current +128-lane reduction tree exactly while retaining Q/K values in registers. +Qwen3.5-4B sm_120 graph-node time is **227.887→122.587 ms (1.859x)** with +byte-identical production tokens and every enclosing axis positive; the +same-tool vLLM residual is **1.135x**. Portable 6/6, CUDA GDN 67/67·4384 and +cached 4B 3/3·1672 pass. It remains opt-in pending repeated and 27B/35B gates; +generic Mamba lifecycle stays `INVENTORIED`. +[Spike and result](specs/sm120-qwen35-postconv-token-tile-2026-08-08.md). + +**2026-08-08 `KERNEL-SSM-MAMBA` causal-conv residual checkpoint.** The +opt-in width-four, one-channel arm removes runtime-width work at the unchanged +64-block feature grid and is byte-exact. Qwen3.5-4B sm_120 graph-node time is +**234.605→219.506 ms (-6.44%)**; dominant waves improve about 149.5→140.1 us +and every observed enclosing axis is slightly positive. CUPTI registers rise +43→52, so the win is not occupancy. The separately measured 256-channel arm +halves the grid but rises to 58 registers and is 4.05% slower than the K4 +one-channel arm; that hypothesis is falsified. Portable 9/9·88, CUDA GDN +67/67·4631 and paged-forward 4/4·8 pass. Arm 1 remains opt-in; generic Mamba +lifecycle stays `INVENTORIED`. +[Spike and result](specs/sm120-qwen35-conv-channel-tile-2026-08-08.md). + ## Count invariants - This table has exactly 35 practical kernel-family rows. diff --git a/.agents/parity-ledger.md b/.agents/parity-ledger.md index 1b751a939..30d9eb54d 100644 --- a/.agents/parity-ledger.md +++ b/.agents/parity-ledger.md @@ -934,3 +934,5 @@ Columns: | 2026-08-07 (`SERVE-CLI-BENCH` + `KERNEL-SSM-MAMBA`; clean sm_120 exact-chunk transplant; local-4B binding only) | Makes the benchmark use production `AsyncLLM`, then ports vLLM's exact `(sequence, BLOCK_M=8 token chunk)` descriptors into shared GDN step metadata and maps one CUDA register-kernel program to each descriptor. `VT_CONV_EXACT_CHUNKS` defaults ON with a same-binary `=0` rollback; `VT_CONV_REG=0` retains tiled/scalar. | Pinned vLLM engine-core queued dispatch `vllm/v1/engine/core.py:200-231,622-669` and causal-conv descriptor mapping `causal_conv1d.py:15-28,71-79,123-124`; local anchors `examples/bench/bench_core.h`, `gdn_attn.{h,cpp}`, `qwen3_5.cpp`, `ops.{h,cpp}`, `cuda_gdn.cu`; [spike/result](specs/sm120-qwen35-conv-chunking-2026-08-07.md). | **ACCEPTED + REBASED-MAIN REPRODUCED.** Contained rebuild; CPU 6/6, CUDA GDN 66/66·4300, cached 4B 3/3·1672; exact/rollback token files identical. On `3d2581551` over `upstream/main` `48a54141f`, same-binary `nsys` reproduces conv **718.704→233.955 ms = 3.072x** and profiled total **6589.65→6739.34 tok/s = +2.272%**; vLLM 145.421 ms leaves **1.609x**. Binding three-pair A/B remains total/output **+2.152%**, TTFT **-2.945%**, TPOT **-1.920%**, E2E **-2.118%**, VRAM unchanged. Sealed-vLLM throughput **1.021246x PASS**; latency/VRAM OPEN. No gate-model extrapolation. [Evidence](../docs/bench-evidence/qwen35-4b-sm120-main-20260807.md). | | 2026-08-08 (`row/ENG-RELEASE-BINARIES`; `ENG-RELEASE-BINARIES` W5; draft PR #141; build-time tooling only) | **Versioned deterministic binary-release manifest contract.** Adds schema v1 plus a stdlib-only generator/validator that records the exact host/backend/static boundary, dependencies, clean source/toolchain/CMake flags/test commands, adaptive CPU tiers and probes, ten-SM CUDA fat inventory with independent runtime/AOT evidence, and preview/stable publication proof. Synthetic CPU/CUDA fixtures are canonical contract inputs, never artifact evidence. | Mirrors the metadata/evidence separation of vLLM's pinned release lanes (`.buildkite/release-pipeline.yaml:1-18,34-170` @ `555967922`) while expressing vllm.cpp's KISS adaptive-CPU/fat-CUDA matrix and explicit external-runtime boundary. No vLLM runtime behavior changes and no release artifact parity is claimed. | RED-first mutation suite 13/13 covers missing/unknown/inconsistent fields, all evidence states, every current x86 tier, all ten SMs/six AOT trees, channel/static/dependency boundaries, duplicate JSON and nondeterminism; accepted release-contract suite remains 30/30. No GPU, archive, runtime, correctness or performance gate ran. | | 2026-08-08 (`row/ENG-RELEASE-BINARIES`; `ENG-RELEASE-BINARIES` W5 scoped re-review repair; draft PR #141; tests/records only) | **Closes three surviving W5 manifest-contract mutations without changing production behavior.** Direct assertions now keep JSON booleans distinct from both the integer schema type and integer constants, while an independent CUDA manifest mutation removes the named `nvidia-driver` dependency and requires the exact missing-driver diagnostic. The existing `schema_version: true` integration rejection and arbitrary-driver rejection remain. | No upstream runtime analogue or behavior delta: this strengthens the local release metadata gate around the already-implemented JSON-type identity and external NVIDIA driver boundary. | Each production removal is RED for its own reason: deleting the integer/bool exclusion fails the direct type assertion; deleting JSON type identity fails the direct constant assertion; deleting `require_external("nvidia-driver", "driver")` leaves the missing-driver manifest error-free and fails the exact diagnostic assertion. Restored source is 19/19; cumulative fresh-review production removals killed are 10/10; accepted release-contract suite remains 30/30. Fixtures remain byte-identical synthetic inputs. No artifact, archive, runtime, correctness, performance, GPU, download, or release-channel claim advances. | +| 2026-08-08 (`KERNEL-SSM-MAMBA`; sm_120 post-conv tile; local-4B only) | Adds opt-in `GdnPostConvTokenTileKernel`: vLLM/FLA's 16-token, per-head, four-warp schedule for 128-wide heads, retaining Q/K in registers while reproducing the shipped 128-lane reduction tree byte-exactly. `VT_GDN_POSTCONV_TOKEN_TILE=1`; unset/`0` keeps the fast megablock. | Pinned FLA `fused_gdn_prefill_post_conv.py:57-149,208-245` and upstream `tests/kernels/test_fused_gdn_post_conv.py:60-208`; local `cuda_gdn.cu`, `gdn_prefill_conv.h`, portable and CUDA tests; [spike/result](specs/sm120-qwen35-postconv-token-tile-2026-08-08.md). | **IMPLEMENTED OPT-IN.** The first faster reduction changed tokens and was rejected. Accepted same-binary graph trace is byte-exact and **227.887→122.587 ms = 1.859x**, leaving **1.135x** to same-tool vLLM 108.035 ms; total/output +0.532%, TTFT -0.850%, TPOT -0.457%, E2E -0.531%. Tests: portable 6/6·50, CUDA GDN 67/67·4384, cached 4B 3/3·1672. Default/repeated and 27B/35B gates OPEN; no extrapolation. | +| 2026-08-08 (`KERNEL-SSM-MAMBA`; sm_120 causal-conv K4/channel-tile residual; local-4B only) | Adds same-binary K=4 causal-conv arms: `VT_CONV_CHANNEL_TILE=1` specializes width at the current 128-channel tile; `=2` also gives each thread two channels/256-channel tile. Unset/`0` keeps the runtime-width kernel. A shared portable/production dispatcher is mutation-gated. | Pinned vLLM `_causal_conv1d_fwd_kernel` compile-time `KERNEL_WIDTH`, `BLOCK_M=8`, `BLOCK_N=256` in `causal_conv1d.py:16-63,78-79,692-742`; local `cuda_gdn.cu`, `gdn_prefill_conv.h`, portable/CUDA tests; [spike/result](specs/sm120-qwen35-conv-channel-tile-2026-08-08.md). | **ARM 1 IMPLEMENTED OPT-IN; ARM 2 FALSIFIED.** Byte-exact rebuilt same-binary trace: arm 1 **234.605→219.506 ms (-6.44%)**, enclosing total/output +0.122%, TTFT -0.282%, TPOT -0.086%; vLLM residual **1.508x**. Arm 2 is 228.401 ms, 4.05% slower than arm 1 and enclosing-neutral/negative. Registers 43→52→58 refute the occupancy thesis. Portable 9/9·88, CUDA GDN 67/67·4631, paged-forward 4/4·8; default/repeated and 27B/35B gates OPEN. | diff --git a/.agents/roadmap_v1.md b/.agents/roadmap_v1.md index d8ae67cca..457497727 100644 --- a/.agents/roadmap_v1.md +++ b/.agents/roadmap_v1.md @@ -100,7 +100,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-chunk conv is default ON: **718.704→233.955 ms (3.072x)** and binding total/output **+2.152%**; sealed-vLLM throughput **1.021246x PASS**, latency/VRAM OPEN. Opt-in post-conv is **1.859x** faster. Opt-in K4 conv arm 1 is byte-exact and **234.605→219.506 ms (6.44%)**; the 256-channel arm is falsified | `GATING` | Repeat opt-in A/B and run unavailable 27B/35B gates before default flips. Latency/VRAM remain open; no 4B extrapolation. Evidence: [exact chunks](../docs/bench-evidence/qwen35-4b-sm120-main-20260807.md), [post-conv](specs/sm120-qwen35-postconv-token-tile-2026-08-08.md), [K4 conv](specs/sm120-qwen35-conv-channel-tile-2026-08-08.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-conv-channel-tile-2026-08-08.md b/.agents/specs/sm120-qwen35-conv-channel-tile-2026-08-08.md new file mode 100644 index 000000000..13fddd9e6 --- /dev/null +++ b/.agents/specs/sm120-qwen35-conv-channel-tile-2026-08-08.md @@ -0,0 +1,173 @@ +# sm_120 Qwen3.5 causal-conv residual — 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 and byte-exact; arm 1 is a locally positive opt-in, arm 2 is +falsified, and default/release-model gates remain pending. + +## Measured selection + +A fresh graph-node `nsys` profile at `7f66792f8`, after the exact-chunk and +opt-in post-conv-tile changes, reports: + +| Family | ours | pinned vLLM | ours / vLLM | local excess | +|---|---:|---:|---:|---:| +| prefill causal conv | **234.255 ms** / 1,728 calls | **145.532 ms** / 1,897 calls | **1.6096x** | **88.723 ms** | +| fused post-conv | 122.511 ms / 1,728 calls | 108.035 ms / 1,923 calls | 1.1340x | 14.476 ms | + +The next metric is therefore **total GPU time of +`CausalConv1dFwdRegKernel` on the exact c32 production workload**, with the +dominant launch shape used as its first micro-metric. Shape-grouping the same +tool's kernel rows divides that total further: + +| Exact programs | Calls | Total | Mean | Share of local conv | +|---:|---:|---:|---:|---:| +| 279 | 576 | 86.052 ms | 149.395 us | 36.7% | +| 280 | 336 | 50.137 ms | 149.218 us | 21.4% | +| 156-158 | 216 | 17.647 ms | 81.699 us | 7.5% | +| all remaining shapes | 600 | 80.419 ms | mixed | 34.3% | + +The 279-280-program waves alone are 912 launches and **136.189 ms (58.1%)**, +so their mean launch time is the selected micro-metric. This avoids guessing +from a whole-run aggregate when a launch-shape-local change is available. + +Artifacts: + +- local report `/tmp/qwen35-next-7f66792f.nsys-rep`, SQLite export + `/tmp/qwen35-next-7f66792f.sqlite`; +- pinned-vLLM report `/tmp/qwen35-async-3f35356e0-vllm.nsys-rep`, SQLite + export `/tmp/qwen35-async-3f35356e0-vllm.sqlite`. + +Nsight Compute is not present in the selected Nix CUDA environment or the host, +so no hardware-counter claim is made. Installing a profiler package is outside +this campaign's permission envelope. Register counts and launch geometry below +come from the same `nsys` CUPTI rows, not an inferred occupancy claim. + +## Whole-chain difference and hypotheses + +Pinned vLLM launches Triton's `_causal_conv1d_fwd_kernel` with +`BLOCK_M=8`, `BLOCK_N=256`, four warps and two stages +(`${VLLM_SOURCE}/vllm/model_executor/layers/mamba/ops/causal_conv1d.py:16-63,78-79,692-742`). +On the dominant wave that resolves to `grid=(279,32,1)`, block 128, and 32 +registers/thread. Its 279-program launches average 88.346 us. + +The local exact-descriptor port already matches `BLOCK_M=8` and enumerates the +same sequence/chunk programs, but retains a runtime convolution width and one +channel per thread: `kConvRegN=128`, `grid=(64,279,1)`, block 128, and 43 +registers/thread (`src/vt/cuda/cuda_gdn.cu:702-854`). Its 279-program launches +average 149.395 us. Qwen3.5-4B has 8,192 convolution channels and width four, +so 64 versus 32 feature blocks is exactly the `BLOCK_N` difference. + +Two independent hypotheses remain and must be measured separately: + +1. **Width specialization.** vLLM's `KERNEL_WIDTH` is a compile-time constant; + local `k` is runtime and reserves arrays through `kConvRegMaxW+1=9` while + guarding every unrolled tap. Dispatching the production `k=4` instantiation + can remove dead taps/branches and reduce register or instruction cost without + changing the grid. +2. **Two channels per thread.** A 128-thread block can process two coalesced + 128-channel stripes, giving a 256-channel feature tile and 32 feature blocks + like upstream. It halves block-level descriptor/control work, but duplicates + each thread's channel-local weights and window. Register pressure can refute + this even when the grid looks better. + +These are ordered discriminators, not one combined patch: first compare the +width-four specialization against the runtime-width kernel at the unchanged +128-channel tile; only then add the two-channel instantiation and compare it +against the specialized one. Unsupported widths and dimensions retain the +current kernel. + +## Port and rollback + +Add compile-time width/channel-count instantiations of the existing register +window, preserving for each channel the exact operation sequence: bias; taps +`j=0..3`; current SiLU; store; window shift; raw-input final-state writeback. +The experiment must not change exact descriptor metadata, scheduler/model +routing, tensor strides, or post-conv dispatch. + +Expose an explicit same-binary experiment selector with three arms: + +- `0`: current runtime-width, one-channel kernel (sealed baseline); +- `1`: width-four specialization, one channel per thread; +- `2`: width-four specialization, two channels per thread. + +Unset retains arm 0 until the evidence supports a default change. Values or +shapes outside the supported experiment fall back to arm 0. Document the +selector in `docs/ENVIRONMENT.md`; factor its parse/launch-contract predicates +into the portable GDN prefill header so CPU-tier tests can kill accidental +default or grid changes. + +## Tests and acceptance + +RED-first coverage must prove: + +1. the selector defaults to arm 0, accepts only the named arms, and invalid + values roll back; +2. the production `C=8192,K=4` launch contract is respectively 64/64/32 feature + blocks with block 128, while partial tiles round up safely; +3. CUDA arms 1 and 2 are byte-identical to arm 0 for output and final state over + BF16/F32 I/O, 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-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..857df7011 100644 --- a/docs/BENCHMARKS.md +++ b/docs/BENCHMARKS.md @@ -43,11 +43,12 @@ The binding comparison. vLLM runs its **production graphed config**, never | 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)) | -### GDN prefill causal-convolution by GPU +### GDN prefill kernels by GPU | GPU | Workload and basis | vllm.cpp | vLLM | Ratio | Status | |---|---|---:|---:|---:|---| -| RTX 5070 Ti (`sm_120`) | Qwen3.5-4B BF16, c32, steady-interval total | 233.955 ms | 145.421 ms | **1.609x slower** | Rebased-main exact chunks ON; rollback 718.704 ms, so local is **3.072x faster** ([evidence](bench-evidence/qwen35-4b-sm120-main-20260807.md)) | +| RTX 5070 Ti (`sm_120`) | Qwen3.5-4B BF16, c32, steady-interval total | 233.955 ms default; **219.506 ms** K4 arm 1 | 145.421 ms | 1.609x default; **1.509x** arm 1 | K4 arm 1 is byte-exact and **6.44% faster**; arm 2 falsified. Arm 1 stays opt-in ([result](../.agents/specs/sm120-qwen35-conv-channel-tile-2026-08-08.md)) | +| RTX 5070 Ti (`sm_120`) | Qwen3.5-4B BF16, c32, fused post-conv total | 122.587 ms | 108.035 ms | **1.135x slower** | Opt-in tile is exact and **1.859x faster** than 227.887 ms fast megablock; all axes improved; default/release gates open ([spec](../.agents/specs/sm120-qwen35-postconv-token-tile-2026-08-08.md)) | | GB10 (`sm_121a`) | Qwen3.6-27B NVFP4, historical normalized prefill | 0.43 us/token/layer | 0.18 us/token/layer | **2.39x slower** | Directional only: unequal token clusters, older pin ([ledger](../.agents/parity-ledger.md)) | | GB10 (`sm_121a`) | Qwen3.6-35B NVFP4, later local kernel A/B | 321.148 us c1; 960.313 us c6 | - | `PENDING` | Register vs tiled improved 4.7%/7.3%; no paired vLLM denominator ([record](../.agents/specs/gdn-prefill-conv-reg-2026-07-18.md)) | | Jetson Thor (`sm_110`), AGX Orin (`sm_87`) | No matched GDN workload | - | - | `PENDING` | Runtime correctness only; no causal-conv speed trace | diff --git a/docs/ENVIRONMENT.md b/docs/ENVIRONMENT.md index 0c58b01bf..97078b027 100644 --- a/docs/ENVIRONMENT.md +++ b/docs/ENVIRONMENT.md @@ -89,6 +89,7 @@ portable/reference path. In normal operation leave them unset. | `VT_GDN_PACKED_DECODE` | on (CUDA GDN) | Unpacked GDN decode path | | `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) | Selects the same-binary register-window causal-conv experiment: `0` is the sealed runtime-width one-channel baseline, `1` specializes the production width `K=4` at the same 128-channel tile, and `2` specializes `K=4` while assigning two channels per thread (256-channel tile). Only the exact values `1` and `2` opt in; unset, invalid values, and non-`K=4` shapes fall back to `0`. Experimental only; no default flip is implied by a microbenchmark win | | `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 71a06b08e..949ec902e 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 and kernel default gates open | Exact chunks ON: local +2.152% total. Opt-in post-conv tile is 1.859x faster. Opt-in causal-conv K4 arm 1 is byte-exact and 234.605→219.506 ms (6.44%); the 256-channel arm is falsified | | 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) | @@ -1291,23 +1291,20 @@ regression. ## Performance detail **Local Qwen3.5-4B plain BF16 direct loader, throughput passed; latency and -VRAM pending:** the benchmark now uses production `AsyncLLM`. Exact -`(sequence, 8-token chunk)` causal-conv metadata is built once per step, shared -across GDN layers and default ON. `VT_CONV_EXACT_CHUNKS=0` is a byte-identical -same-binary rollback. Rebased-main graph-node `nsys` confirms the mechanism: -causal-conv falls from **718.704 to 233.955 ms (3.072x)**, leaving **1.609x** -to vLLM. The profiled whole run improves **2.272%** with identical token files. - -The enclosing A/B improves total/output **2.152%**, TTFT **2.945%**, TPOT/ITL -**1.920%** and E2E latency **2.118%** without a local VRAM regression. Against -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. - -This local 4B diagnostic does not establish 27B/35B support. Exact evidence and -reproduction: -[Qwen3.5-4B exact-chunk outcome](bench-evidence/qwen35-4b-sm120-main-20260807.md). +VRAM pending:** production `AsyncLLM` uses default-ON exact `(sequence, +8-token chunk)` causal-conv dispatch. Rebased-main graph-node `nsys` measures +**718.704→233.955 ms (3.072x)** and +2.272% profiled throughput, leaving 1.609x +to vLLM; the binding A/B improves total/output 2.152%, TTFT 2.945%, TPOT 1.920% +and E2E 2.118%. Against sealed vLLM, throughput is **1.021246x PASS**; TTFT +1.085812x, TPOT 1.024597x and peak VRAM 13053.3/12820 MiB remain OPEN. + +The next same-tool target, fused post-conv, now has an opt-in upstream-shaped +16-token/four-warp CUDA tile. It is byte-exact and cuts **227.887→122.587 ms +(1.859x)**, moving the vLLM gap from 2.112x to **1.135x**; the one enclosing +profile improves total/output 0.532%, TTFT 0.850%, TPOT 0.457% and E2E 0.531%. +It remains opt-in pending repeated A/B and unavailable 27B/35B gates. Evidence: +[exact chunks](bench-evidence/qwen35-4b-sm120-main-20260807.md), +[post-conv tile](../.agents/specs/sm120-qwen35-postconv-token-tile-2026-08-08.md). There is no front-page race clip yet; when one is produced it will follow the LocalAI house style (side-by-side, identical output, honest measured ratios). diff --git a/scripts/env-doc-allowlist.txt b/scripts/env-doc-allowlist.txt index 5abe69189..d343e20bf 100644 --- a/scripts/env-doc-allowlist.txt +++ b/scripts/env-doc-allowlist.txt @@ -92,6 +92,7 @@ VT_GDN_PACKED_DECODE_TRITON VT_GDN_PACKED_REG_TILE VT_GDN_POSTCONV_FAST VT_GDN_POSTCONV_SPLIT +VT_GDN_POSTCONV_TOKEN_TILE VT_GDN_SLACK_MEMSET VT_GDN_STATE_BF16 VT_GDN_TILE_PIPE diff --git a/src/vt/cuda/cuda_gdn.cu b/src/vt/cuda/cuda_gdn.cu index ae4d83f67..f05709c75 100644 --- a/src/vt/cuda/cuda_gdn.cu +++ b/src/vt/cuda/cuda_gdn.cu @@ -842,6 +842,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 +962,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 +1019,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 +1065,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 +1565,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 +1732,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 +1755,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 +1777,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(), 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/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..efb0a26f7 100644 --- a/tests/vt/test_ops_gdn.cpp +++ b/tests/vt/test_ops_gdn.cpp @@ -707,6 +707,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()); @@ -2392,14 +2411,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 +2441,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 +2468,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 +2484,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 +2493,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 +2540,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 +2576,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 +2600,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 +2631,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");