diff --git a/.agents/roadmap_v1.md b/.agents/roadmap_v1.md
index e9e3a4d37..f6a8655c5 100644
--- a/.agents/roadmap_v1.md
+++ b/.agents/roadmap_v1.md
@@ -138,6 +138,10 @@ issue is not yet placed. Keyed record: update in place, never append.
 | [#277](https://github.com/mudler/vllm.cpp/issues/277) | `SERVE-METRICS` | `/metrics` is dead on the production server: AsyncLLM never folds IterationStats into the logger | bug |
 | [#321](https://github.com/mudler/vllm.cpp/issues/321) | `SERVE-METRICS` | AsyncLLM output thread outlives metrics logger and uses freed PromRegistry | bug |
 | [#81](https://github.com/mudler/vllm.cpp/issues/81) | `SPEC-MTP` | Speculative decoding: MTP k>1, dynamic and adaptive depth | feature |
+| [#436](https://github.com/mudler/vllm.cpp/issues/436) | `SPEC-DSPARK` | W7: the sequential Markov sample loop is HOST-side — 28% of the draft step, 14.9 MB downloaded per step | perf |
+| [#442](https://github.com/mudler/vllm.cpp/issues/442) | `SPEC-DSPARK` | W8: capture the T=1+k speculative VERIFY shape (the measured 0.870x-0.981x gap) | perf |
+| [#513](https://github.com/mudler/vllm.cpp/issues/513) | `SPEC-DSPARK` | DFlash/DSpark draft loading hardcodes the VL weight prefix, so a text-only Qwen3.5 target cannot share its embedding | bug |
+| [#536](https://github.com/mudler/vllm.cpp/issues/536) | `ROAD-V1-C3` | C3's named tail records DSpark as unspiked three days after it landed, and files TLI under the wrong lane | bug |
 | [#83](https://github.com/mudler/vllm.cpp/issues/83) | — | Auto-size memory to the workload; pre-flight error instead of OOM | feature |
 | [#85](https://github.com/mudler/vllm.cpp/issues/85) | — | In-source `cmake .` collides every examples/ target with its own source directory and dies in the linker | bug |
 | [#146](https://github.com/mudler/vllm.cpp/issues/146) | — | Pascal support | feature |
@@ -196,7 +200,7 @@ models we already ship + benchmark. Full seam map + M0–M5 W-plan:
 | 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/<gpu>.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(<obj>.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%**. The reviewed combined default-off K4 causal-conv + 16-token post-conv A/B is exact and improves total throughput **0.7175%**, TTFT **1.340%**, TPOT **0.556%** and E2E **0.716%**. Atomic all-or-zero AsyncLLM wave admission now matches vLLM's add-all-before-step behavior and fixes the pretokenized exactness confound. Corrected three-repetition cross-engine result: throughput **6831.71 vs 6643.40 tok/s (1.0283x PASS)**; host PSS passes; TTFT **1016.51 vs 936.59 ms (1.0853x slower)**, TPOT **34.477 vs 33.917 ms (1.0165x slower)**, E2E **5395.13 vs 5244.07 ms (1.0288x slower)** and VRAM **+118.7 MiB** remain open. Exact BV16+swizzle+REGSTATE and slack-memset local opt-ins improve all enclosing axes. Geometric argmax scratch (wait migration, +1% TPOT) and BF16 vector writeback (only 0.143% y800; unstable) are rejected and removed | `GATING` | Preserve the throughput stack; use same-tool interval profiling to split remaining TTFT intake/prefill and then target the largest attributable interval. VRAM and 27B/35B correctness remain gates; no 4B-to-gate-model extrapolation. Evidence: [campaign spec](specs/sm120-qwen35-pareto-2026-08-09.md), [production baseline](../docs/bench-evidence/qwen35-4b-sm120-main-20260807.md) |
-| 3 | `ROAD-V1-C3` | MTP k=1 + GDN speculative path, then DFlash, DSpark and heterogeneous-vocabulary TLI | [engine matrix](engine-matrix.md), [coverage view §8](feature-matrix.md#8-speculative-decoding) | MTP and DFlash specs exist. **M-mtp-0 CLOSED 2026-07-24** - the standalone MTP draft head is oracle-parity-proven on BOTH gate checkpoints (op-level parity vs a dumped k=1 vLLM oracle, not a token-generation SACRED gate). **I2 SCHEDULER-HALF LANDED 2026-07-24** ([mtp-spec-decode §2.7](specs/mtp-spec-decode.md)): host-side spec-decode scheduler/engine plumbing + the FROZEN spec-metadata ABI that I3 (rejection sampler) and I5 (verify/propose runner) build against - `SpeculativeConfig`, `DraftTokenIds`, `Request::spec_token_ids`/`NumTokensWithSpec`, the first population of `scheduled_spec_decode_tokens`, `Scheduler::update_draft_token_ids`, the `take_draft_token_ids` seam, `EngineCore::post_step`, `InputBatch::num_accepted_tokens`/`update_req_spec_token_ids`; DEFAULT-OFF and INERT (no `SpeculativeConfig` => `num_lookahead_tokens == 0` => byte-identical engine). `SPEC-MTP` **STAYS `GATING`** because M-mtp-1..4 (greedy rejection, GDN spec slots, k>1, CUDA graphs) are still open, so spec decode remains user-invisible. DSpark is user-promoted scope with DeepSeek-V4/Qwen3 draft models, reduced-vocabulary handling and full-CUDA-graph behavior inventoried under `SPEC-DSPARK`; tokenizer-agnostic target<->draft mapping is separately inventoried as `SPEC-TLI`. Their dedicated spikes are not written **I3 GREEDY REJECTION SAMPLER LANDED 2026-07-24** (`SPEC-REJECTION` `READY` -> `ACTIVE`): per-request logits expansion to `1 + k_i` rows plus the greedy accept rule (accept a draft iff it equals the target argmax; on the first mismatch emit the target argmax and stop; bonus token when all k accept), CUDA==CPU bit-exact at vocab 248320. **I4 GDN-HALF LANDED 2026-07-24** (`SPEC-GDN-SEGMENTS` -> `ACTIVE`): the GDN spec metadata split + decode->prefill reclassification, the `T>1`/`IS_SPEC` recurrence with per-timestep snapshots, the conv sliding window honouring `num_accepted`, k+1 slot allocation - bit-exact rollback. **I5a GDN LAYER ROUTING + runner spec-metadata upload LANDED 2026-07-24** (`CLAIM-SPEC-MTP-I5A`): `GdnBlockPaged` now routes a pure-spec batch through `vt::GdnSpecDecode`/`vt::CausalConv1dSpecUpdate` and the runner uploads I4's six spec device tensors - first of the scoped M-mtp-1 sub-increments (I5a GDN wiring -> I5b prepare_prefill -> I5c MTP paged propose -> I5d config+runner-loop+27B token gate, spec §5), DEFAULT-OFF INERT, bit-exact vs the I4 ops, no e2e loop yet. **I5b `prepare_prefill_inputs` LANDED 2026-07-24** (`CLAIM-SPEC-MTP-I5B`): the drafter prefill input-prep - shift each request's `input_ids` left one within its query span, splice the just-sampled next token, `query_len -= num_rejected`, emit last-token index / query_start_loc / seq_lens into the `SpecPrefillInputs` struct; a HOST routine (no new CUDA kernel; mirrors our DEVICE-NEUTRAL `prepare_inputs`/`combine_sampled_and_draft_tokens` family), unit-gated 7 cases / 27 assertions RED-first, DEFAULT-OFF INERT, additive by construction. **I5c MTP PAGED PROPOSE + DRAFT KV LAYER LANDED 2026-07-24** (`CLAIM-SPEC-MTP-I5C`): `Qwen3_5MTPModel::ForwardPaged` runs the head + one full_attention decoder layer over the head's OWN paged draft KV layer (ReshapeAndCache + PagedAttention over the target's block table / slot mapping); `MakeQwen3_5KVCacheSpec(num_spec>0)` adds that draft KV layer (`fa_draft` FullAttentionSpec group, index num_hidden_layers); `ForwardDeviceTap` exposes the `[T,H]` post-final-norm hidden tap (INERT); and `MtpProposePrefill` is the callable k=1 propose (I5b shift-splice -> one paged forward -> argmax at last_token_indices, early-exit). CORE PROOF: the paged forward reproduces I1's standalone head logits/argmax on BOTH gate checkpoints; a two-step drive proves the draft-KV write/read (RED control diverges). DEFAULT-OFF INERT (no spec config -> draft KV layer not allocated, tap nullptr, target forward byte-identical); NOT wired into the runner step loop. **I5d-pre REGISTRY/FORWARD-SEAM ENABLING REFACTOR LANDED 2026-07-25** (`CLAIM-SPEC-MTP-I5D-PRE`): a scoping pass found the model seam is fully TYPE-ERASED, so the runner cannot reach the concrete target weights / hidden-state tap / loaded MTP weights the I5d loop needs. Four ADDITIVE, inert-when-spec-off access paths + one latent-bug fix - the `hidden_tap` out-field on the type-erased `ModelForwardInput` (routes to the existing `ForwardDeviceTap`), a `LoadedModel::BuildMtpDraft` virtual (typed path to the draft, null for non-MTP), MTP weight loading + shard retention in `FromModelDir` behind `EngineParams::speculative_config`, and the `GPUModelRunner` ctor widened with optional draft/draft-KV/`SpeculativeConfig`; PLUS the latent `initialize_kv_cache` fix (select the FIRST non-eagle full-attn group as the target so a third `fa_draft` group can't displace it; byte-identical at num_spec==0). DEFAULT-OFF INERT, unit-gated RED-first, spec-off SACRED gates byte-identical. **I5d CONFIG RUNTIME + VERIFY/PROPOSE RUNNER LOOP LANDED as a spec-off-byte-identical PARTIAL 2026-07-25** (`CLAIM-SPEC-MTP-I5D`): `--speculative-config` JSON parse -> `EngineParams` -> `LoadedEngine` resolution (widened KV `MakeQwen3_5KVCacheSpec(num_spec>0)`, `BuildMtpDraft`, forced sync scheduling, `MakeScheduler(spec)`, `EngineCore(check_for_draft=true)`) + the full runner loop (draft splice, hidden-tap capture, GDN builder spec-overload feed, k+1 GDN state-slot remap + widened conv cache + draft-KV alloc, `MtpProposePrefill`, `take_draft_token_ids`, acceptance telemetry). CUDA `-Werror` 0 warnings, cutlass-ON; spec-OFF SACRED byte-identical (27B 235/235, 35B 315/315, Coder 138/138 + spec unit tests ALL PASS). **The three-way 27B token gate is NOT yet passing**: the spec-ON engine RUNS the loop end to end and MEASURES the blocker (`test_qwen27_spec_decode`) - it throws on the FIRST prefill step at `gdn_state_gather: working/cache row shapes must match` (`src/vt/ops.cpp:1773`) because I4's spec conv rollback needs the conv row widened to `(K-1)+num_spec` while the non-spec GDN conv ops assume `(K-1)`. `SPEC-MTP` STAYS `GATING`. **I5e LANDED 2026-07-25 (`CLAIM-SPEC-MTP-I5E`): the non-spec GDN conv ops made widened-cache-aware (mirror vLLM `state_len=KERNEL_WIDTH-1` + physical `stride_conv_state_tok`, leading `(K-1)` sub-window, byte-identical at `num_spec==0`) AND the async input-combine forced off under spec (it overwrote the verify batch's draft position with the committed token -> 0 acceptance, RCA'd on the real 27B). **THE THREE-WAY 27B SINGLE-REQUEST GREEDY GATE PASSES**: our-spec-ON == vLLM `--speculative-config mtp` greedy == our-spec-OFF token-for-token, **acceptance 16/16 drafts accepted** (~16 target steps saved); spec-OFF SACRED byte-identical (27B 235/235, 35B 315/315, Coder 138/138), compute-sanitizer 0 on the spec step. `SPEC-MTP` LEAVES `GATING` (single-request greedy correctness PROVEN); NOT `DONE` - the MIXED `GdnBlockPaged` concurrency split/merge + the throughput A/B vs vLLM same-config are I6. **I6 LANDED 2026-07-25 (`CLAIM-SPEC-MTP-I6`, `benchmark_binding=true`): the §5 c1 THROUGHPUT GATE — ours spec-ON AT/ABOVE vLLM spec-ON on EVERY measured axis at c1** (TPOT 66.2/62.95 vs 69.1/65.3 ms prose/code, ours ~1.04x faster; output tput +4.6%/+3.9%; ITL/TTFT lower; acceptance ours 0.85/0.92 vs vLLM 0.838, within noise; spec helps both ~1.5-1.6x TPOT; ours ~4% faster spec-OFF too), via an additive example-only `--speculative-config` bench flag (NO engine code touched). STAYS `ACTIVE`: the c>1 mixed-batch path is still refused + owes a c>1 A/B, and no server-facing spec flag yet. | `ACTIVE` | M-mtp-0, I2 scheduler-half, I3 rejection sampler, I4 GDN spec slots, I5a GDN layer routing, I5b prepare_prefill, I5c MTP paged propose + draft KV, I5d-pre the registry/forward enabling seam, and I5d config runtime + verify/propose runner loop (spec-off byte-identical) are landed; next (before `SPEC-MTP` leaves `GATING`) is closing the measured I5d gate blocker - make the non-spec GDN conv ops widened-cache-aware (mirror vLLM `causal_conv1d` `state_len=width-1+(seqlen-1)`) + the MIXED `GdnBlockPaged` split/merge - then the passing M-mtp-1 27B k=1 greedy three-way token gate + acceptance, then M-mtp-2 35B, then DFlash, the DSpark spike/gates and TLI. **DFlash D0-redo + D1 LANDED 2026-07-26 (`CLAIM-DFLASH-D0D1`, [dflash-spec-decode §0](specs/dflash-spec-decode.md)): `SPEC-DFLASH` UNBLOCKED + `ACTIVE` on the advanced pin `555967922`/vLLM 0.26.0.dev0.** The prior 0.25.0 ORACLE-BLOCKED verdict is SUPERSEDED — under `VLLM_USE_V2_MODEL_RUNNER=1` (vllm#40898 resolved) the mixed-SWA/full z-lab 27B draft CONSTRUCTS and the drafter is ALIVE (acceptance 2.21/8.80/4.75/4.57 > 1, `num_spec=16`, flashinfer-native fp8-KV; goldens committed). Gate FORM measured STRICT MODE-MATCHED (vLLM-ON run-deterministic K>=3 but != vLLM-OFF at k=16 near-ties — NOT the MTP three-way identity). **D1 `DF-AUX-TAPS` DONE:** the single hidden tap is generalized to the multi-tap `[T,H×taps]` (`ForwardDeviceMultiTap` capturing `(hidden+res)` at `target_layer_ids`), config-gated byte-identical off; unit gate 598 assertions (RED-first), CUDA 697/697 + sanitizer 0, 27B MTP e2e 9/9 + 27B SACRED 235/235 byte-identical (inertness). **D2-D5 LANDED 2026-07-26 (`CLAIM-DFLASH-D2`/`D3`/`D4D5`/`D5`):** the drafter model + the project's first non-causal in-block attention (D2, GPU parity vs the real vLLM draft), context-KV precompute + `prepare_dflash_inputs` (D3, GPU numeric-parity 61/61), the non-autoregressive whole-block propose brick + `dflash` config-select (D4), and the RUNNER-LOOP INTEGRATION + 27B e2e (D5): the full verify/propose loop is wired (separate z-lab draft load + target-shared bf16 embed/lm_head, aux-tap capture, per-request combined-feature context accumulation honoring num_rejected, `propose_drafts_dflash`) and RUNS end to end - `test_qwen27_dflash_spec_decode` 2/4 STRICT token-exact vs the vLLM-DFlash-ON golden + acceptance ~ vLLM on ALL 4 (19/39/29/25 vs 17/39/30/25). The 2 divergences are SINGLE bf16 near-tie flips (ratified near-tie ROOT = the D3 inline context-KV recompute envelope), NOT a wiring bug; inertness SACRED 235/235 + MTP 9/9 byte-identical; CUDA `-Werror` clean, no new kernel. NOT a clean strict-4/4 pass - STRICT 4/4 token-identity + the speed A/B = D6 (persistent paged draft-KV bit-matching vLLM's fused projections + the uniform-1+k FULL CG). Capture tool + goldens: `scripts/spec/d{0,2,3}_dflash_*.py`, `tests/parity/goldens/dflash_27b{,_draft,_kvprep}/`. **D6-D9 SPEED CAMPAIGN 2026-07-27:** D6 c1 A/B + STRICT-4/4 bf16-irreducibility RCA; D7 device-resident within-step forward (bit-identical); D8 acceptance RCA + FINAL golden A/B (ours 0.69× vLLM). **D9 (`CLAIM-DFLASH-D9`) PERSISTENT PAGED DRAFT-KV LANDED (bit-identical, +22.7%): `AppendContextKVHost` + `ForwardBlockLogitsWithPrecomputedKV` replace the O(context²) per-step recompute with an append-only per-request store; ours-ON 20.99→25.75 tok/s = 0.917× vLLM-ON (28.09, was 0.69×); e2e 27/27 SAME tokens, SACRED 235/235 + MTP 9/9 byte-identical, CUDA `-Werror` clean, no new kernel. D8's "bf16 acceptance ceiling" REFUTED — same-trajectory per-step acceptance == vLLM (ratio 1.00) and ours realized acceptance (3.68/step) > vLLM (3.31); the SOLE residual (~8%) is the FULL uniform-(1+k) CUDA graph (eager-vs-graphed), a closeable increment. SPEC-DFLASH stays `ACTIVE` (speed not yet ≥ vLLM).** **SPEC-MTP → `DONE` 2026-07-26 (`CLAIM-SPEC-MTP-DONE`, records-only, closing commit I7 `72f9fb1`):** the user ratified the c>1 near-tie+SPEED criterion, closing both I6-owed items (mixed-batch concurrency + server/CLI/C-ABI `--speculative-config`); MTP k=1 is COMPLETE + gated. **M-mtp-2 CLOSED 2026-07-26 (`CLAIM-SPEC-MTP-M-MTP-2`): the 35B `Qwen3_5MoeMTP` full e2e three-way token gate PASSES** — our spec-ON == our spec-OFF == vLLM 0.25.0 `--speculative-config mtp` greedy == vLLM spec-OFF, 16/16 vs the `greedy_ids` anchor (STRICT, c1), acceptance 16/16 both sides; c1 spec-ON 1.19x TPOT / +16.3% output-tput vs spec-OFF (acceptance 0.908) — MoE speedup transfers; spec-OFF byte-identical (test+docs-only). MTP is now `DONE` on BOTH gate models (`MODEL-SPEC-qwen3-5-mtp-qwen3-5-moe-mtp` `GATING`→`DONE`). **DFlash D11+D12 2026-07-27 — the FULL uniform-(1+k) CUDA graph is being built in three parts:** D11 (`CLAIM-DFLASH-D11`) landed Part A (the device-store primitive, CPU-gated); **D12 (`CLAIM-DFLASH-D12`) landed A-wire (the D11 device store is now the PRODUCTION path; GPU-gated e2e 27/27 all-exact acceptance 19/39/29/25 + SACRED 235/235 + MTP 9/9 byte-identical) + Part B (`vt::DFlashPagedBlockAttention`, the capture-safe paged kernel; `test_ops_dflash_paged_block_attn` 795648/795648 CPU==CUDA + cross-check vs materialized `DFlashBlockAttention` + compute-sanitizer 0; NO function-local host `cu_seqlens` upload = capture-UAF fixed).** Speed UNCHANGED 0.917× (A-wire eager + Part B not yet wired). The SOLE remaining piece is Part C (static-shape capture + device mask-scatter + `BeginCapture`/replay + the ≥vLLM c1 A/B); if ours-ON-graphed ≥ vLLM-ON → SPEC-DFLASH DONE → C3 complete. C3 stays `ACTIVE` (DFlash Part C + DSpark/TLI remain) |
+| 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)~~ **SUPERSEDED 2026-08-12 ([#536](https://github.com/mudler/vllm.cpp/issues/536)): DFlash Part C LANDED and `SPEC-DFLASH` is `DONE` (D13/D14, `489a7544`); "their dedicated spikes are not written", earlier in this cell, is likewise superseded — the DSpark spike landed 2026-08-09 (`2b342620e`) and `SPEC-DSPARK` went `ACTIVE`.** **DSPARK W1–W8 LANDED + GPU-GATED (`CLAIM-SPEC-DSPARK`, [spec](specs/dspark-spec-decode.md) §§6a–6aa):** Markov head + sequential sampler + native AND Speculators-format loading + `d2t` reduced draft vocab + runner/one-surface wiring; W7 ([#436](https://github.com/mudler/vllm.cpp/issues/436)) moved the sequential Markov sample on device (text byte-identical, sampling −11%/−15%); W8 ([#442](https://github.com/mudler/vllm.cpp/issues/442)) captured the T=1+k VERIFY by mirroring vLLM's uniform-decode predicate (`uniform_decode_query_len = 1 + num_speculative_tokens`, `cudagraph_dispatcher.py:37`) — same-binary A/B +12.2%/+3.5%, text byte-identical, all four e2e spec suites green with capture ON and OFF. It also fixed an ENGINE-WIDE defect no landed gate could see: `EngineCoreProc` never threaded `check_for_draft_tokens`, so every speculator's drafts were dropped on the CLI and server paths. **Cross-engine, pinned clocks, interleaved reps: 35B-A3B MoE 0.975x code cell (non-overlapping distributions) / 1.012x prose cell — NOT parity, and not claimed as such.** The residual is localised to `marlin_moe_wna16::Marlin` and attributed to a **12.9% effective-DRAM-bandwidth** gap on byte-equivalent machine code (94 regs / 3664 SASS both sides, upstream doing 4.4% MORE work per launch); every source-level explanation is eliminated in §§6v–6z, and upstream `ncu` is BLOCKED in both replay modes. C3 stays `ACTIVE`: DSpark owes the SACRED-corpus token gate under the ratified near-tie protocol, the 27B dense re-measure, Gemma4's `1 + N` layout, and padded/multi-request capture shapes; `SPEC-TLI` remains untouched and its host row is `SPEC-DRAFT-MODEL`, not this lane (see the C3 area row) |
 | 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 |
 | 6 | `ROAD-V1-C6` | Priority and async/overlap scheduling + AsyncLLM streaming serving — RE-PRIORITIZED 2026-07-10: `SERVE-ASYNC-LLM` re-promoted T1→T0 because it blocks order 0 and async scheduling is vLLM's default at the pin (B3) | [engine matrix](engine-matrix.md), [coverage view §1](feature-matrix.md#1-engine-core--scheduling), [async-serving spike](specs/async-serving.md) | W1 `ENG-CORE-BUSY-LOOP`, W2 `SERVE-ASYNC-LLM`, W4 `ENG-PRIORITY-SCHED` implemented/`GATING`; **W3 `ENG-ASYNC-SCHED` host-side machinery + runner device-input half + sampler-OUTPUT half all LANDED + CPU-gated 2026-07-16 (`ACTIVE`, `CLAIM-ASYNC-SCHED-W3`):** `AsyncScheduler` placeholders + depth-2 `step_with_batch_queue` + `async_scheduling` config resolution + `VT_ASYNC_SCHED` rollback; the runner `combine_sampled_and_draft_tokens` device-input path; PLUS the sampler-OUTPUT half — `vt::Backend` event/pinned primitives, `AsyncGPUModelRunnerOutput` non-blocking sampled-ID D2H on a copy queue + event (`get_output()` waits only that event, MAIN queue never blocked), `Sampler::forward(sampled_ids_out)` device-resident greedy, `sample_tokens_async`/`runner_supports_async`, and the `Executor`+`step_with_batch_queue` async-output seam. All `VT_ASYNC_RUNNER` default OFF, sync path byte-identical, RED→GREEN; CPU ctest 111/111, tools 164/164. **The ENABLE-FLIP now LANDED + CPU-gated (2026-07-16, same claim):** `LoadedEngine` reorders `runner_` before the scheduler and builds `AsyncScheduler`+mcb=2 (threaded to `AsyncLLM`→`EngineCoreProc`→`step_with_batch_queue`) when `ResolveAsyncScheduling(runner_.runner_supports_async())` ON, else the byte-identical `Scheduler`+depth-1 (vLLM enable/disable log mirrored); the device combine/scatter kernel (`cuda_combine_tokens.cu`, 1:1 with `_combine_sampled_and_draft_tokens_kernel`) runs main-stream-ordered on the CUDA async path, deleting the pre-scatter `Synchronize`; construction-matrix CPU-tested RED→GREEN. Together they target the ~3.25 ms/step idle. **W3 async scheduling is now the DEFAULT (flipped ON `a0013a2`, 2026-07-17, `ENG-ASYNC-SCHED` DONE):** the TTFT-premium discriminator (`c63a1ec`, evidence `dgx:~/work/vllm.cpp-w3-discriminator/6ea7856…`) proved the +26–31% TTFT premium is **vLLM's OWN async trade** — vLLM async-ON vs its own sync pays the same envelope at −0.7% throughput, and upstream defaults it ON regardless — so the old "throughput-neutral ⇒ keep OFF" framing is retired as mis-calibrated. W3-on **flips both binding ITL-tail anomalies** (c16 p99_itl 1.024, c32 p90_itl 1.020 + p99_itl 1.026, now PASS) and nets positive on the axis arithmetic; the flip is token-neutral (both models, all arms bit-identical). `VT_ASYNC_RUNNER=0` / `VT_ASYNC_SCHED=0` are the same-binary rollbacks | `DONE` | `ENG-ASYNC-SCHED` closed; residual decode-mean parity (the ~1% deficit) now rides the RMSNorm/kernel-efficiency levers, not async |
@@ -263,7 +267,7 @@ live in [feature-matrix.md](feature-matrix.md).
 |---|---|---|
 | C1 | **Kernel drop-in alignment**: reshape `vt::` CUDA/ROCm adapter entry points around upstream `csrc` raw-pointer/shape/stride/stream signatures so copied kernels bind with only the Torch tensor wrapper replaced ([backends.md §drop-in](backends.md#post-mvp-drop-in-kernel-compatibility-with-upstream)) | 🚧 [implementation spike accepted](specs/dropin-kernel-abi.md); `BACKEND-ABI-VT` W0 is CPU-green/`GATING` with CUDA and consolidation debts named, then each family migrates with an independent checkpoint |
 | C2 | Dense/MoE model families: Llama, Qwen3 dense, Mixtral, then Qwen3-Next | ☐ T1; [feature matrix §4](feature-matrix.md#4-model-families) |
-| C3 | MTP speculative decode, starting k=1 on 27B and including the GDN path, then DFlash, DSpark and TLI | 🚧 **MTP k=1 e2e works on the 27B GDN hybrid: the three-way single-request greedy gate PASSES (2026-07-25, I5e) — our-spec-ON == vLLM `--speculative-config mtp` greedy == our-spec-OFF token-for-token, acceptance 16/16 drafts accepted; `SPEC-MTP` left `GATING`.** **I6 (2026-07-25, `benchmark_binding=true`): the c1 THROUGHPUT A/B vs vLLM same-config LANDED — ours spec-ON is 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). **I7 (2026-07-25): the MIXED spec+non-spec GDN batch (concurrency), the server/CLI `--speculative-config` flag, and the c>1 A/B LANDED** — mixed `GdnBlockPaged` split/merge (new `vt::IndexSelect`/`IndexCopy` op) proven MODEL-INDEPENDENTLY bit-exact (mixed == pure spec + pure prefill, 27B/35B); ours spec-ON keeps helping ~1.43-1.57x at c2/c4/c8 (does not go neutral). Honest c>1 caveat: the 27B greedy is bf16-batch-nondeterministic (spec-OFF max_seqs 4-vs-1 differs 2/3), so exact c>1 token identity is a MODEL impossibility — correctness rests on the bit-exact split/merge proof + acceptance parity, with exact identity at c1. **SPEC-MTP → `DONE` 2026-07-26 (`CLAIM-SPEC-MTP-DONE`, records-only): the user RATIFIED the c>1 criterion (near-tie-distributional + SPEED delta, since strict c>1 token-exact is a proven bf16-batch-nondeterminism model impossibility affecting vLLM identically), closing both I6-owed items; MTP k=1 spec-decode is COMPLETE + gated (27B model row `DONE`; c1 token-exact + above vLLM, c2-c8 on-par-or-above; spec-OFF byte-identical).** 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-ON == our-OFF == vLLM 0.25.0 spec-ON == vLLM spec-OFF, 16/16 vs the `greedy_ids` anchor, acceptance 16/16 both sides; c1 spec-ON 1.19x TPOT / +16.3% tput vs spec-OFF, acceptance 0.908) — MTP now `DONE` on BOTH gate models. Remaining for C3: DFlash; **DSpark explicitly in scope** under `SPEC-DSPARK`; heterogeneous-vocabulary mapping is independently inventoried under `SPEC-TLI`; their spikes queue after parity/MTP. **DFlash readiness RE-ASSESSED 2026-07-25 (`CLAIM-SPEC-DFLASH-READINESS`, design-only) against the landed MTP machinery — verdict GREEN, dispatch-ready, no HW/oracle/download blocker:** `SPEC-DFLASH` gets the frozen ABI + greedy rejection + GDN spec/mixed-batch slot path + draft-KV pattern + runner loop FREE (and `num_lookahead=k+1` already coded), extends the single hidden tap to multi-tap, and builds NEW only the `qwen3_dflash` drafter + the project's first non-causal in-block attention + context-KV precompute + `prepare_dflash_inputs` + uniform-1+k FULL CG; the landed rejection/GDN paths are k-general (DFlash's k=15 blocks need validation/memory-sizing, not mechanism); both z-lab drafts exist on HF (27B 1.73 GB / 35B 368 MB, fit the pool) and the dgx 0.25.0 oracle constructs DFlash. `SPEC-DFLASH` stays `READY` with the D0-D6 W-plan in its spec. **DFlash D0 RUN-VERIFIED the readiness → verdict OVERTURNED to ORACLE-BLOCKED 2026-07-26 (`CLAIM-DFLASH-D0`):** the "oracle constructs DFlash" claim was config-only; a real `LLM(...)` run on dgx ABORTS constructing the mixed-SWA/full z-lab draft at `qwen3_dflash.py:93` `NotImplementedError` (mixed sliding/full via `layer_types` unsupported, upstream vllm#40898) — both Qwen3.6 drafts are mixed + no all-full variant exists, so it is NOT config-fixable. Draft download ungated + config confirmed, reuse map re-verified at `8a379182`. D1–D6 BLOCKED (no vLLM arm for the D5 gate); UNBLOCK = pin > 0.25.0 resolving vllm#40898. `SPEC-DFLASH` → **SPIKE/oracle-BLOCKED** (same class as Gemma-4/OLMo-3). Capture tool ready `scripts/spec/d0_dflash_oracle_capture.py`. **DFlash D0-redo + D1 LANDED 2026-07-26 (`CLAIM-DFLASH-D0D1`) on the advanced pin `555967922`/vLLM 0.26.0.dev0 — UNBLOCKED, `SPEC-DFLASH` `ACTIVE`** (mixed-attn draft constructs + drafter alive under `VLLM_USE_V2_MODEL_RUNNER=1`; D1 `DF-AUX-TAPS` multi-tap done). **DFlash D2 `DF-DRAFT-MODEL` CODE LANDED + CPU-GATED 2026-07-26 (`CLAIM-DFLASH-D2`):** the `qwen3_dflash` draft model, the project's FIRST non-causal / bidirectional in-block attention primitive (`vt::DFlashBlockAttention`, a SEPARATE op — causal attention byte-identical), the fc aux-combine, mask-embed, and the z-lab loader; CPU gate green (op 12/12 incl. RED non-causal + model forward 95/95 incl. RED causal-flip + block isolation + fc RED), existing causal attention/dense-forward tests unchanged. **D2 GPU promotion GREEN on dgx + D3 `DF-DRAFT-KV-PREP` DONE 2026-07-26 (`CLAIM-DFLASH-D3`):** context-KV precompute + `prepare_dflash_inputs` + context-aware draft forward pass `test_qwen3_dflash_kvprep_parity` 61/61 vs the REAL vLLM draft (prepare INTEGER bit-exact vs vLLM's Triton kernel, context-KV K/V rel-L2 0.31%/0.26%, 13 STRICT + 3 near-tie = 16/16), 235/235 + 9/9 + D2 37/37 byte-identical. **D4 `DF-ENGINE-INTEGRATION` propose brick + `dflash` config-select CODE LANDED + CPU-GATED 2026-07-26 (`CLAIM-DFLASH-D4D5`):** `DflashProposeBlock`/`SampleDflashBlockDrafts` (non-autoregressive whole-block propose composing D3 `ForwardBlockLogitsWithContext` + greedy per-mask argmax, anchor not sampled) + `ParseSpeculativeConfigJson`/`ResolveDflash` accept `method:"dflash"`; CPU gate `test_dflash_propose` 5/19 GREEN (RED-first anchor-read fails 4/5); additive + config-gated ⇒ MTP/non-spec byte-identical by construction. **D5 e2e LANDED + RUNS on dgx 2026-07-26 (`CLAIM-DFLASH-D5`):** the runner verify/propose loop is wired and `test_qwen27_dflash_spec_decode` runs — 2/4 STRICT token-exact vs the vLLM-DFlash-ON golden + acceptance ~ vLLM on all 4 (the 2 divergences are single bf16 near-tie flips = the ratified near-tie root, not a wiring bug); spec-OFF byte-identical (SACRED 235/235 + MTP 9/9). **D6 2026-07-27 (`CLAIM-DFLASH-D6`, records-only, NO code):** the **c1 speed A/B is DONE** — our DFlash-ON = **2.50x TPOT / 2.48x output-tput** over our OFF (median TPOT 40.4 vs 101.2 ms, tput 24.4 vs 9.86 tok/s, acceptance 0.22 = 3.56/16, rep-stable <1.5%), `benchmark_binding=true`; vs vLLM-DFlash-ON graphed vLLM-DFlash-ON graphed = 28.5 tok/s / 35.1 ms TPOT / acceptance_len 4.30 (same 8 prompts, `VLLM_USE_V2_MODEL_RUNNER=1`, mm-off, gpu_util 0.30), so OURS IS ~14% BELOW vLLM-DFlash-ON on output throughput (24.4 vs 28.5 tok/s) - both ~on-par at spec-OFF (9.86 vs 9.83 tok/s), but vLLM extracts a larger DFlash speedup (2.90x vs our 2.47x) because its draft step is fully device-resident + CUDA-graphed (ours host-orchestrates 13 downloads/step) + slightly higher acceptance (~4.3 vs ~3.6 draft tokens/step). The DONE speed bar (ours >= vLLM) is NOT met; closing it = the device-resident draft rewrite + FULL CG (D6 part 2). **STRICT-4/4 proven bf16-IRREDUCIBLE** (draft KV is bf16 not fp8 per `torch_utils.py:398`, the D3 golden already compares pre-storage bf16, and a fused KV GEMM is per-element invariant to per-layer ⇒ the ratified near-tie gate is the FINAL correctness form). The **FULL CG + persistent-paged-KV are BLOCKED on a device-resident draft-path rewrite** (13 host downloads/step) — the remaining throughput-parity increment. `SPEC-DFLASH` stays `ACTIVE` (correctness-final + c1-speed-measured); DSpark/TLI remain. **D7-D12 SPEED CAMPAIGN 2026-07-27:** D7 device-resident within-step forward (bit-identical, download hypothesis refuted); D9 (`CLAIM-DFLASH-D9`) persistent paged draft-KV LANDED (bit-identical, +22.7%, ours-ON 0.917× vLLM-ON, acceptance-ceiling REFUTED); D11 (`CLAIM-DFLASH-D11`) FULL-CG Part A device-store primitive; **D12 (`CLAIM-DFLASH-D12`) A-wire (device store is PRODUCTION, e2e 27/27 + SACRED 235/235 + MTP 9/9) + Part B (`vt::DFlashPagedBlockAttention` capture-safe kernel, 795648/795648 CPU==CUDA + cross-check + sanitizer-0) LANDED + GPU-gated.** Speed UNCHANGED 0.917×; Part C (static-shape capture + ≥vLLM c1 A/B) is the sole remaining piece → then SPEC-DFLASH DONE |
+| C3 | MTP speculative decode, starting k=1 on 27B and including the GDN path, then DFlash, DSpark and TLI | 🚧 **MTP k=1 e2e works on the 27B GDN hybrid: the three-way single-request greedy gate PASSES (2026-07-25, I5e) — our-spec-ON == vLLM `--speculative-config mtp` greedy == our-spec-OFF token-for-token, acceptance 16/16 drafts accepted; `SPEC-MTP` left `GATING`.** **I6 (2026-07-25, `benchmark_binding=true`): the c1 THROUGHPUT A/B vs vLLM same-config LANDED — ours spec-ON is 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). **I7 (2026-07-25): the MIXED spec+non-spec GDN batch (concurrency), the server/CLI `--speculative-config` flag, and the c>1 A/B LANDED** — mixed `GdnBlockPaged` split/merge (new `vt::IndexSelect`/`IndexCopy` op) proven MODEL-INDEPENDENTLY bit-exact (mixed == pure spec + pure prefill, 27B/35B); ours spec-ON keeps helping ~1.43-1.57x at c2/c4/c8 (does not go neutral). Honest c>1 caveat: the 27B greedy is bf16-batch-nondeterministic (spec-OFF max_seqs 4-vs-1 differs 2/3), so exact c>1 token identity is a MODEL impossibility — correctness rests on the bit-exact split/merge proof + acceptance parity, with exact identity at c1. **SPEC-MTP → `DONE` 2026-07-26 (`CLAIM-SPEC-MTP-DONE`, records-only): the user RATIFIED the c>1 criterion (near-tie-distributional + SPEED delta, since strict c>1 token-exact is a proven bf16-batch-nondeterminism model impossibility affecting vLLM identically), closing both I6-owed items; MTP k=1 spec-decode is COMPLETE + gated (27B model row `DONE`; c1 token-exact + above vLLM, c2-c8 on-par-or-above; spec-OFF byte-identical).** 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-ON == our-OFF == vLLM 0.25.0 spec-ON == vLLM spec-OFF, 16/16 vs the `greedy_ids` anchor, acceptance 16/16 both sides; c1 spec-ON 1.19x TPOT / +16.3% tput vs spec-OFF, acceptance 0.908) — MTP now `DONE` on BOTH gate models. Remaining for C3: DFlash; **DSpark explicitly in scope** under `SPEC-DSPARK`; heterogeneous-vocabulary mapping is independently inventoried under `SPEC-TLI`; their spikes queue after parity/MTP. **DFlash readiness RE-ASSESSED 2026-07-25 (`CLAIM-SPEC-DFLASH-READINESS`, design-only) against the landed MTP machinery — verdict GREEN, dispatch-ready, no HW/oracle/download blocker:** `SPEC-DFLASH` gets the frozen ABI + greedy rejection + GDN spec/mixed-batch slot path + draft-KV pattern + runner loop FREE (and `num_lookahead=k+1` already coded), extends the single hidden tap to multi-tap, and builds NEW only the `qwen3_dflash` drafter + the project's first non-causal in-block attention + context-KV precompute + `prepare_dflash_inputs` + uniform-1+k FULL CG; the landed rejection/GDN paths are k-general (DFlash's k=15 blocks need validation/memory-sizing, not mechanism); both z-lab drafts exist on HF (27B 1.73 GB / 35B 368 MB, fit the pool) and the dgx 0.25.0 oracle constructs DFlash. `SPEC-DFLASH` stays `READY` with the D0-D6 W-plan in its spec. **DFlash D0 RUN-VERIFIED the readiness → verdict OVERTURNED to ORACLE-BLOCKED 2026-07-26 (`CLAIM-DFLASH-D0`):** the "oracle constructs DFlash" claim was config-only; a real `LLM(...)` run on dgx ABORTS constructing the mixed-SWA/full z-lab draft at `qwen3_dflash.py:93` `NotImplementedError` (mixed sliding/full via `layer_types` unsupported, upstream vllm#40898) — both Qwen3.6 drafts are mixed + no all-full variant exists, so it is NOT config-fixable. Draft download ungated + config confirmed, reuse map re-verified at `8a379182`. D1–D6 BLOCKED (no vLLM arm for the D5 gate); UNBLOCK = pin > 0.25.0 resolving vllm#40898. `SPEC-DFLASH` → **SPIKE/oracle-BLOCKED** (same class as Gemma-4/OLMo-3). Capture tool ready `scripts/spec/d0_dflash_oracle_capture.py`. **DFlash D0-redo + D1 LANDED 2026-07-26 (`CLAIM-DFLASH-D0D1`) on the advanced pin `555967922`/vLLM 0.26.0.dev0 — UNBLOCKED, `SPEC-DFLASH` `ACTIVE`** (mixed-attn draft constructs + drafter alive under `VLLM_USE_V2_MODEL_RUNNER=1`; D1 `DF-AUX-TAPS` multi-tap done). **DFlash D2 `DF-DRAFT-MODEL` CODE LANDED + CPU-GATED 2026-07-26 (`CLAIM-DFLASH-D2`):** the `qwen3_dflash` draft model, the project's FIRST non-causal / bidirectional in-block attention primitive (`vt::DFlashBlockAttention`, a SEPARATE op — causal attention byte-identical), the fc aux-combine, mask-embed, and the z-lab loader; CPU gate green (op 12/12 incl. RED non-causal + model forward 95/95 incl. RED causal-flip + block isolation + fc RED), existing causal attention/dense-forward tests unchanged. **D2 GPU promotion GREEN on dgx + D3 `DF-DRAFT-KV-PREP` DONE 2026-07-26 (`CLAIM-DFLASH-D3`):** context-KV precompute + `prepare_dflash_inputs` + context-aware draft forward pass `test_qwen3_dflash_kvprep_parity` 61/61 vs the REAL vLLM draft (prepare INTEGER bit-exact vs vLLM's Triton kernel, context-KV K/V rel-L2 0.31%/0.26%, 13 STRICT + 3 near-tie = 16/16), 235/235 + 9/9 + D2 37/37 byte-identical. **D4 `DF-ENGINE-INTEGRATION` propose brick + `dflash` config-select CODE LANDED + CPU-GATED 2026-07-26 (`CLAIM-DFLASH-D4D5`):** `DflashProposeBlock`/`SampleDflashBlockDrafts` (non-autoregressive whole-block propose composing D3 `ForwardBlockLogitsWithContext` + greedy per-mask argmax, anchor not sampled) + `ParseSpeculativeConfigJson`/`ResolveDflash` accept `method:"dflash"`; CPU gate `test_dflash_propose` 5/19 GREEN (RED-first anchor-read fails 4/5); additive + config-gated ⇒ MTP/non-spec byte-identical by construction. **D5 e2e LANDED + RUNS on dgx 2026-07-26 (`CLAIM-DFLASH-D5`):** the runner verify/propose loop is wired and `test_qwen27_dflash_spec_decode` runs — 2/4 STRICT token-exact vs the vLLM-DFlash-ON golden + acceptance ~ vLLM on all 4 (the 2 divergences are single bf16 near-tie flips = the ratified near-tie root, not a wiring bug); spec-OFF byte-identical (SACRED 235/235 + MTP 9/9). **D6 2026-07-27 (`CLAIM-DFLASH-D6`, records-only, NO code):** the **c1 speed A/B is DONE** — our DFlash-ON = **2.50x TPOT / 2.48x output-tput** over our OFF (median TPOT 40.4 vs 101.2 ms, tput 24.4 vs 9.86 tok/s, acceptance 0.22 = 3.56/16, rep-stable <1.5%), `benchmark_binding=true`; vs vLLM-DFlash-ON graphed vLLM-DFlash-ON graphed = 28.5 tok/s / 35.1 ms TPOT / acceptance_len 4.30 (same 8 prompts, `VLLM_USE_V2_MODEL_RUNNER=1`, mm-off, gpu_util 0.30), so OURS IS ~14% BELOW vLLM-DFlash-ON on output throughput (24.4 vs 28.5 tok/s) - both ~on-par at spec-OFF (9.86 vs 9.83 tok/s), but vLLM extracts a larger DFlash speedup (2.90x vs our 2.47x) because its draft step is fully device-resident + CUDA-graphed (ours host-orchestrates 13 downloads/step) + slightly higher acceptance (~4.3 vs ~3.6 draft tokens/step). The DONE speed bar (ours >= vLLM) is NOT met; closing it = the device-resident draft rewrite + FULL CG (D6 part 2). **STRICT-4/4 proven bf16-IRREDUCIBLE** (draft KV is bf16 not fp8 per `torch_utils.py:398`, the D3 golden already compares pre-storage bf16, and a fused KV GEMM is per-element invariant to per-layer ⇒ the ratified near-tie gate is the FINAL correctness form). The **FULL CG + persistent-paged-KV are BLOCKED on a device-resident draft-path rewrite** (13 host downloads/step) — the remaining throughput-parity increment. `SPEC-DFLASH` stays `ACTIVE` (correctness-final + c1-speed-measured); DSpark/TLI remain. **D7-D12 SPEED CAMPAIGN 2026-07-27:** D7 device-resident within-step forward (bit-identical, download hypothesis refuted); D9 (`CLAIM-DFLASH-D9`) persistent paged draft-KV LANDED (bit-identical, +22.7%, ours-ON 0.917× vLLM-ON, acceptance-ceiling REFUTED); D11 (`CLAIM-DFLASH-D11`) FULL-CG Part A device-store primitive; **D12 (`CLAIM-DFLASH-D12`) A-wire (device store is PRODUCTION, e2e 27/27 + SACRED 235/235 + MTP 9/9) + Part B (`vt::DFlashPagedBlockAttention` capture-safe kernel, 795648/795648 CPU==CUDA + cross-check + sanitizer-0) LANDED + GPU-gated.** Speed UNCHANGED 0.917×; Part C (static-shape capture + ≥vLLM c1 A/B) is the sole remaining piece → then SPEC-DFLASH DONE. **RECONCILED 2026-08-12 ([#536](https://github.com/mudler/vllm.cpp/issues/536)): Part C LANDED, `SPEC-DFLASH` is `DONE` (D13/D14, `489a7544`), so the two clauses above that say "DSpark explicitly in scope … their spikes queue after parity/MTP" and "DSpark/TLI remain" are SUPERSEDED, not deleted.** **DSpark: spike 2026-08-09 (`2b342620e`), `SPEC-DSPARK` `ACTIVE`, W1–W8 implemented and GPU-gated** — Markov head, sequential sampler (device-side since [#436](https://github.com/mudler/vllm.cpp/issues/436)), native AND Speculators-format loading, `d2t` reduced draft vocab, runner/one-surface wiring, and the T=1+k verify capture ([#442](https://github.com/mudler/vllm.cpp/issues/442)) mirroring vLLM's `uniform_decode_query_len = 1 + num_speculative_tokens`; plus the engine-wide `EngineCoreProc` `check_for_draft_tokens` fix that had been silently dropping EVERY speculator's drafts on the CLI/server paths. Cross-engine under pinned clocks: **35B-A3B MoE 0.975x code / 1.012x prose vs the pinned graphed oracle — NOT parity**, residual attributed to a 12.9% effective-DRAM-bandwidth gap inside `marlin_moe_wna16::Marlin` on byte-equivalent machine code (spec §§6t–6aa). **`SPEC-TLI` is genuinely untouched — no commit, code, spec or issue — and it is NOT a DSpark tail:** upstream's `use_heterogeneous_vocab` (`config/speculative.py:150`) + `VocabMapping` (`v1/spec_decode/vocab_mapping.py:68`) are consumed only by `SpecDecodeBaseProposer` (`v1/spec_decode/llm_base_proposer.py`) and `v1/spec_decode/draft_model.py:19`, never by the V2-runner speculators our DFlash/DSpark port mirrors, so its host row is `SPEC-DRAFT-MODEL` and it is prerequisite-blocked behind that row's W3. DSpark's `d2t` does not cover it (one tokenizer's vocab offset table vs a cross-tokenizer-family string intersection) |
 | C4 | FP8 W8A8 quantization breadth | ☐ T1; [feature matrix §5](feature-matrix.md#5-quantization) |
 | C5 | Sliding-window KV/attention and YaRN long-context scaling | 🚧 [joint spike accepted](specs/sliding-local-yarn-long-context.md): all W1-W8 implementation leaves are `GATING`; W5-W8 typed YaRN/MRoPE/Llama 3/Phi-3 LongRoPE/dynamic-NTK cache construction passes fifteen pinned-source CPU oracle fixtures and sanitizers. CUDA compile/runtime and all model/trace/every-axis gates remain open; gate Qwen3.6 checkpoints are default interleaved MRoPE regressions, with separate feature-positive e2e gates |
 | C6 | Priority + async/overlap scheduling + AsyncLLM streaming serving | 🚧 joint spike written: [async-serving.md](specs/async-serving.md); W1/W2/W4 are implemented and `GATING`. W2 fixed capacity is CPU/sanitizer-green and GPU-classified healthy/steady-state-neutral; broader every-axis parity remains. **W3 host-side + runner device-input + sampler-OUTPUT half + the ENABLE-FLIP all LANDED + CPU-gated (2026-07-16, `ACTIVE`)** — `combine_sampled_and_draft_tokens` (input) + `vt::Backend` event/pinned + `AsyncGPUModelRunnerOutput` non-blocking D2H + `sample_tokens_async` + the depth-2 async-output seam + the `LoadedEngine` `AsyncScheduler`+mcb=2 construction switch + the `cuda_combine_tokens.cu` device combine/scatter kernel (all `VT_ASYNC_RUNNER` default OFF; no-env production byte-identical). The FULL W3 DGX proof RAN at `f086b64` (5/5 gates PASS; c16 W3-on TPOT −5.4 ms/step but TTFT +36 %/+730 ms, throughput neutral) — the TTFT regression is the closed-loop Little's-law consequence of neutral throughput + faster decode, NOT an admission bug (`test_async_admission_timing.cpp` proves depth-2 admits 1:1 with sync/vLLM). W3 stays **default-OFF** until a depth-2 throughput lever lands; `runner_supports_async` default-ON stays DGX-gated. Order-0 FP4 closure precedes W3 default-ON overlap; [feature matrix §1](feature-matrix.md#1-engine-core--scheduling) |
@@ -397,8 +401,12 @@ method `"dspark"`, V2-runner-only) with draft models for BOTH our registered
 target families (Qwen3 `qwen3_dspark.py:95`, Gemma4 `gemma4_dspark.py:134`) —
 re-grounded at the pin and slotted beside our landed MTP + DFlash lanes:
 [dspark-speculator-note.md](specs/dspark-speculator-note.md), row
-`SPEC-DSPARK` (engine matrix, `INVENTORIED`); full scope is its own future
-spike (`planned: specs/dspark-spec-decode.md`).
+`SPEC-DSPARK` (engine matrix, ~~`INVENTORIED`~~ **`ACTIVE` since 2026-08-09**);
+~~full scope is its own future spike (`planned: specs/dspark-spec-decode.md`)~~
+— **that spike was WRITTEN (`2b342620e`) and its W1–W8 are implemented and
+GPU-gated: [dspark-spec-decode.md](specs/dspark-spec-decode.md)**
+(superseded in place 2026-08-12,
+[#536](https://github.com/mudler/vllm.cpp/issues/536)).
 
 ## vLLM feature-gap sweep (2026-07-28, `CLAIM-FEATURE-GAP-SPIKE`)
 
diff --git a/.agents/specs/dspark-spec-decode.md b/.agents/specs/dspark-spec-decode.md
index d08b789d5..d41687e47 100644
--- a/.agents/specs/dspark-spec-decode.md
+++ b/.agents/specs/dspark-spec-decode.md
@@ -17,7 +17,7 @@
 | Dependencies | Landed: `SPEC-DFLASH` (`DONE`), `SPEC-REJECTION` verify half, `SPEC-GDN-SEGMENTS`. External, PENDING developer authority: checkpoint downloads (2.79-8.80 GB), dgx GPU time, push/draft-PR. Blocking unknown: R1, whether the pinned oracle runs DSpark at all. |
 | Work breakdown | §4 — W1 config, W2 Markov head + draft model, W3 loader (native + Speculators), W4 speculator (anchor layout + sequential sampling), W5 runner + one-surface, W6 gates. W1-W4 are CPU-gateable. |
 | Risks/decisions | §6 — R1 oracle runnability (V2 runner), R2 Speculators format is a new subsystem, R3 the community 27B checkpoint's `attn_output_gate`, R4 the `k >= dspark_block_size` garbling trap, R5 sequential sampling vs CUDA-graph capture, R6 greedy before probabilistic, R7 GB10 host-RAM pressure. |
-| Status | W1-W5 LANDED and DSpark now genuinely speculates on the 35B gate model (real acceptance, 6.78 -> 41.89 tok/s) after fixing an engine-wide `check_for_draft_tokens` wiring bug that silently disabled EVERY speculator on the CLI/server path. W6 PARTIAL: spec-ON output is token-identical to spec-OFF and reproducible on the 35B gate model (§6b), but speed is ~2% BEHIND spec-off and the cross-engine + acceptance-band gates are still owed. R1 answered (§6a). |
+| Status | **`ACTIVE`. W1-W8 LANDED and GPU-gated** (last refreshed 2026-08-12, [#536](https://github.com/mudler/vllm.cpp/issues/536); §8 records why this field had drifted). W1-W5 gave a working speculator on the 35B gate model after fixing an engine-wide `check_for_draft_tokens` wiring bug that silently disabled EVERY speculator on the CLI/server path (§6b); W7 ([#436](https://github.com/mudler/vllm.cpp/issues/436)) moved the sequential Markov sample on device, byte-identical (§6k); W8 ([#442](https://github.com/mudler/vllm.cpp/issues/442)) captured the T=1+k verify by mirroring vLLM's uniform-decode dispatch predicate, +12.2%/+3.5% same-binary (§§6m-6n). **Cross-engine under pinned clocks the 35B-A3B MoE lane is 0.975x (code cell, non-overlapping) / 1.012x (prose cell) — NOT parity**, with the residual localised to `marlin_moe_wna16::Marlin` and attributed to a 12.9% effective-DRAM-bandwidth gap on byte-equivalent machine code (§§6s-6aa). R1 answered (§6a). Still owed for a binding W6: the SACRED-corpus token gate under the ratified near-tie protocol, the 27B dense re-measure, the Gemma4 `1 + N` layout on real weights, padded/multi-request capture shapes, and the next bandwidth lever (§6aa). ~~W1-W5 LANDED … W6 PARTIAL: … speed is ~2% BEHIND spec-off …~~ (the 2026-08-10 text, superseded by §§6c-6aa; the "~2% behind" was a cold single-shot reading corrected in §6c). |
 | Goal (developer, 2026-08-09) | a FULL DSpark implementation in vllm.cpp, mirrored from vLLM |
 
 ## 0. Verdict
@@ -1592,3 +1592,46 @@ and attributed to achieved memory bandwidth rather than left as "unexplained".
 
 Sources: pin files cited inline at `555967922`; HF API queried 2026-08-09;
 our anchors cited against `bc6e3d72`.
+
+## 8. Record reconciliation (2026-08-12, [#536](https://github.com/mudler/vllm.cpp/issues/536))
+
+`ROAD-V1-C3`'s named tail still read "DSpark (`SPEC-DSPARK`) +
+heterogeneous-vocabulary TLI (`SPEC-TLI`) **unspiked** — overlaps `ROAD-V1-D3`"
+three days after this row's W1-W8 landed and were measured cross-engine. Every
+clause was wrong, and a punch-list reader would have dispatched a port that
+already exists.
+
+| Surface | Said | Now |
+|---|---|---|
+| [roadmap-v1-completion.md](roadmap-v1-completion.md) §2 C3 row, §3 item 17 | DSpark + TLI "unspiked", "overlaps D3", Size M | superseded in place; DSpark is a perf tail plus owed gates (S-M), TLI is a separate row (M) |
+| [roadmap_v1.md](../roadmap_v1.md) rows 3 and C3 | "DFlash Part C + DSpark/TLI remain"; "their dedicated spikes are not written" | `SPEC-DFLASH` is `DONE` (D13/D14, `489a7544`); this spike exists (`2b342620e`) |
+| [roadmap_v1.md](../roadmap_v1.md) §DSpark grounding note | `SPEC-DSPARK` "(engine matrix, `INVENTORIED`)", spike "future" | `ACTIVE`, spike written |
+| [spec-decode-inventory.md](spec-decode-inventory.md) `dspark` row + lifecycle summary | **INVENTORIED** | **ACTIVE** |
+| [spec-decode-inventory.md](spec-decode-inventory.md) §HF speculators | "no `speculators`-format adapter" | W3 shipped the DSpark one (`qwen3_dspark.cpp:227-300`) |
+| [docs/STATUS.md](../../docs/STATUS.md) method surface | `dspark` INVENTORIED, contradicting the same page's own DSpark paragraph | ships DSpark |
+| This spec's `Status` field | 2026-08-10 "W1-W5 … W6 PARTIAL … ~2% BEHIND spec-off" | W1-W8, 0.975x/1.012x cross-engine |
+| [roadmap_v1.md](../roadmap_v1.md) Open issues | #436, #442, #513 absent from the intake table | listed |
+
+**The TLI half is filed under the wrong row, and that is the more useful
+finding.** Upstream TLI is `use_heterogeneous_vocab`
+(`config/speculative.py:150`) plus `VocabMapping`
+(`v1/spec_decode/vocab_mapping.py:68`), and both are consumed ONLY by
+`v1/spec_decode/llm_base_proposer.py:432-495,688-691,831-837`
+(`SpecDecodeBaseProposer`) and `v1/spec_decode/draft_model.py:19,34-58`. The
+V2-runner speculators this row ports (`v1/worker/gpu/spec_decode/dspark/`,
+`.../dflash/`) contain no heterogeneous-vocab path at all, so nothing DSpark
+lands moves TLI, and TLI's host is `SPEC-DRAFT-MODEL` — locally a CPU propose
+brick with no runner construction. In particular this row's `d2t` work is NOT
+TLI: `d2t` offsets ids inside ONE tokenizer's vocabulary
+(`draft_id + d2t[draft_id]`, §2 D), while `VocabMapping` builds a string-level
+intersection ACROSS tokenizer families, probing the space prefix at init to
+handle a BPE draft against a SentencePiece target.
+
+**Why the drift happened, since the row can be told from the outside.** The
+punch-list is written from row summaries, and this row's summaries were the
+thing not updated: the spike went straight from "planned" to twenty-odd
+`measure(SPEC-DSPARK)` commits without either the inventory's lifecycle summary
+or `STATUS.md`'s one-line method surface following. Both are derived statements
+about a row that lives elsewhere, which is exactly the shape AGENTS.md's Records
+section warns about — a fact stored away from the thing it describes drifts
+silently, because nothing fails when it does.
diff --git a/.agents/specs/roadmap-v1-completion.md b/.agents/specs/roadmap-v1-completion.md
index c8184459f..eef4fd172 100644
--- a/.agents/specs/roadmap-v1-completion.md
+++ b/.agents/specs/roadmap-v1-completion.md
@@ -37,7 +37,7 @@ carry more than one class list the dominant one first.
 | `ROAD-V1-C1` extensibility | **DONE** (cornerstone) | Drop-in kernel ABI W0, Platform seam, model self-registration, and the **portable op-fusion framework ORDER-1 milestone** (W0–W4 merged+gated, `KERNEL-FUSION-FRAMEWORK`); consistency-audit CI check landed. `BACKEND-ABI-VT`/`BACKEND-CUDA-ARCH-ADDITIVITY` seams gated on sm_121a. | Row stays SPIKE-open only for **non-blocking** tail: Tier-1 fusion perf interpreter (composite-only → single-launch), `FUSION-DENSE-MIGRATE` (route 5 drift models — CLOSED 2026-08-10, [#299](https://github.com/mudler/vllm.cpp/issues/299)), a real Metal/Vulkan catalog realization (M4-reachable / HW-blocked), and migrating a production kernel family onto the common adapter. Correctness cornerstone is closed. |
 | `ROAD-V1-C2` model families | **RI** (+HW/EXT sub) | First additive model (Qwen3 dense) + a broad **text sweep correctness-complete + SACRED-gated**: Qwen3/Qwen3Moe/Coder, Llama/Yi/InternLM3, Mistral, GLM-4-9B/GLM-4.7-Flash, Gemma-1/2/3, OPT, DeepSeek-V2-Lite (MLA), OLMo-2, Phi-3/4, Phi-1/2, Granite-3, StableLM, InternLM2, MiniCPM, MiniCPM3 (MLA). 20 ACTIVE model rows. | **SPEED close** on every one (all 20 are "correctness-complete, speed pending"). MoE/SSM breadth (Qwen3-Next, Falcon, Falcon-H1, GraniteMoe*, Cohere2Moe, PhiMoE, Mamba/Jamba/Zamba2/NemotronH) = RI (INVENTORIED/SPIKE). Frontier: Kimi-Linear-48B fits (RI, +KDA kernel); DeepSeek-V3/GLM-5/MiniMax-M2/M3/Kimi-K2 = HW (>119 GiB); Command-R = EXT (HF token). |
 | `ROAD-V1-C2-LOCAL-BF16` | **RI** (S) | Local Qwen3.5-4B plain-BF16 diagnostic rebased onto current additive seams; CPU/CUDA + direct ON/OFF token-equivalence green; H32 AOT / plain-BF16 graphs / ratio-4 FA2 landed + trace-proven. | Port device-resident sampled-token mapping to discrete CUDA (remove the measured main-stream wait) and rerun the exact 4B series. Small. |
-| `ROAD-V1-C3` spec-decode | **DONE** (core) | **MTP k=1 DONE + gated on BOTH gate models** (`SPEC-MTP`, c1 token-exact + above vLLM, c2–c8 on-par-or-above); **DFlash DONE + speed gate MET** (`SPEC-DFLASH` D14, our-ON ≥ vLLM-ON). | Named tail only: DSpark (`SPEC-DSPARK`) + heterogeneous-vocabulary TLI (`SPEC-TLI`) unspiked — overlaps `ROAD-V1-D3`. Core spec-decode is gate-closed. |
+| `ROAD-V1-C3` spec-decode | **DONE** (core) | **MTP k=1 DONE + gated on BOTH gate models** (`SPEC-MTP`, c1 token-exact + above vLLM, c2–c8 on-par-or-above); **DFlash DONE + speed gate MET** (`SPEC-DFLASH` D14, our-ON ≥ vLLM-ON). | ~~Named tail only: DSpark (`SPEC-DSPARK`) + heterogeneous-vocabulary TLI (`SPEC-TLI`) unspiked — overlaps `ROAD-V1-D3`.~~ **SUPERSEDED 2026-08-12** ([#536](https://github.com/mudler/vllm.cpp/issues/536), see §3 item 17): `SPEC-DSPARK` is `ACTIVE` — W1–W8 implemented and GPU-gated, 35B-A3B MoE **0.975x** code / **1.012x** prose vs the pinned graphed oracle (#442), remaining work a perf tail plus owed gates. `SPEC-TLI` is genuinely untouched and belongs under `SPEC-DRAFT-MODEL`, whose W3 blocks it. `ROAD-V1-D3` excludes both by its own spec, so it overlaps nothing here. Core spec-decode is gate-closed. |
 | `ROAD-V1-C4` quantization | **RI** | **3 schemes DONE**: NVFP4-MO-W4A16, NVFP4-CT-W4A4, FP8-MO-STATIC (all R/M/C/E/P). **GGUF CPU vs llama.cpp is CLOSED** (2026-07-22, aarch64 binding host): decode **at parity** (1.03× behind, inside llama.cpp's ±1.8% run spread — the elementwise f16/bf16 GEMM lever `KERNEL-GEMM-CPU-ELEM` E1-E4 `18094ee2` took it 3.38×→1.03×), prefill **1.18× ahead** (q8_0 repack-at-load G7), RSS **1.01×**, byte-identical greedy tokens. | NVFP4-CT-W4A16 perf gate. FP8-generic dispatch (static/dyn × tensor/channel/token/block). Breadth: AWQ/GPTQ/Marlin-wiring, i-quants, MXFP4/MX, bitsandbytes, KV-quant — all INVENTORIED. (GGUF-vs-llama.cpp speed is no longer an open C4 blocker.) |
 | `ROAD-V1-C5` sliding/YaRN | **RI** | Joint spike accepted; all W1–W8 leaves implemented and CPU/oracle/sanitizer green. **CUDA GPU CLOSURE 2026-07-27 (`CLAIM-ROADMAP-C5`, dgx GB10 sm_121a, clean build of `489f7771`, oracle vLLM 0.26.0.dev0):** shared scaled-RoPE + local-mask CUDA path compiles `-Werror`-clean + RUNS on GB10; feature-positive correctness gates PASS — SWA Gemma-2/Gemma-3 48/48, LongRoPE Phi-4-mini 16/16 (RED-first), llama3 Llama-3.2-1B 16/16, dynamic-NTK InternLM2 16/16; both RoPE 0.26-oracle recaptures BIT-IDENTICAL to goldens. `ATTN-SLIDING-WINDOW`/`ATTN-ROPE-{LLAMA3,LONGROPE,DYNAMIC-NTK}`/`ATTN-YARN` → `ACTIVE`; `ATTN-CHUNKED-LOCAL` + `KV-*-SPEC` honest. | **Honest residual (vehicle-blocked, not skipped):** YaRN model e2e (no cached Nomic/gpt-oss consumer) + chunked-local model e2e (no Llama4 row) REACHABLE-BLOCKED; long-context positive-mask (prompt > W) SWA e2e + KV-memory G8; every-axis **SPEED** tail (all leaves correctness-complete, speed-pending). |
 | `ROAD-V1-C6` async/priority serving | **RI** | `ENG-ASYNC-SCHED` **DONE** (`6ea7856`, default-ON, DGX token-neutral). W1/W2/W4 landed. | `SERVE-ASYNC-LLM` (GATING → prod-ON, blocks the SGLang floor + `ROAD-V1-A`), `ENG-PRIORITY-SCHED` + `ENG-CORE-BUSY-LOOP` GPU gates (GATING, held behind SERVE-GATE-ONLINE). |
@@ -213,7 +213,52 @@ gate → size (S/M/L) → vehicle model. `[H]` = user-directed headline.
     leaving only entries that need the shared layer extended. The dgx paged-engine
     confirmation for those five is OWED.
 16. **`ROAD-V1-C9` 0.26 denominators/goldens refresh** (recurring). **Size S, ongoing.**
-17. **`ROAD-V1-C3` DSpark + TLI** (core spec-decode done; overlaps D3). **Size M.**
+17. ~~**`ROAD-V1-C3` DSpark + TLI** (core spec-decode done; overlaps D3). **Size M.**~~
+    **RECONCILED 2026-08-12** ([#536](https://github.com/mudler/vllm.cpp/issues/536)):
+    the item was written when both halves were untouched. Neither half is what
+    it says, and they are not one item.
+    **DSpark is not unspiked.** Its spike spec landed 2026-08-09 (`2b342620e`,
+    [dspark-spec-decode.md](dspark-spec-decode.md)), `SPEC-DSPARK` has been
+    `ACTIVE` since, and W1–W8 are implemented and GPU-gated: the Markov head,
+    the sequential sampler, native **and** Speculators-format loading, the `d2t`
+    reduced draft vocab, the runner/one-surface wiring, the device sequential
+    sample ([#436](https://github.com/mudler/vllm.cpp/issues/436)) and the T=1+k
+    verify capture ([#442](https://github.com/mudler/vllm.cpp/issues/442),
+    mirroring vLLM's `uniform_decode_query_len = 1 + num_speculative_tokens`).
+    En route it fixed an engine-wide defect: `EngineCoreProc` never threaded
+    `check_for_draft_tokens`, so EVERY speculator's drafts were dropped on the
+    CLI and server paths. Measured against the pinned graphed oracle under
+    pinned clocks the 35B-A3B MoE lane is **0.975x** (code cell,
+    non-overlapping distributions) and **1.012x** (prose cell), with the
+    residual localised to one kernel and attributed to a **12.9%
+    effective-DRAM-bandwidth** gap on byte-equivalent machine code (94 registers
+    / 3664 SASS instructions on both sides, spec §§6t–6aa). **Remaining is a
+    perf tail plus owed gates, not a port: Size S–M**, and its next lever is
+    named (`cudaMemAdvise`/placement on the expert slab; upstream `ncu` counters
+    are BLOCKED in both replay modes, so a standalone `moe_wna16_marlin_gemm`
+    harness is the only remaining route). Owed for a binding W6: the SACRED-corpus
+    token gate under the ratified near-tie protocol, the 27B dense re-measure
+    (its earlier cells were never like-for-like), the Gemma4 `1 + N` layout on
+    real weights, and padded/multi-request spec capture shapes.
+    **TLI is untouched — and it is not a DSpark tail.** No commit, no code, no
+    spec, no issue; `SPEC-TLI` is `INVENTORIED`. Upstream TLI is
+    `use_heterogeneous_vocab` (`config/speculative.py:150`) + `VocabMapping`
+    (`v1/spec_decode/vocab_mapping.py:68`), consumed ONLY by
+    `v1/spec_decode/llm_base_proposer.py` (`SpecDecodeBaseProposer`) and
+    `v1/spec_decode/draft_model.py:19`. The V2-runner speculators our
+    DFlash/DSpark port mirrors (`v1/worker/gpu/spec_decode/{dflash,dspark}/`)
+    have no heterogeneous-vocab path at all, so TLI's host is `SPEC-DRAFT-MODEL`
+    — a CPU propose brick with no runner construction — and TLI is
+    prerequisite-blocked behind that row's W3, not merely unspiked. DSpark's
+    `d2t` does not cover it: `d2t` is an offset table inside ONE tokenizer's
+    vocabulary (`draft_id + d2t[draft_id]`), while TLI builds a string-level
+    intersection ACROSS tokenizer families (BPE `Ġ` vs SentencePiece `▁`,
+    probed at init). **Size M, and it should be re-filed under `SPEC-DRAFT-MODEL`.**
+    **The D3 overlap is backwards.**
+    [spec-decode-breadth-d3.md](spec-decode-breadth-d3.md) §Scope puts DSpark and
+    TLI explicitly *out of* `ROAD-V1-D3` and back under `ROAD-V1-C3`; D3's
+    landing covers no part of this tail. What DSpark reused is C3's own
+    MTP/DFlash verify/reject loop.
 
 ## 4. Bottom line
 
diff --git a/.agents/specs/spec-decode-inventory.md b/.agents/specs/spec-decode-inventory.md
index 49d3ff854..c916e4f0f 100644
--- a/.agents/specs/spec-decode-inventory.md
+++ b/.agents/specs/spec-decode-inventory.md
@@ -42,7 +42,7 @@ vLLM, it is not an addition vLLM lacks.
 | `mtp` (canonical) | `EagleProposer` via `use_eagle()`; special `Gemma4Proposer`/`Step3p5MTPProposer` (`gpu_model_runner.py:627-630`); MRV2 `MTPSpeculator` | head-on-base per family (nextn layer on target) | V1 default + MRV2 | **DONE** `SPEC-MTP` (Qwen3.5/3.6 k=1); family breadth `SPEC-MTP-FAMILY` |
 | `deepseek_mtp` ... `inkling_mtp` (20 family strings) | `MTPModelTypes` `speculative.py:37-59` -> deprecate-remap to `mtp` (`:686-690`); models `deepseek_mtp.py`, `glm4_moe_mtp.py`, `ernie_mtp.py`, `nemotron_h_mtp.py`, ... | head-on-base per family | V1 default (dispatched by `draft_model_config.hf_config.model_type`) | **ACTIVE** DeepSeek-V4 (`MODEL-SPEC-deepseek-v4-deep-seek-v4-mtp`, W1, weight-blocked); rest **INVENTORIED** (`SPEC-MTP-FAMILY`, model-matrix `MODEL-SPEC-*-mtp`) |
 | `dflash` | `vllm/v1/spec_decode/dflash.py` (DFlashProposer); models `laguna_dflash.py`, `qwen3_dflash.py`; speculators algo `algos.py:93` | separate block-diffusion draft + aux | V1 default + MRV2 | **DONE** `SPEC-DFLASH` (ported from vLLM; +GGUF `SPEC-DFLASH-GGUF`) |
-| `dspark` | MRV2 `vllm/v1/worker/gpu/spec_decode/dspark/speculator.py:37`; models `gemma4_dspark.py`, `qwen3_dspark.py`; speculators algo `algos.py:133` | separate SAR block draft + aux | V1 default (via `use_eagle()`) + MRV2 | **INVENTORIED** `SPEC-DSPARK` |
+| `dspark` | MRV2 `vllm/v1/worker/gpu/spec_decode/dspark/speculator.py:37`; models `gemma4_dspark.py`, `qwen3_dspark.py`; speculators algo `algos.py:133` | separate SAR block draft + aux | V1 default (via `use_eagle()`) + MRV2 | ~~**INVENTORIED**~~ **`ACTIVE`** `SPEC-DSPARK` (W1–W8 implemented + GPU-gated; 35B-A3B MoE 0.975x code / 1.012x prose vs the pinned graphed oracle, #442; superseded here 2026-08-12, #536) |
 
 ## User-visible config axes (apply on top of a method)
 
@@ -51,25 +51,34 @@ vLLM, it is not an addition vLLM lacks.
 | Acceptance / rejection sampler | `speculative.py:77,216` (`rejection_sample_method`) | `standard` \| `synthetic` \| `block` | `standard` **ACTIVE** `SPEC-REJECTION`; `synthetic`/`block` **ABSENT** `SPEC-ACCEPT-VARIANTS` |
 | Draft sampling | `speculative.py:78,283` (`draft_sample_method`) | `greedy` \| `probabilistic` | `greedy` only; `probabilistic` **ABSENT** `SPEC-ACCEPT-VARIANTS` |
 | Dynamic k (per batch size) | `speculative.py:1336-1337`; `vllm/v1/spec_decode/dynamic/utils.py` (`num_speculative_tokens_per_batch_size`) | schedule list | **ABSENT** `SPEC-DYNAMIC` |
-| Heterogeneous draft/target vocab | `vllm/v1/spec_decode/vocab_mapping.py` | shared-token / ID translation | **INVENTORIED** `SPEC-TLI` |
+| Heterogeneous draft/target vocab (TLI) | `vllm/config/speculative.py:150` (`use_heterogeneous_vocab`); `vllm/v1/spec_decode/vocab_mapping.py:68` (`VocabMapping`); consumed ONLY by `llm_base_proposer.py:432-495,688-691,831-837` (`SpecDecodeBaseProposer`) and `draft_model.py:19,34-58` | cross-tokenizer-family string intersection, target↔draft ID translation, constrained draft logits | **INVENTORIED** `SPEC-TLI` — untouched, and it hangs off the V1 `SpecDecodeBaseProposer`, NOT the V2-runner DFlash/DSpark speculators we ported, so its host row is `SPEC-DRAFT-MODEL` and it is prerequisite-blocked behind that row's W3 (2026-08-12, #536). DSpark's `d2t` does NOT cover it: `d2t` offsets ids inside ONE vocabulary |
 | GDN spec metadata + slot-snapshot rollback | `vllm/v1/attention/backends/gdn_attn.py` | — | **ACTIVE** `SPEC-GDN-SEGMENTS` |
 
 ## HF `speculators` checkpoint-format integration
 
 `vllm/transformers_utils/configs/speculators/algos.py` registers
 `SUPPORTED_SPECULATORS_TYPES = {eagle3, peagle, dflash, dspark}` (`:15,55,93,133`)
-and rewrites a HF `speculators`-format config into a vLLM draft arch. Our side
+and rewrites a HF `speculators`-format config into a vLLM draft arch. ~~Our side
 loads native draft configs directly (no `speculators`-format adapter); tracked
 under the EAGLE/DFlash/DSpark rows as a loader-format residual, not a separate
-method.
+method.~~ **SUPERSEDED 2026-08-12 (#536): `SPEC-DSPARK` W3 landed the DSpark
+adapter** — `speculators_model_type == "dspark"` detection, the
+`speculators_config` proposal-method unwrap, and the `algos.py:133-165` field
+translation (`aux_hidden_state_layer_ids` → `target_layer_ids = [i-1]`,
+`sample_from_anchor`, `draft_vocab_size`, `mask_token_id`, `markov_rank`,
+`block_size`), in `src/vllm/model_executor/models/qwen3_dspark.cpp:227-300`.
+It is the path the `RedHatAI/*speculator.dspark` gate-model drafts load through.
+The residual is now the EAGLE3/PEAGLE/DFlash algos, not the whole subsystem.
 
 ## Our lifecycle summary (engine matrix §Speculative decoding)
 
 - **DONE**: `SPEC-MTP`, `SPEC-MTP-GGUF`, `SPEC-DFLASH`, `SPEC-DFLASH-GGUF`.
 - **ACTIVE**: `SPEC-REJECTION`, `SPEC-GDN-SEGMENTS`, `SPEC-NGRAM`,
-  `SPEC-DRAFT-MODEL`.
+  `SPEC-DRAFT-MODEL`, **`SPEC-DSPARK`** (moved out of INVENTORIED here
+  2026-08-12, #536: it has been `ACTIVE` in the engine matrix since the
+  2026-08-09 spike, with W1–W8 implemented and GPU-gated).
 - **SPIKE**: `SPEC-MEDUSA`. **BLOCKED**: `SPEC-EAGLE3`.
-- **INVENTORIED**: `SPEC-DSPARK`, `SPEC-TLI`, and the nine enumerated here:
+- **INVENTORIED**: ~~`SPEC-DSPARK`,~~ `SPEC-TLI`, and the nine enumerated here:
   `SPEC-NGRAM-GPU`, `SPEC-SUFFIX`, `SPEC-EAGLE`, `SPEC-MTP-FAMILY`,
   `SPEC-ACCEPT-VARIANTS`, `SPEC-DYNAMIC`, `SPEC-CUSTOM-CLASS`,
   `SPEC-EXTRACT-HIDDEN`, `SPEC-MLP-SPECULATOR`.
diff --git a/docs/STATUS.md b/docs/STATUS.md
index e3ed99f03..cb23cca73 100644
--- a/docs/STATUS.md
+++ b/docs/STATUS.md
@@ -169,7 +169,7 @@ concurrency-1 A/B our-on 29.32 tok/s vs vLLM-on 29.24, non-overlapping bands,
 vLLM 0.26.0.dev0 stack (which resolves vllm#40898), and it remains gated behind
 a spike while its user-facing serving surface is finalized.
 
-**Method surface (enumerated from vLLM source 2026-08-06, `.agents/specs/spec-decode-inventory.md`).** Of the 13 vLLM `SpeculativeMethod` strings we ship MTP (k=1), DFlash and n-gram; draft_model is a CPU brick and Medusa a spike; EAGLE1/EAGLE3, ngram-gpu, suffix, dspark, custom_class, extract_hidden_states, dynamic-k and the synthetic/block acceptance variants are INVENTORIED; mlp_speculator is upstream-deprecated (no V1 proposer). Draft DEPTH (k>1, dynamic, adaptive) unbuilt (`ROAD-V1-D3-SPEC-K`, #81).
+**Method surface (enumerated from vLLM source 2026-08-06, `.agents/specs/spec-decode-inventory.md`).** Of the 13 vLLM `SpeculativeMethod` strings we ship MTP (k=1), DFlash, DSpark and n-gram; draft_model is a CPU brick and Medusa a spike; EAGLE1/EAGLE3, ngram-gpu, suffix, custom_class, extract_hidden_states, dynamic-k and the synthetic/block acceptance variants are INVENTORIED; mlp_speculator is upstream-deprecated (no V1 proposer). Draft DEPTH (k>1, dynamic, adaptive) unbuilt (`ROAD-V1-D3-SPEC-K`, #81).
 
 **DeepSeek-V4 native MTP** (`DeepSeekV4MTPModel`, ACTIVE — W1 self-spec wiring,
 2026-07-30) has its nextn draft head wired to the same lossless spec-decode path.