Kernel-family matrix

August 24, 2026 · View on GitHub

This is the canonical practical kernel-family inventory for roadmap C1 and the backend expansion blocks. The accepted inventory spike is specs/kernel-family-inventory.md. Upstream references are pinned to vLLM 555967922 (0.26.0.dev0 + transformers 5.14.1; advanced 2026-07-26 from the prior e24d1b24/0.25.0 pin, see specs/pin-advance.md) and include the dependency that owns execution when vLLM only orchestrates dispatch.

ANCHOR-BACKFILL means legacy code exists but the new per-row evidence and cross-architecture contract is not complete. PARTIAL means the working subset is explicit. INVENTORIED is source knowledge only. Runtime-selected kernels remain unverified until an nsys trace on the declared workload identifies the actual family.

Current order-0 kernel leaf: finalized 179a0fc proves the FP4 tactic family already matches and the merged BA topology is structurally closed. The packed-decode spike is CLOSED on EQUIVALENCE (KERNEL-GDN-PACKED-DECODE → DONE, owner e47b4d6): correctness-green at f344dec, structurally accepted from 7ff713e/24cea4f, the c16 slot defect fixed and proven at c172336, and W1D3's G3 resolved over eight sealed component roots + the 8-pair locked c16 A/B (−0.205% ± 0.30, <1σ) + the 24-window trace attribution (packed GPU-cheaper, no attributable packed-side cost). There is no STABLE regression on any axis; the packed path stays the default (VT_GDN_PACKED_DECODE=0 rollback), with no complete-pass marker and no speed credit claimed. qkvz (KERNEL-GEMM-BF16 W2A) is now implemented (2026-07-15) and GATING on the orchestrator's DGX gates; the exact grid stays authorized (fresh vLLM denominators; explicit --mamba-ssm-cache-dtype float32; cite 702f481). The authorized exact-grid rerun has since landed: the current binding is 9ecd9d0 114/124 (benchmark_binding refers here, superseding 3f256ab's 55/124, retained immutable); see roadmap_v1.md and engine-matrix.md SERVE-GATE-ONLINE.

The C1 implementation spike is accepted at dropin-kernel-abi.md. Its additive BACKEND-ABI-VT W0 spine is CPU-green and GATING; no row below is promoted by that test-only probe. Until its CUDA/cross-build debts close, every new vt CUDA kernel still notes its upstream csrc/dependency counterpart in a source comment. Afterward, each family migrates incrementally to the common raw pointer/shape/stride/scalar/layout/ workspace/stream adapter and completes its own correctness, trace, every-axis performance, and memory checkpoint before the next migration stacks.

Current KERNEL-GDN-PACKED-DECODE extension checkpoint (2026-07-25): the historical Hv=32-launcher-rejection text embedded in that completed 27B row is superseded for dense 4B only. CLAIM-LOCAL-BF16-H32-AOT adds the current raw-packed/FP32-state gdn_decode_h32 specialization; 35B remains inert through the dense-only model selector. Full GDN is 66/66 (4,242/4,242), the flag test is 10/10, the real graph/direct/eager 4B gate is 3/3 (1,672/1,672), and sm_120 plus sm_121a manifest drift checks pass. Same-binary AOT versus rollback is +4.5906% total/output throughput, and the final graph-node trace proves the specialization executes. The extension is complete; its parent 4B loader row remains speed-pending.

Current KERNEL-GDN-PACKED-DECODE / KERNEL-GDN-AOT-BF16 per-arch AOT checkpoint (2026-07-28, CLAIM-TRITON-AOT-PER-ARCH, base 308c312a): the vendored Triton-AOT GDN fast-path cubins (decode + delta_h + chunk_o + kkt/tril/ wu, h48+h32) previously existed for sm_121a ONLY, so the MEASURED codegen-win GDN decode (Triton REG:205/0-spill vs hand-CUDA REG:255+STACK:48 spills) was runtime-real on GB10 alone; on sm_80/86/89/90a/100a a -DVLLM_CPP_TRITON=ON build had no vendored tree and every committed cross-family arch build shipped -DVLLM_CPP_TRITON=OFF (GDN decode on the SPILLING hand kernel). The full GDN AOT set is now regenerated + vendored for all five arches via the sanctioned -DVLLM_CPP_TRITON_REGEN=ON -DVLLM_CPP_TRITON_VENDORED_ARCH=sm_XX pipeline on dgx (Triton 3.6.0 / ptxas 12.8, cross-compiled — no target board needed): 57 artifacts + MANIFEST per arch matching the sm_121a fileset, cuobjdump proves real per-target SASS (sm=80/86/89/90/100, decode REG 209–217 / 0-spill, all under the hand-CUDA spill floor). The builder path (-DVLLM_CPP_TRITON=ON, no REGEN) selects + integrity-verifies each new tree (vendored … (no Python)). Honest signal: DERIVED+BUILD-VERIFIED — no non-sm_121 board runs a GDN model here, so the decode-parity CLAIM on those arches is build-verified, not runtime-measured. The build-time cubin selection (sm_${VLLM_CPP_CUDA_ARCHITECTURES}) is already additive; sm_121a is byte-untouched (SACRED 27B/35B gate structurally unchanged), check-triton-aot-drift.sh rc=0 across all six trees. Full recipe + per-arch table: per-arch Triton AOT.

Current KERNEL-ATTN-FA2 ratio-4 checkpoint (2026-07-25): the existing generic head-dim-256 split-KV launcher now admits the Qwen3.5-4B exact Hq/Hkv=16/4 topology behind VT_FA2_DECODE_4B (default ON, =0 rollback). The composed reference suite is 25/25 (454,474/454,474). Same-binary FA2 versus generic paged fallback is +1.6004% total/output throughput. Node-mode execution is trace-proven at 180.28 us/call, within 1.1% of the matched vLLM FA2 kernel at 178.40 us/call. Evidence: 2026-07-25 4B repair.

KERNEL-ATTN-FA2 d128 varlen decode GQA group-swap (ported 2026-08-06 row/KERNEL-FA2-GQA-SWAP; FLIPPED DEFAULT-ON 2026-08-06 row/KERNEL-FA2-GQA-SWAP-FLIP, VT_FA2_DECODE_GQA_SWAP, =0 opts out): vLLM's seqlenq_ngroups_swapped decode grid ported into LaunchDecodeVarlenFA2Bf16 (cuda_flash_attn_fa2.cu) so the Qwen3-dense d128 decode launches (batch, kv_heads) not (batch, hq) — the ngroups query heads pack into seqlen_q, KV read once per group, heuristic sees batch*kv_heads — halving the CTA count at batch>=2 (#47: ours over-waved 192 CTAs c2 / ~5 waves c8). Presented WITHOUT a materialized transpose via kv-major-group-minor strides, mirroring the shipped d256 LaunchDecodeFA2Bf16 swap (the vendored get_lse_tile/combine already honor the flag). Gated so =0 is byte-identical to the shipped plain-varlen reduction; ON is non-byte-exact only when num_splits>1 (split reduction order → near-tie, toward vLLM). FLIP CAMPAIGN on GB10 (CUDA 13.0, sm_121a, HEAD 1f446fd7): op RED-first test 280/280 assn GREEN (both ratios, batch 1/2/4/8, short+long; swap_launches==1; wrong-stride RED = 26528 violations), compute-sanitizer 0-err/0-leak; SACRED Qwen3-0.6B/4B greedy near-tie gate 16/16 both, TOKEN-IDENTICAL swap-ON vs plain (no token flip) at default, =1, and =0; #44 MXFP4 smoke token-exact at the new default. nsys ours c2 decode-flash grid (1,3,64)=b×q_heads → (1,5,16)=b×kv_heads, per-call 63.7→45.3us (default no-env reproduces the swap grid). Binding q3mxfp4 grid swap-ON total tok/s c1 0.990 / c2 0.922 / c4 0.930 / c8 0.942 (was 0.989/0.911/0.919/0.913, c2-c8 outside per-rep noise, c1 flat, NO regression), median TTFT at/above parity, peak GPU mem 2.614x LESS. Flipped per parity-enablers (improves, no regression, correctness holds). Still <1.0x tput/TPOT — residual grouped-Marlin +7-9% + ~0.7ms host/sched. Detail: state KERNEL-FA2-GQA-SWAP-FLIP.

ID	Item	Upstream	Our code	Tests/evidence	Spike/spec	State	Owner
`KERNEL-CPU-A76-Q8-DOT`	Cortex-A76 Q8_0 x Q8_0 SDOT, exact-order C++ and scheduled AAPCS64 variants for the Pi decode/prefill keep-quant GEMM	llama.cpp `ggml/src/ggml-cpu/quants.c:400` portable dot and `arch/arm/quants.c:1076-1160` Arm DotProd @ project pin `237ad9b96`; Qwen semantics remain vLLM `555967922`	ACLE SDOT + HWCAP/MIDR selector, AAPCS64 two-block leaf, Q8 dispatch, same-binary variants	direct exact variant/edge tests; QEMU 20/20, 150258 assertions; physical Pi exact checksums/tokens; assembly vs compiler SDOT +3.66% M1/T1, +5.08% M128/T1, +3.69% M128/T4, but −2.43% M1/T4. Same-file llama.cpp floor measured/NOT MET: vllm.cpp 0.461x prefill / 0.653x decode+E2E, 0.758x peak RSS (better), exact-prompt output equal; assembly evidence, competitor evidence	A76 Q8 dot spec	`GATING`	-
`KERNEL-ACCEL-PROVIDER-SELECT`	WHICH implementation of an op runs, when more than one exists on a device — the selection layer above every kernel family in this matrix. Distinct from `KERNEL-CUDA-DISPATCH-AOT`, which is about which ARCH a CUDA kernel is compiled/selected for; this is about which PROVIDER (ours, a vendor library, MLX, llama.cpp) serves the op at all	no single upstream file: this is the shape vLLM's runtime chain uses everywhere — flashinfer tactic registries, cuBLASLt/CUTLASS per-call heuristics, and torch's backend selection — rather than compile-time pinning	`vt::OpProvider` op_provider.h + op_provider.cpp; the flat `[OpId][DeviceType]` `void` table it replaces is gone from ops.cpp with the ~70 op wrappers untouched. Providers registered today: `vt-native` (every backend kernel in the tree, priority 0, unconditional — behaviour preserved exactly), `mlx` (priority 100, Metal `kMatmul`/`kMatmulBT`, build-gated `VLLM_CPP_MLX`, metal_mlx_provider.mm), and — NEW 2026-07-23 (`CLAIM-BACKEND-SEAM-S5-1`, work row `S5`) — `vt-cpu-ref` (priority −1000, the portable reference tier)*: the CPU kernel installed LAZILY as a negative-priority fallback on a UNIFIED-MEMORY device's first `GetOp` miss, mirroring `custom_op.py:138 forward_native`, so a partial backend runs an op it lacks natively instead of throwing. Native always wins (priority); gated on `Backend::UnifiedMemory()` (a discrete GPU never gets it — a CPU kernel on true device memory is corruption); observable via `GetReferenceTierHits()` + a one-time loud stderr line	test_op_provider.cpp 11 cases / 47 assertions — deterministic selection under REVERSED registration order, name tie-break, duplicate rejection, capability predicate, caps re-resolution, decline-and-fall-back, stats, runtime disable; test_metal_backend.cpp 9 cases / 108 assertions on the M4 with MLX ON, including MLX-vs-MSL-vs-CPU NMSE per op at real shapes and an end-to-end DECLINE; test_reference_tier.cpp (S5): discrete-device refusal + unified-device zero-native-kernel fallback correctness + native-wins + observability, hardware-free via a fake backend on `kXPU`. Linux CPU 156/156; dgx regression set ALL UNCHANGED — anchor `tests/vt/test_op_provider.cpp:64`	Metal/MLX reuse study §6; accelerator-seam-audit §10; drop-in kernel ABI (the complementary ARGUMENT half)	`ACTIVE` — mechanism landed and gated with THREE provider kinds (`vt-native`, `mlx`, `vt-cpu-ref`); the CUDA/CPU/Vulkan vendor provider rows it was designed for are not yet populated (so the row is deliberately left open)	`CLAIM-BACKEND-ACCEL-PROVIDER-1`
`KERNEL-CUDA-DISPATCH-AOT`	CUDA runtime, streams, graphs, per-arch dispatch, and generated/AOT artifact selection	platform selection `vllm/platforms/cuda.py:205-493`; gencode handling `CMakeLists.txt:201-220`; JIT packages `cmake/external_projects/triton_kernels.cmake:1-28`	cuda_backend.cu:20, CMakeLists.txt:37, TritonAOT.cmake:57	CUDA backend tests; the Triton GDN AOT tree is now vendored per-arch — `sm_121a` (RUNTIME-VERIFIED) plus `sm_80/86/89/90a/100a` (DERIVED+BUILD-VERIFIED, `CLAIM-TRITON-AOT-PER-ARCH`, 2026-07-28), each a real per-target cubin (`cuobjdump sm=80/86/89/90/100`); the CUDA fast-path arch fan-out (non-GDN) remains as evidenced in `backend-matrix.md`	inventory; per-arch Triton AOT	`ANCHOR-BACKFILL`	-
`KERNEL-FUSION-FRAMEWORK`	Portable automatic op-fusion framework (roadmap_v1 ORDER-1 extensibility cornerstone) — declare-once `constexpr FusedRecipe` catalog (backend-agnostic, above `vt::`) transcribing vLLM's finite fusion-pass set; realize-per-backend through the `vt::` op table (Tier-0 composite = CPU oracle inherited free; Tier-1 interpreter = one kernel/backend lights up every recipe). A new vLLM fusion PR = ONE declaration; a new GPU = one catalog realization; a new model = additive pattern bindings. NOT a perf lever (35B ceiling ~3.5%/step compute-bound); primary value = extensibility + mechanical upstream-sync + CPU/CUDA oracle-drift elimination	pass manager `vllm/compilation/passes/pass_manager.py:86,138-200`; the finite fusion-pass set `passes/fusion/{rms_quant_fusion.py:618,act_quant_fusion.py:283,qk_norm_rope_fusion.py:188,attn_quant_fusion.py:362}`; declare-once analog `vllm/model_executor/custom_op.py:103` (`forward_native`/`forward_cuda`/`enabled()` :138-207); combo_kernels OUT-OF-SCOPE `vllm/config/compilation.py:967`	Phase-0 skeleton LANDED: fused_recipe.h:72, recipes.h:35, kFusedChain dispatch, Tier-0/1 CUDA cuda_ops.cu:1062 + CPU cpu_ops.cpp:1427; catalog-to-be = the hand-fused ops (`kRmsNormQuantFp8`/`kSiluMulFp4Quant`/`kAttnQkNormRopeGate`/`kSigmoidGateFp4Quant`/`kRmsNormGatedQuantFp8`). 2026-07-19 (`CLAIM-FUSION-FRAMEWORK-W0`): W0 ADOPTED at ONE production site — `RunLayerPaged` post_attention_layernorm (35B MoE paged, `test_qwen36_paged_engine` hot path) now routes the plain add+residual+gemma-RMSNorm through `vt::FusedChain(kFusedAddRmsNorm)` behind `FusedChainAdoptEnabled()` (`VT_FUSED_CHAIN_ADOPT`, default ON / `=0` restores the exact prior `vt::RmsNorm(residual)` hand-call). BEHAVIOR-PRESERVING by construction (Tier-0 composite dispatches to the same `vt::RmsNorm(residual)` primitive; perf-neutral, NOT a speedup) — proves the declare-once/realize-per-backend seam end-to-end in production. The other 3 identical plain sites (`RunLayer`/`RunDenseLayer`/`RunDenseLayerPaged` post-attn + the input-layernorm plain path) are left on the hand-call; W1's generalized POD (multi-input operands + quant/silu/rope opcodes) is required only for the quant-fused sites (`kSiluMulFp4Quant`/`kRmsNormQuantFp8`/…), NOT for these plain add+rmsnorm sites — the current 3-opcode/4-role POD expresses them byte-identically. 2026-07-20 (`CLAIM-FUSION-FRAMEWORK-W1`): W1 GENERALIZED the POD (infrastructure only, perf-neutral, NO call site changed). `FOp` grew from `{kAdd,kMul,kRmsNorm}` to add `kSilu`/`kSigmoid` (Tier-1 elementwise) + fused-primitive opcodes mapping 1:1 to standalone ops so the composite is byte-exact by construction — `kSiluMul`→`vt::MoeSiluMul`, `kSigmoidGate`→`vt::SigmoidGateBf16`, `kRmsNormGated`→`vt::RmsNormGated`, `kRope`→`vt::RopeFromCache`, `kQuantFp8`→`vt::QuantFp8Static`, `kQuantFp4`→`vt::ScaledFp4Quant`, `kAttnQkNormRopeGate`→`vt::AttnQkNormRopeGate`. The fixed `{kIn,kResidual,kWeight,kOut}` 4-role model became an INDEXED `FOperandSlot operands[8]` table (`FStep` names input indices + `out`/`out2`); tensors bind positionally via `FusedBinding`, scalars via `FusedParams`; intermediates are caller-bound scratch (byte-exact to the unfused sequence). Tier-0 composite is now ONE device-agnostic walker in `ops.cpp` (per-opcode dispatch to the standalone op, self-dispatching per device — eliminates CPU/CUDA oracle drift); the per-backend `kFusedChain` op is now ONLY the Tier-1 interpreter over the canonical `[x,weight,residual,out]` shape. All five W2 target chains DECLARED as `constexpr` recipes (recipes.h, citing each vLLM pass). TIER (honest): composite realizes all five byte-exact; Tier-1 covers only the elementwise/rmsnorm subset (W0 `kFusedAddRmsNorm`) — the four quant/activation chains + attn preamble are COMPOSITE-ONLY. Backend-negotiation (§3b/§6): fp8 terminal `vt::QuantFp8Static` CUDA-only (fp8 chains' composite CUDA-only; CPU covers up to the bf16 norm); `kAttnQkNormRopeGate` a composite-only MACRO (per-head 3-D operands → dispatches the whole preamble to the single `vt::AttnQkNormRopeGate` op). ALL FIVE byte-exact expressible, none forced. Anchors: generalized POD fused_recipe.h, binding/params/general entry ops.h, device-agnostic composite `FusedChainComposite`, Tier-1 interpreters cuda_ops.cu+cpu_ops.cpp. 2026-07-20 (`CLAIM-FUSION-FRAMEWORK-W2`): W2 MIGRATED the bespoke hand-fusions to `vt::FusedChain(recipe)` — the framework now OWNS the fusion dispatch, perf-neutral by construction. New realization-binding infra (the "fast realization" tier the spike §3c reserved): `FusedRecipe.fast_op` (fused_recipe.h) carries the raw `OpId` of the EXISTING single-launch bespoke fused kernel that realizes the whole recipe (backend-AGNOSTIC — names WHICH abstract op, realized per-backend via the op table); a non-throwing `OpRegistered` probe + a per-recipe `DispatchFusedFast` adapter switch + convenience overloads in ops.cpp/ops.h. `FusedChain(recipe,…)` dispatches recipe→fast_op WHEN the backend registers that OpId (the SAME kernel the model called directly pre-migration → byte-identical AND perf-neutral; no extra kernel, no per-forward getenv/alloc), else falls through to the byte-exact Tier-0 composite (graceful degradation). The composite is the exposed oracle (`FusedChainComposite`) the fast realization is validated against. Bound recipes: `kRmsNormQuantFp8`→`OpId::kRmsNormQuantFp8`, `kRmsNormGatedQuantFp8`→`OpId::kRmsNormGatedQuantFp8`, `kSiluMulFp4Quant`→`OpId::kSiluMulFp4Quant`, `kSigmoidGateFp4Quant`→`OpId::kSigmoidGateFp4Quant` (recipes.h); `kAttnQkNormRopeGate` has NO fast_op (its composite MACRO already dispatches to the single `vt::AttnQkNormRopeGate` launch — perf-neutral trivially). SIX call sites migrated in qwen3_5.cpp behind `VT_FUSED_CHAIN_ADOPT` (`=0` restores the exact prior hand-calls, same binary): `kSiluMulFp4Quant` (MoE down-proj), `kSigmoidGateFp4Quant` (full-attn o-proj), `kRmsNormGatedQuantFp8` ×2 (GDN out-proj), `kRmsNormQuantFp8` (input-layernorm), `kAttnQkNormRopeGate` ×2 (attn preamble). qwen3_5.cpp net +61 lines (−21/+82): each migrated fused-op call became one `FusedChain` call (call-for-call neutral), but the mandated same-binary rollback `else` branch + explanatory comment at each of the 6 sites grows the raw count — the true shrinkage lands when the rollback branches and bespoke `OpId`s retire; the dispatch surface is already framework-OWNED. 2026-07-20 (`CLAIM-FUSION-FRAMEWORK-W3`): W3 MECHANICAL-UPSTREAM-SYNC PROOF — ported a NEW, previously-unported vLLM fusion-pass variant, `SiluMulFp8StaticQuantPattern` (`vllm/compilation/passes/fusion/act_quant_fusion.py:81` → `_C.silu_and_mul_quant`; registered unconditionally at `act_quant_fusion.py:296`), as ONE `constexpr FusedRecipe kSiluMulQuantFp8` (recipes.h) + its byte-exact test. The static-per-tensor-FP8 sibling of `kSiluMulFp4Quant`; its Tier-0 composite is realized ENTIRELY from EXISTING standalone `vt::` ops (`kSiluMul`→`vt::MoeSiluMul`, then `kQuantFp8`→`vt::QuantFp8Static`) — NO new primitive, NO kernel, NO composite-walker case, NO dispatch edit, NO model-site edit. `fast_op = kNoFastOp` (no bespoke silu·mul→static-fp8 fused OpId exists in-tree; a fast kernel is a separate later perf step). Additivity headline: adding a whole new fusion pattern touched exactly 2 shared files — `include/vt/recipes.h` (+1 declaration) + `tests/vt/test_ops_fused_chain.cpp` (+1 byte-exact test); `git diff --stat main` = those 2 files only. This IS the framework's primary-value proof (spike §4/§10): a new upstream fusion PR = one declaration. 2026-07-20 (`CLAIM-FUSION-FRAMEWORK-W4`): W4 — the EXECUTABLE BACKEND-ADDITIVITY PROOF, closing the W-series. Makes the claim executable: "a new backend registers `kFusedChain` (Tier-0 composite) ONCE and inherits the ENTIRE catalog correct, ZERO per-recipe work." Proof approach (spec §10 W4, PREFERRED option b): treat the EXISTING CPU backend AS the 'second backend' relative to CUDA — NO mock `DeviceType` (which would edit the core enum + every switch, ironically non-additive). New test `test_fused_chain_additivity.cpp` enumerates the WHOLE catalog (`kCatalog`, all 7 recipes) and in ONE generic loop asserts each runs correct on the CPU backend via the Tier-0 composite (`vt::FusedChainComposite`), BYTE-EXACT vs the standalone-op-sequence golden, over the CPU-expressible scope: 4 CPU-FULL end-to-end (`kFusedAddRmsNorm`/`kSiluMulFp4Quant`/`kSigmoidGateFp4Quant`/`kAttnQkNormRopeGate`), 3 CPU-PREFIX (`kRmsNormQuantFp8`/`kRmsNormGatedQuantFp8`/`kSiluMulQuantFp8` — static-fp8 terminal `vt::QuantFp8Static` CUDA-only per §3b/§6; the prefix is byte-exact and the FULL composite is asserted to THROW on CPU, documenting the backend-negotiated tail, not silently skipping). ADDITIVITY EVIDENCE (catalog GROWS ⇒ backend does NOT): catalog recipes.h 1→6→7 recipes (W0→W1→W3); the composite walker `FusedChainCompositeImpl` is ONE function with a per-OPCODE switch (12 `FOp::` cases, NOT per-recipe); the CPU/CUDA `kFusedChain` registration is ONE line each; `src/vt/cpu/cpu_ops.cpp` does NOT `#include recipes.h` (the backend never sees the catalog); W3's whole new recipe `kSiluMulQuantFp8` appears in ZERO backend TUs + ZERO in the `ops.cpp` dispatch — it runs purely through the shared walker (that is WHY W3 touched 2 files and no `src/vt/`). KERNEL-FUSION-FRAMEWORK ORDER-1 milestone DONE (W0 adopt + W1 POD + W2 migrate + W3 mechanical-sync + W4 backend-additivity). HONEST DEFERRED / future (named, non-blocking the ORDER-1 milestone): (a) the Tier-1 PERF interpreter covers only the elementwise/rmsnorm subset — the quant/activation chains realize composite-only (correct/byte-exact, but the single-pass fused perf kernel for the quant chains is future); (b) a REAL Metal/Vulkan realization of the catalog needs the M4 dev-box (HW-blocked); (c) per-recipe fast single-launch kernels for composite-only recipes are separate later perf steps. Perf is NOT the framework's value (§11: 35B ceiling ~3.5%/step, compute-bound). 2026-07-25 (`CLAIM-FUSION-CONSISTENCY-AUDIT`): CONSISTENCY AUDIT + ENFORCEMENT — read-only static audit of catalog USAGE across all 12 model forwards (fusion-consistency-audit.md). Verdict MOSTLY consistent: qwen3_5 family fully migrated; qwen3/qwen3_moe/deepseek_v2 adopted `kFusedAddRmsNorm{,Std}`; DRIFT — gemma/gemma2/gemma3/glm4/phi3 hand-call the residual `vt::RmsNorm(...,&res)` add+RMSNorm chain and never route through the catalog though the recipe already exists (one-line adoption each = follow-on `FUSION-DENSE-MIGRATE`). olmo2/granite/opt deliberately-not-fused (post-norm/LayerNorm). GDN glue + MoE-combine remain bespoke (spike §10 W2 remainder). NEW enforcement CI check `scripts/check-fusion-consistency.py` (+ mutation test, wired into the `agent-record` CI job) flags any model that hand-fuses add+RMSNorm without the catalog unless allowlisted — the 5 drift models on `fusion-consistency-allowlist.txt` keep the gate GREEN while blocking new silent bypasses; removing a stem is the gate closing. 2026-08-10 (`FUSION-DENSE-MIGRATE`, issue #299, spec fusion-dense-migrate.md): the MERGED-GEMM allowlist's no-blocker half is CLOSED. The five entries whose only stated reason was `pending FOLD-MIGRATE` — commandr, glm4, minicpm, minicpm3, phi3, all plain bf16 dense SwiGLU MLPs — now route gate/up through `layers::UnquantizedMlpGateUpMethod` (the direct arm, not the `MakeMlpGateUpMethod` factory: no loader for the three `Qwen3DenseMlpWeights` models ever populates `_fp4`, so the factory would add an untestable quantized branch), and their entries are GONE from `merged-gemm-consistency-allowlist.txt`: 11 → 6 entries, and the six survivors are exactly the Group-B blockers that need the SHARED LAYER extended (f32 activations + rank-1 biases, ggml block-quant weights vs the explicitly UNQUANTIZED arm, an up-front-device-staging vs `OwnedTensor`/`ResidentWeight` residency difference, a GeGLU arm on the SwiGLU-only grouped MoE op, an NVFP4-Marlin-resident arm for weights held as raw device pointers) — each its own row, none touched here. MECHANICAL NOTE the issue got wrong: a fully-folded TU has NO hand-call left, so it drops OUT of the detector's denominator rather than moving into its numerator — the checker reads `15 scanned / 6 routed / 11 allowlisted` → `10 scanned / 6 routed / 6 allowlisted`, NOT the `11 routed / 15` the issue predicted; drift is 0 either way and that is what the gate asserts. Shared-layer TUs (`linear.h`/`nvfp4.h`) and every `vt::` op are UNTOUCHED, so no already-routed model can move. BYTE-EXACT BY CONSTRUCTION: each replaced body was the seam's own `{ResidentWeight; MatmulBT[2I,H]; SiluAndMul}` with `M` spelled `T`, and every call site passes a `DBuf{T,H}` so `x.shape[0] == T` identically. The OP sequence is identical; the ALLOCATION sequence is NOT, and that is stated rather than hidden — `DBuf` is pooled and frees on destruction, so the `[T,2I]` gate_up block is now released when `Apply` returns and the `[T,H]` output may reuse it; no value can move (`vt::MatmulBT` writes every output element from a fresh f32 accumulator) and peak pool use is unchanged or lower, and every earlier fold has the same shape. CPU-gated (no GPU on this claim): new byte-exact case in `test_linear_method` for the DIRECTLY-constructed arm at BOTH the decode shape (M=1) and a prefill shape (M=4) vs the standalone sequence, RED-first PROVEN (mutating the seam's `SiluAndMul`→`GeluAndMul` fails 2 cases / 11 assertions; GREEN on byte-for-byte restore) → 6 cases / 76 assertions GREEN; two new mutation cases in `test_check_fusion_consistency.py` make BOTH regressions RED (re-allowlisting a folded stem, and reverting a fold) → 20/20. EXECUTED COVERAGE OF THE FOLDED TUs (added on review finding F4): NEW `test_dense_gate_up_seam_forward` drives the REAL forward of four of the five folded TUs (commandr/glm4/minicpm/phi3) over synthetic in-memory weights on CPU — no checkpoint, no GPU — pinning the gate/up split ANALYTICALLY (zeroing the UP half and zeroing `down_proj` both make the MLP contribute exactly zero, so they must be byte-identical, which holds only at a split of exactly `I`), the half order, and a vacuity guard → 4 cases / 1940 assertions GREEN. RED-first EXECUTED both ways: `I` → `I - 1` at all four call sites fails all four cases (phi3 alone fails only its own, so each TU is genuinely reached — before this file the mutation survived 176 CPU tests), and a SHAPE-PRESERVING `vt::SiluAndMul` CPU-kernel mutation (`silu up` → `silu + up`) fails all four on the split assertion while `test_linear_method` stays GREEN at 76/76 — a defect class the seam unit test structurally cannot see. Both restored byte-for-byte. HONEST/OWED: all five paged-engine SACRED gates are checkpoint-gated dgx-only and emit a loud SKIP on a CPU box, so the ORACLE token-exact confirmation is PENDING for the next GPU holder — the same shape A1 carried for Granite/StableLM — tracked by #337. (The row's original "no e2e evidence is possible on a CPU box" framing was too strong and is retired: ORACLE evidence needs the GPU, SELF-CONSISTENCY evidence does not.) `minicpm3` is the one folded TU the new harness does not drive — MLA attention plus the load-time `kv_b_proj` absorption belong to the DeepSeek-V2 synthetic harness.	test_ops_fused_chain.cpp (byte-exact composite==interp==golden; W1 extends it: each new recipe's Tier-0 composite == its standalone-op-sequence golden byte-exact, fp4+attn CPU+CUDA, fp8 CUDA-only. W2 adds the fast==composite==golden triple arm per recipe — CPU 228 assertions, CUDA-on-dgx 420, memcheck 0). W3 adds `RunSiluMulQuantFp8Cuda` (kSiluMulQuantFp8: FusedChain==composite==MoeSiluMul+QuantFp8Static golden, byte-exact, CUDA — the newly-ported pass's test), and the no-token-regression gate `test_qwen27_paged_engine` 235/235 + `test_qwen36_paged_engine` 315/315 unchanged (the new recipe is DECLARED, not wired into any model — the engine is untouched); W4 adds `test_fused_chain_additivity.cpp` (the whole-catalog additivity loop: every recipe byte-exact on the CPU 'second backend' via the Tier-0 composite, CPU-full end-to-end / fp8-prefix + negotiated-tail-throw; count guard `==7`) — 1 case / 17 assertions CPU green, sibling `test_ops_fused_chain` 228/228 unchanged, engine byte-identical so 27B 235/235 + 35B 315/315 structurally unchanged; to port: `tests/compile/passes/test_fusion.py:249-287` oracle discipline (byte-exact, stricter than upstream `assert_close`) — anchor `tests/vt/test_ops_fused_chain.cpp:124`	portable-fusion-framework.md	`ANCHOR-BACKFILL` (W-series ORDER-1 PROOF MILESTONE DONE: W0 adopt + W1 POD + W2 migrate + W3 mechanical-sync + W4 backend-additivity, all merged+gated; row stays SPIKE-open because the framework retains named future work — the Tier-1 quant-chain perf interpreter (composite-only today) + a real Metal/Vulkan catalog realization (M4 HW-blocked) + per-recipe fast kernels, §10 Wn)	`CLAIM-FUSION-FRAMEWORK-W4`
`KERNEL-GEMM-BF16`	BF16 dense GEMM, including torch-Linear/TN-equivalent layout and Qwen GDN merged input projections	Qwen mapper/packing `vllm/model_executor/models/qwen3_5.py:200-210,278-288`; merged linear `vllm/model_executor/layers/linear.py:580-808`; GDN calls `vllm/model_executor/layers/mamba/gdn/qwen_gdn_linear_attn.py:908-943` (qkvz projection + mixed/z split `:923-936`, construction `:481-496`); unquantized dispatch `vllm/model_executor/layers/utils.py:92-99,332-338`; runtime cuBLASLt heuristic	W1 one-owner BA + W2A one-owner QKVZ loader:165, merged/fallback qkvz dispatch forward:2372, eligibility seam internal:52, stride-aware consumers:930, unchanged cuBLASLt path:188, exact oracle generator:1, and explicit trace contracts:113	W1 BA closed through the packed-decode equivalence closure (`f344dec` 235/235 both arms; 35B/GGUF inert). W2A qkvz implemented 2026-07-15 (test-first): merged owner in exact [q,k,v,z] rows, ONE BF16 GEMM + strided mixed/z views on CUDA default, split rollback from the same owner (`VT_GDN_MERGED_QKVZ=0` / `VT_GDN_MERGED_PROJ=0`); CPU tier green — merged-view battery:2367, loader/CPU-exactness:456, eligibility:451, full CTest 107/107, tools 162/162, clean -Werror rebuild. DGX gates at `baea3ec`: default/2a-qkvz-rollback/35B-inertness arms PASS; the `VT_GDN_MERGED_PROJ=0` arm exposed a non-mode-aware gate-test expectation (engine correct — master-off deselects packed decode by the designed BA coupling) → fixed test-first via `detail::PackedGdnDecodeEnvSelected` (truth table:491, 16/16); 2b re-run, memcheck (first run PATH-only) and the 145→97 BF16 trace pending; `benchmark_binding=false`, no speed credit	merged GDN projections; packed decode	`ANCHOR-BACKFILL`	CLAIM-GDN-BA-ROUNDING-1
`KERNEL-GEMM-FP8`	FP8/INT8 scaled-mm plus static activation quant	C2x/C3x dispatch `CMakeLists.txt:737-863`; stable quant sources `:383-388`; vLLM cuBLASLt fp8 reuses an in-graph plan (nvjet_sm121_qqtst_* kernels) — no per-call heuristic	cuda_matmul_fp8_cutlass.cu:312, cuda_matmul.cu:345; Hopper `sm_90a` arch coverage (2026-07-28, `CLAIM-CUDA-SM90-C3X`): the Hopper C3x FP8 scaled-mm is a SEPARATE build-verify TU cuda_scaled_mm_c3x_sm90.cu (faithful 1:1 port of vLLM `cutlass_3x_gemm_sm90_fp8`, `ArchTag=Sm90`+`KernelTmaWarpSpecializedFP8FastAccum` → wgmma/TMA), gated by its own `scaledmm-c3x-sm90` FEATURE-TABLE cell (90a-only) — DERIVED+BUILD-VERIFIED, no H100/H200 board ran it; datacenter-Blackwell `sm_100a` arch coverage (2026-07-28, `CLAIM-CUDA-SM100-C3X`, DC3):* the sm100 C3x FP8 scaled-mm is a SEPARATE build-verify TU cuda_scaled_mm_c3x_sm100.cu (faithful 1:1 port of vLLM `cutlass_3x_gemm_sm100_fp8`, `ArchTag=Sm100`+`KernelScheduleAuto` → 5th-gen tcgen05 collective; `sm100_fp8_config_{default,M256,M64}`, 2SM `ClusterShape<_2,_2,_1>` default), gated by its own `scaledmm-c3x-sm100` FEATURE-TABLE cell (100a-only) — DERIVED+BUILD-VERIFIED (cuobjdump `sm_100a` cubin + `Sm100TmaUmma`/`SM100_MMA_F8F6F4`/`TMEM` symbols), no B200 board ran it (see backend-matrix `BACKEND-CUDA-SM100`); the sm_12x production body here (`ArchTag=Sm120`) is UNCHANGED (see backend-matrix `BACKEND-CUDA-SM090`); opt-in per-device plan cache fp8_plan_cache.h + GetOrBuildCachedFp8Plan/BuildFp8Plan (`VT_FP8_PLAN_CACHE=1` opt-in; DEFAULT OFF — bit-exact but measured production-NEUTRAL, premise not reproduced, see ledger 2026-07-18)	FP8 tests; byte-exact cached==fresh test_ops_fp8_cutlass.cpp#L387 + `VT_FP8_PLAN_CACHE=1` on-arm ctest; CPU key/flag test_fp8_plan_cache.cpp; 27B 235/235 + 35B 315/315 both flags. Merged-QKV FP8 sub-lever (`CLAIM-FP8-MERGED-QKV-1`, `VT_FP8_MERGED_QKV` opt-in): extend the fp4-only merged-QKV fusion to 35B FP8 — ONE fp8 GEMM over the N-concatenated Q/K/V operand + per-column dequant, replacing 3 separate per-shard GEMMs (10 attn layers, 30→10 GEMMs/step). fp8 is PER-TENSOR scaled so a single-alpha concat is incorrect; realized as concat RAW bytes + GEMM alpha=1 + resident per-column alpha vector via NEW `vt::MulColVecF32` (ops.h, cpu_ops.cpp, cuda_glue.cu); model glue `ResidentFp8Qkv`/`MergedFp8QkvD`/`MergedFp8QkvEligible` + `ProjectFullAttnQkv` branch (qwen3_5.cpp), resident fields (qwen3_5_weights.h); byte-exact CPU tests test_ops_glue.cpp. CPU gates GREEN (glue 10/10, fp8_cutlass 6/6, matmul 7/7, clean -Werror); DGX GREEN @ `e9ce593` (clean CUDA -Werror 0 warn): 35B 315/315 token-exact both arms + 27B 235/235 both arms (inert), merge proven to fire; in-situ TPOT A/B NEUTRAL (c1/c8 ~0%, c2/c4 −0.5%, all ≤0.9% within rep noise) ⇒ landed OPT-IN (`VT_FP8_MERGED_QKV` default OFF, token-exact but not measurably faster; the merged-QKV sub-lever is complete — the broad FP8 row stays `ANCHOR-BACKFILL` for its remaining scope). Spec fp8-merged-qkv-projection.md. Per-column alpha epilogue sub-lever (`PERF-FP8-ALPHA-FOLD`, #402 §3 "Lever B", `VT_FP8_ALPHA_VEC_EPILOGUE` opt-in, DEFAULT OFF): new `vt::MatmulFp8CublasLtAlphaVec` (ops.h, cuda_matmul.cu) applies the merged operand's per-output-column folded alpha INSIDE the cuBLASLt epilogue via `CUBLASLT_POINTER_MODE_ALPHA_DEVICE_VECTOR_BETA_ZERO` (our fp8 D is col-major [N,M], so cuBLASLt's output ROWS are our output COLUMNS and the resident `alpha_vec` already fits), instead of a second full-tensor `vt::MulColVecF32` pass measured at 122.99 ms/req over 48 calls = 43.6% of the 27B T=4096 prefill deficit. Algo-capability gated with a two-launch fallback; `scale_mode` keeps the plans from aliasing. CPU plumbing tests test_fp8_plan_cache.cpp 7/7 + 63 assertions; the bitwise CUDA case and both SACRED gates are OWED (implementer host had no nvcc/GPU). Spec perf-fp8-alpha-fold.md	inventory; fp8-merged-qkv	`ANCHOR-BACKFILL`	-
`KERNEL-GEMM-NVFP4-W4A4`	NVFP4 W4A4 dense quant, merged/fused projections, runtime bucketing, SM12 tactics, v0.25 persistent plan selection and model-owned alpha	SM10/11/12 FP4 families `CMakeLists.txt:940-1002`; CT alpha parameter `compressed_tensors_w4a4_nvfp4.py:95-141`; executed FlashInfer pass-through `kernels/linear/nvfp4/flashinfer.py:97-176`; fused selection `act_quant_fusion.py:36-40,128-181,283-300`; stable fused body `activation_nvfp4_quant_fusion_kernels.cu:30-163`, packed helpers `nvfp4_utils.cuh:25-36,118-329`, vector loads `cuda_vec_utils.cuh:123-175,264-288`; FlashInfer 0.6.13 device pointer `gemm_base.py:1307-1350`, `fp4_gemm_cutlass_sm120.cu:52-77,82-105,135-175`; v0.25 cache lifecycle sources retained	Existing W3-C/W3-F anchors remain. W3-I1 adds the default-off packed fused producer and dispatch, public zero-lifecycle contract, and candidate/fallback/graph/alignment tests. Datacenter-Blackwell arch coverage (2026-07-28, `CLAIM-CUDA-SM100-NVFP4`): the sm_100a tcgen05 block-scaled NVFP4 GEMM is a SEPARATE build-verify TU cuda_matmul_nvfp4_sm100.cu (faithful 1:1 port of vLLM `Fp4GemmSm100`, `ArchTag=Sm100`+`KernelScheduleAuto`), gated by its own `cutlass-nvfp4-sm100` FEATURE-TABLE cell (100a-only) — DERIVED+BUILD-VERIFIED, no B200 board ran it; the sm_12x production body here is UNCHANGED (see backend-matrix `BACKEND-CUDA-SM100`). The trace-only controller, driver, batch-keyed validator, fail-closed c2 finalizer, and finalizer tests support exact c2 without changing production builds	Clean W3-I1 remains default-off after 27/40 timing + 3/8 memory. Finalized `179a0fc` proves all 12 local ranges and 1,522 steady oracle windows resolve the same 128 Stream-K 128x64x256 + 80 static-persistent 128x32x256 split. Diagnostic local/oracle FP4 medians are 52.508720 / 52.734326 ms, so FP4 GEMM is not the positive c2 residual. Status `9e0143fa…7b57` is `complete-diagnostic`; `3f256ab` stays 55/124 and no new speed credit exists. 2026-07-17 (`CLAIM-FP4-QUANT-FAST-1`): two NUMERICS-NEUTRAL decode-glue vectorization sub-levers landed OPT-IN (default OFF; the underlying quant kernels stay implemented, no row-state change). Bit-identical vectorized-load+store fast kernels behind `VT_FP4_QUANT_FAST` (ScaledFp4QuantFastKernel) + `VT_SILU_FP4_FAST` (SiluAndMulFp4QuantFastKernel): each thread does ONE 16-byte `uint4` load (vs 16 scalar) + ONE 64-bit packed store (vs eight 1-byte), memory-access-pattern change ONLY (exact `CastToFp4NibbleDev`/`F32ToFp8Dev`/`fmaxf`-amax/bf16-SiLU math unchanged), grounded 1:1 in vLLM `nvfp4_quant_kernels.cu:56-80,98` @ `e24d1b24`; the numerics-changing hw-cvt + `__hmax2` reduction (vLLM's other ~0.4× of the edge) stays out of scope (that is the non-bit-identical `VT_FP4_FUSED_VEC` native kernel). BIT-IDENTITY PROVEN: byte-exact new-vs-old nibbles+scales, 60/60 adversarial parity asserts + full suite 24/24 (26,976), flag header + CPU test 20/20. Isolated nsys per-launch (swizzled bf16): ScaledFp4Quant K=5120 1.12-1.18× / K=17408 1.44-1.62×, SiluAndMul I=17408 1.14× (c2) / 1.38× (c16-c32) — PARTIAL vs the ≥1.3× flip bar (clears the larger shapes, misses the dominant K=5120 / c2; swizzled small-M is padding-thread-dominated), so BOTH stay default OFF; the orchestrator owns the combined in-situ A/B. Engine token gate both-flags-ON PASSED: 27B 235/235 (16/16 token-exact vs vLLM) + 35B 315/315. `benchmark_binding=false`, binding 52/124. 2026-07-18 (`CLAIM-CONV-UPDATE-FAST-1`): both flags flipped DEFAULT OFF→ON per the parity-enabler policy (bit-identical ⇒ never-slower + token-safe; under the strict ≥1.0 gate every fraction counts). Predicate parse in fp4_quant_fast.h inverted to default-ON `=0`-rollback; CPU flag test RED→GREEN 20/20; CUDA byte-exact test scalar baseline arm → `=0`, re-verified byte-exact 25/25 (26,976). No kernel-body change. Full default set (both flags default ON) 27B 235/235 + 35B 315/315; combined `=0` rollback 235/235 + 315/315. Binding grid re-measures the combined in-situ effect. 2026-07-19 (`CLAIM-SIGMOID-GATE-FOLD-1`): full-attention sigmoid-gate → o_proj activation-quant fusion landed OPT-IN (`VT_FUSE_SIGMOID_QUANT=1`, default OFF). NEW `vt::SigmoidGateFp4Quant` op (cuda `SigmoidGateFp4QuantKernel`, cpu composite, ops decl) folds `attnsigmoid(gate)` into the o_proj NVFP4 activation quant — one kernel, no bf16 `gated` intermediate — mirroring vLLM Inductor `triton_poi_fused_mul_scaled_fp4_quant_sigmoid` and the `SiluMulFp4Quant` precedent; model dispatch `SigmoidGateOProjD` (qwen3_5.cpp) fires only on the 27B true-W4A4 o_proj (35B W4A16-Marlin/fp8 reads bf16 acts ⇒ inert, keeps `SigmoidGateBf16`). BIT-IDENTICAL to `SigmoidGateBf16`+`ScaledFp4Quant`: byte-exact op test 14/14* (CPU f32/bf16 + CUDA, 3 shapes). DGX-GREEN (production flags, clean CUDA `-Werror` 0 warn): 27B 235/235 both arms + 35B 315/315 (inert). In-situ 27B TTFT A/B (input-1024, 3 reps): c1 −0.15% / c2 −0.03% — NEUTRAL within rep noise (o_proj is a small slice of 27B prefill) ⇒ OPT-IN. Spec glue-fusion — anchor `tests/vt/test_ops_nvfp4_fp4.cpp:113`	small-M spike; W3-E spike; W3-C spike; W3-F device-alpha spike; W3-G spike; W3-H normal-producer spike; W3-I fused-producer spike	`ANCHOR-BACKFILL`	CLAIM-SERVE-GATE-1
`KERNEL-GEMM-MARLIN-W4A16`	FP4 W4A16 Marlin dense/grouped GEMM and repack	Marlin generation/targets `CMakeLists.txt:548-679,1168-1274`; capability floor `marlin_utils_fp4.py:29-35`; DENSE marlin_gemm `csrc/libtorch_stable/quantization/marlin/marlin.cu:326-541,545`	cuda_marlin_repack.cu:129, cuda_moe_marlin.cu:156; `KERNEL-MARLIN-DENSE-PORT` (gated OFF `VT_MARLIN_DENSE`): vLLM's OWN dense W4A16 GEMM cuda_marlin_dense.cu, lifted dispatcher marlin_mm_dense.cu, op ops.h `kMarlinDenseGemm`, routing dense_nvfp4_gemm.h — byte-preserving E=1 replacement for the single-expert MoE route (no par-regroup ULP)	dense tests, MoE + DENSE tests; 35B gate; dense-port unit RED-first battery (dense-vs-ref-vs-grouped, M=1..8, NVFP4+MXFP4, wrong-stride RED); 3 new dense `.cu` compile-clean dgx sm_121a; GPU exec PENDING	inventory	`ANCHOR-BACKFILL`	`KERNEL-MARLIN-DENSE-PORT`
`KERNEL-GEMM-INT-LOWBIT`	AWQ/GPTQ/integer Marlin, Machete, and AllSpark	`CMakeLists.txt:468-731`; Marlin types/capability `marlin_utils.py:43-149`	-	-	inventory	`INVENTORIED`	-
`KERNEL-GEMM-W4A8`	CUTLASS W4A8 dense/MoE	`CMakeLists.txt:1004-1035`; upstream test `tests/kernels/quantization/test_cutlass_w4a8_moe.py`	-	-	inventory	`INVENTORIED`	-
`KERNEL-GEMM-QUTLASS-MX`	QuTLASS NVFP4/MXFP4 block-scaled GEMM	`cmake/external_projects/qutlass.cmake:56-147`; upstream tests `test_mxfp4_qutlass.py`, `test_nvfp4_qutlass.py`	-	-	inventory	`INVENTORIED`	-
`KERNEL-GEMM-DEEPGEMM`	DeepGEMM dense, batched, and MoE JIT paths	`cmake/external_projects/deepgemm.cmake:61-105`; upstream `tests/kernels/moe/test_deepgemm.py`	-	-	inventory	`INVENTORIED`	-
`KERNEL-GEMM-CPU-ELEM`	The elementwise (f32/f16/bf16) CPU GEMM — `kMatmul` (`[K,N]` weight) and `kMatmulBT` (`[N,K]` weight) for every non-block-quantized operand. Distinct from `QUANT-GGUF-CIQ-GEMM`, which owns the block-quantized `kMatmulBTQuant`: this row is what runs for EVERY safetensors CPU path and for every f16/f32 tensor of a mixed GGUF (60 % of the bench file's weight bytes, including the tied `token_embd`/`lm_head`)	llama.cpp `ggml/src/ggml-cpu/vec.cpp:139` (`ggml_vec_dot_bf16`, incl. the `:172` bf16 shift-widen), `:264` (`ggml_vec_dot_f16`), `vec.h:72-73`, `simd-mappings.h` (`GGML_F16_VEC_LOAD`), `ggml-cpu.c:1155-1243` (chunk worker) @ `237ad9b96`	cpu_matmul_elem.{h,cpp} — per-dtype micro-kernel tier table `{bt, nk, btm, mr}` with a runtime ISA probe: portable 16-accumulator tier (always built), AArch64 NEON (`vshll_n_u16`/`vcvt_f32_f16` widen, `vtrnq_f32` 4x4 transpose, MR=4), x86-64 SSE2 + probed F16C (MR=2); the rewritten chunk worker keeps the original loop verbatim as `MatmulOneChunkRef` and widens the activation row to f32 once per 16-row tile	test_ops_matmul_elem.cpp — 5 cases / 654 assertions, the gate is `memcmp` BYTE-IDENTITY against an independent in-test scalar reference (not NMSE): 3 activation x 3 weight x 2 output dtypes x 2 orientations x 14 shapes covering ragged K and N, row-strided activations, thread counts 1/2/4/8, and the ENTIRE 16-bit f16/bf16 widening domain incl. every inf/NaN pattern; green on x86-64 AND dgx aarch64 and under all three `VT_CPU_MATMUL_TIER` settings. DGX: clean CUDA `-Werror` build 0 warnings, SACRED set ALL UNCHANGED each STANDALONE (27B 235/235, 35B 315/315, Coder 6/6, Qwen3-dense 16/16, OPT 6/6, DeepSeek-V2 8/8, gguf 28/28), goldens md5 `2965ef5772b556d3f3f86fedf4221b2f` identical before/after. Binding A/B (idle dgx aarch64, one flock, SAME binary, 3 reps): prefill 21.67 -> 73.97 t/s (3.41x), decode 7.649 -> 23.79 t/s (3.11x), RSS unchanged; vs llama.cpp decode 3.21x -> 1.03x behind (parity within 3.1 %), prefill 8.00x -> 2.34x behind. Tokens byte-identical across BEFORE/AFTER/`VT_CPU_REF=1` (md5 `d235db12f2cd304007530286a1755c95`) — anchor `tests/vt/test_ops_matmul_elem.cpp:124`	elementwise CPU GEMM leaf	`ACTIVE`	`CLAIM-KERNEL-CPU-ELEM-GEMM-1`
`KERNEL-QUANT-CIQ-IQUANT`	Keep-quant `vec_dot` for the ~2-3-bit codebook encodings IQ2_XXS, IQ3_XXS and Q2_K — the DeepSeek-V4 single-Spark GGUF memory enabler (W8). Extends `QUANT-GGUF-COMPUTE`'s six-type `kMatmulBTQuant` so the `unsloth/DeepSeek-V4-Flash-GGUF UD-IQ2_XXS` routed experts (IQ2_XXS gate/up, IQ3_XXS down; Q2_K = the UD-Q2_K_XL sibling) stay COMPRESSED and dot directly against Q8_K activations — never expand to bf16 (~91 GiB vs the ~316 GiB bf16 OOM). CPU tier only, exactly like the six existing k-quants: `kMatmulBTQuant` is registered on `kCPU` alone (there is NO CUDA keep-quant vec_dot for ANY k-quant), so on GB10 it runs on the 20 ARM cores against the unified pool	llama.cpp `ggml/src/ggml-cpu/quants.c:514` (`ggml_vec_dot_q2_K_q8_K_generic`), `:855` (`ggml_vec_dot_iq2_xxs_q8_K_generic`), `:999` (`ggml_vec_dot_iq3_xxs_q8_K_generic`) + `ggml-common.h:288-299,371-374,385-400` (block structs), `:499,:503,:550,:1007` (kmask/ksigns/iq2xxs/iq3xxs grids) + `ggml-quants.c:2503` (`dequantize_row_iq3_xxs`) @ `237ad9b96`	cpu_quant_dot.cpp (`VecDotQ2_KQ8_K`/`VecDotIQ2_XXSQ8_K`/`VecDotIQ3_XXSQ8_K`), cpu_quant_iq_tables.h (shared grids), cpu_quant_dequant.cpp (`DequantIQ3_XXS`), cpu_quant_traits.cpp (three Q8_K rows), dtype.{h,cpp} (`kIQ3_XXS`), cpu_quant_blocks.h, gguf_reader.cpp (id 18 sizing)	test_ops_quant_dot.cpp — the 3 new types added to `kWeightCases`, gated by the SAME machinery as the six k-quants: `vec_dot` vs an INDEPENDENT f64 dequant-then-dot (≤1e-5·L1), `MatmulBTQuant` NMSE ≤5e-4 vs dequant-f32, bit-exact across thread counts, ragged-K rejection. 19 cases / 130444 assertions GREEN (CPU Debug, `-Werror` clean on the new TUs); RED-first PROVEN: perturbing the IQ2_XXS `0.125` fold fails 2 cases / 18 assertions, revert restores 19/130444	gguf i-quant DSV4; DeepSeek-V4 §W8	`SPIKE`	`CLAIM-DEEPSEEK-V4-W8`
`KERNEL-QUANT-CIQ-GEMM-CUDA`	CUDA keep-quant GGUF k-quant GEMM (MMVQ-style) — the kCUDA provider for `kMatmulBTQuant` (the FIRST CUDA keep-quant GEMM for any GGUF k-quant). Runs the DeepSeek-V4 routed-expert / MLA GEMMs ON THE GPU with the weights kept COMPRESSED in the unified pool (no bf16 expansion, which would OOM the 119 GiB pool at ~316 GiB): quantize the activation tile to Q8_K on-GPU, integer-dot it against the compressed block scales/codebook (dequant-in-kernel), warp-per-output. Serves the whole Q8_K family (Q2_K/Q3_K/Q4_K/Q5_K/Q6_K/IQ2_XXS/IQ3_XXS — DeepSeek-V4's experts are IQ2_XXS/IQ3_XXS/Q2_K); legacy Q8_0-activation types (Q4_0/Q8_0) fall back to the CPU keep-quant kernel over the same unified tensors. Registering it flips `GgufQuantComputeAvailable` TRUE on `kCUDA` so the loader keeps blocks compressed on a CUDA device and `vt::MatmulBT`→`kMatmulBTQuant` dispatches to the GPU instead of the unified-memory CPU reference tier (experts move OFF the 20 ARM cores)	The landed CPU keep-quant ORACLE it mirrors: cpu_quant_gemm.cpp (`MatmulBTQuantKernel`), cpu_quant_dot.cpp (the 7 `VecDot*Q8_K`), cpu_quant_act.cpp (`QuantizeRowQ8_K`) — themselves 1:1 ports of llama.cpp @ `237ad9b96` `ggml-cpu/quants.c:514/:566/:645/:720/:800/:855/:999` + `ggml-quants.c:2696`, `ggml-cpu.c:1245-1443` (mul_mat wiring). NOT a copy of llama.cpp's CUDA `ggml-cuda/mmvq.cu`+`vecdotq.cuh` (that path uses a Q8_1 activation → would NOT reproduce our Q8_K-based CPU oracle) — the STRUCTURE (MMVQ warp-per-row, dequant-in-kernel) is ported, the Q8_K numerics are the CPU oracle's	cuda_quant_dot.cu (`QuantizeQ8KKernel` + the 7 `Dot*` per-super-block device fns + `QuantDotGemmKernel` warp-per-output + `MatmulBTQuantKernelCuda` + the `RegisterOp(kMatmulBTQuant, kCUDA)` registrar; per-stream grow-only Q8_K scratch RETIRED on growth for cudagraph safety via graph_safe_scratch.h), cuda_quant_iq_tables.cuh (`__constant__` iq2/iq3/sign grids auto-derived from the CPU header)	test_cuda_quant_dot.cpp — CUDA vs the CPU keep-quant oracle (NMSE ≤1e-6, proving the Q8_K quant + INTEGER dot are bit-identical; only the per-super-block float scale sum reassociates) AND vs an independent f64 dequant reference (NMSE ≤5e-4, the test-backend-ops band) over the 7 types × M{1,4,32,512} × N{1,7,16}. RUNTIME-VERIFIED on the DGX GB10 (sm_121): 2/2 cases · 92401/92401 assertions GREEN, compute-sanitizer memcheck 0 errors; RED-first PROVEN (perturbing the IQ2_XXS `0.125` fold → 24 assertions fail, nmse 0.0064; revert restores 92401/92401, source md5 `9890f7e1…`). DeepSeek-V4 experts-on-GPU tok/s is the FOLLOW-ON (benchmark lane `a35f6be0`); the provider-registered + GPU-dispatch mechanism is gated here	CUDA keep-quant GEMM; DeepSeek-V4 §W8; gguf i-quant DSV4	`ACTIVE`	`CLAIM-CUDA-KEEPQUANT-GEMM`
`KERNEL-EW-NORM-ACT`	Activations, RMS/layer norm, embedding, casts, and glue	stable sources `CMakeLists.txt:383-397`; vLLM vectorized/CUB reference `csrc/libtorch_stable/layernorm_kernels.cu:54-173,251-360`; the decode ORACLE golden (enforce_eager, pip-vllm:0.24.0) runs the EAGER csrc `cub::BlockReduce<float,1024>` kernel — NOT generated Triton (that only appears under torch.compile, which the golden disables)	cuda_ops.cu:889, cuda_glue.cu:262, gated-fast RmsNormGatedRowFastKernel + TryLaunchRmsNormGatedFast + flag header; fused gated-norm→fp8-quant RmsNormGatedQuantFp8KernelCuda + CPU oracle + op	RMSNorm correctness contracts remain green. Finalized `179a0fc` has 177 calls on every local and steady oracle B=2 window; diagnostic medians are local 2.249728 ms versus oracle generated partitions 0.439491 ms, a +1.810237 ms residual. Finalizer status `9e0143fa…7b57` is `complete-diagnostic`; a whole-chain spike follows the larger BF16-GEMM launch residual, and cross-profiler timing is not speed credit. 2026-07-16 (`CLAIM-EW-NORM-QUANT-RECONCILE`): the refuted decode norm/quant FUSION lever reassigns its residual here as an EFFICIENCY headroom. Decode-shape RMSNorm microbench `~/work/vllm.cpp-ewnorm-spike` (sm_121a, flock): the shipped `RmsNormRowKernel` is 6-9 µs isolated at M=16 (H 2048/3072/4096), a single-pass shared-staged bf16x2-vectorized variant is 1.27-1.49× faster but reorders the f32 reduction ⇒ occasional 1-ULP (token-exactness hazard, per-arch like the attn preamble). Real recoverable headroom is a NON-bit-exact ≤1.5× ⇒ ~0.3-0.5 ms/step c16 ceiling on a ~168 ms TPOT (~0.3%), gated on the 27B-fp4 token gate + an in-situ c16 A/B (isolated-fast ≠ in-situ-fast; the reg-tile lever proved this). Spike first; do not speculatively implement. 2026-07-16 (`CLAIM-EW-NORM-ACT-1`): Phase-1 same-profiler adjudication CONFIRMS the lever, PORTED default-OFF. The 129 standalone decode RMSNorm launches are ALL at hidden `H=5120` (64 input_layernorm + 64 post_attention_layernorm + 1 final; q/k head norms are fused into the attn preamble; the reconcile microbench's H=2048-4096 was the WRONG shape). nsys pure-kernel BOTH sides, isolated, at M×5120 (evidence `dgx:~/work/vllm.cpp-ewnorm-phase1`): ours `RmsNormRowKernel` 8.44-8.53 µs/launch vs vLLM `triton_red_fused…rms_norm` 2.37-2.68 µs = 3.18-3.56× (honest Δ ≈0.77 ms/step at c2 AND c16; c2's is ~33% of the ~2.4 ms c2 gap). Cross-profiler confound removed (it was in-situ only: ours in-situ nsys 15.5 µs ≈1.84× contention-inflated over 8.46 µs isolated). Port `RmsNormRowFastKernel` = 1:1 vLLM `fused_add_rms_norm_kernel<bf16,8>` (1024-thread block, 16-byte `_f16Vec` loads, block reduce): isolated ~parity with vLLM (nsys 2.83 µs, 2.24-2.50× over V0), bf16-EXACT vs shipped at c2-c16 (2/163840 elements 1-ULP at c32). Reordered reduction ⇒ NOT bit-identical (token-exactness hazard, 35B fp8 ULP-sensitive) ⇒ shipped OFF, then the DGX proof PASSED and the default FLIPPED ON (2026-07-17): gate3 token gates ALL PASS both flags/models (27B 17/17+84/84, 35B 4/4+8/8); gate4 corrected-build c16 A/B fast +1.1% tput / −1.68/−1.90 ms meanTPOT on 2 clean pairs (legacy-r3 VOID interference anomaly; c2 pooled medians parity — lottery). 2026-07-17 (`CLAIM-ASYNC-SCHED-W3`): the fast-kernel DEFAULT was ROLLED BACK to OFF. The async-default-flip DGX gate ran the fuller `test_qwen27_paged_ENGINE` 16-token production stream (the `696a991` flip gate used only `paged_FORWARD` 17/17, which does not exercise it) and found fast-ON DIVERGES from the pip-vLLM oracle golden — 234/235 fast-ON vs 235/235 fast-OFF, async-independently (the reordered 1024-thread reduction flips a documented 27B whitespace/near-tie greedy argmax at token 7; vLLM's real oracle runs an Inductor-Triton rmsnorm, NOT the csrc kernel this port mirrors, so oracle bit-parity was never guaranteed). Token-exactness vs the oracle is sacrosanct ⇒ the shipped `RmsNormRowKernel` (oracle-exact) is the default again; `RmsNormRowFastKernel` is OPT-IN (`VT_RMSNORM_DECODE_FAST=1`). The fast-kernel default-ON perf lever REOPENS pending a numerical fix that matches the Inductor-Triton stream. Evidence `dgx:~/work/vllm.cpp-async-flip` (fast-OFF 235/235 both models × 3 async arms; fast-ON 234/235 3/3); closing record parity-ledger.md#L503. 2026-07-17 (`CLAIM-EW-NORM-ACT-2`): NUMERICS REWORK — real cub reduction; re-flipped default ON. The rollback premise was WRONG: the oracle golden is generated with `LLM(..., enforce_eager=True)` (`tools/parity/dump_qwen36.py:242`, pip-vllm:0.24.0) ⇒ the EAGER csrc op with `cub::BlockReduce<float,1024>`, NOT Inductor-Triton. `RmsNormRowFastKernel` reverted to the csrc per-element math (packed bf16 add, f32 sum_squares of bf16, 1024 threads) and swapped the hand warp-shuffle for the ACTUAL `cub::BlockReduce<float,1024>.Reduce(v, CubAddOp{}, blockDim.x)` (`#include <cub/cub.cuh>`), reproducing the oracle's exact reduction order (guard H≥1024 so csrc `block==1024`). DGX proof (`dgx:~/work/vllm.cpp-ewnorm-numerics`, corrected build): `test_qwen27_paged_ENGINE` 235/235 fast-ON (token 7 = 198, the tier that caught the regression) + `qwen36_paged_engine` 315/315 fast-ON; both rollback arms 235/235 + 315/315; paged_forward 84/84 + 8/8; CUDA parity 132/132; perf nsys pure-kernel 2.66 µs median vs shipped 8.66 µs (~3.2×), within vLLM's 2.37-2.68 µs. c16 in-situ A/B (w0+3 pairs) = NO WIN: fast −0.60% tput / +0.34 ms meanTPOT, 3/3 pairs (contradicts gate4's +1.1%; the shipped-arm's ~2% run-variation dominates ⇒ NULL within noise, as the spec anticipated). Default STAYS OFF (opt-in); NOT re-flipped — the token-exactness blocker is FIXED (the rework's achievement) but the flip acceptance (confirmed c16 win / no regression) is not met, so `RmsNormRowFastKernel` lands opt-in (`VT_RMSNORM_DECODE_FAST=1`, now token-safe + the true vLLM mirror); the default flip awaits an in-situ win (c2 target). 2026-07-17 (`CLAIM-SERVE-GATE-2`): the c2 in-situ win ARRIVED — DEFAULT FLIPPED ON. Phase-0 preflight of the authorized binding-grid rerun (one flock, a321d7c hard-verified CUTLASS+FA2 production build, binding c2 corpus, interleaved w0-discard + 3 pairs, house pooled per-request-median convention, 18 requests/arm): pooled-median TPOT fast 101.900 vs legacy 102.812 ms = −0.912 ms (−0.887%), paired −1.237/−1.211/−0.843 3/3 fast-faster; total throughput +1.446% (167.83 vs 165.43), 3/3 pairs fast-higher; no void signature. Evidence `dgx:~/work/vllm.cpp-online-gate/preflight-rmsnorm-c2-a321d7c…/`. Predicate now default-ON / `'0'`-rollback (house `gdn_packed_decode_triton.h` convention), flag test inverted RED→GREEN 10/10; the fast-ON engine token gates stand at `e68c518` (235/235 + 315/315). 2026-07-17 (`CLAIM-SERVE-GATE-2` Phase-0 CORRECTION): flip REVERTED — combination numerics. The campaign's 27B engine sanity gate on the FULL default set (async + GDN cubin + RMSNorm-fast) fails 233/235 at the documented token-7 near-tie, and the combined stream EQUALS the fixture's `want_emu` (pip-vLLM EAGER) golden 16/16 — each kernel is 235/235 ALONE (e68c518 fast-ON/cubin-off; a321d7c cubin-ON/fast-off) but the PAIR's combined ≤1-ulp roundings cross a tie vLLM itself decides differently between graphed and eager modes; the pair was never gate-tested together. Default back to OPT-IN (`VT_RMSNORM_DECODE_FAST=1`); LESSON: flip batteries must run the FULL prospective default set. Re-flip candidate: the `CLAIM-EW-NORM-ACT-2` campaign's Triton-faithful RMSNorm variant + cubin (production-numerics-consistent pair, untested together). The binding grid runs from the reverted SHA (async + cubin defaults). Closing record parity-ledger.md#L504; spec rework addendum. Prior fast-ON evidence `dgx:~/work/vllm.cpp-ewnorm-act-src`; ledger#L500. 2026-07-17 (`CLAIM-EW-NORM-ACT-3`): BIT-SAFETY rework — bit-identical output; DEFAULT FLIPPED ON. Resolves the `a875397` revert cause by CONSTRUCTION: `RmsNormRowFastKernel` now reproduces the shipped `RmsNormRowKernel`'s float op sequence EXACTLY — residual add `__float2bfloat16(f32(x)+f32(res))` (== shipped `ResRound`, was `__hadd2`), variance summed in shipped's exact kBlock=256 strided-partial + shared-tree ORDER (1024-thread vectorized Pass 1 stages per-element f32 squares to shared, then the 256-thread tree reduces them; was `cub::BlockReduce<float,1024>` = a different order), `inv=1.0f/sqrtf` (was `rsqrtf`); only the element-independent normalize pass is vectorized. Output is therefore BIT-IDENTICAL (0-ulp) to shipped ⇒ `fast+cubin ≡ shipped+cubin ≡ 198`. DGX (`dgx:~/work/vllm.cpp-ewnorm-bitsafe`, clean `-Werror` build, CUTLASS+FA2 hard-verified, one flock): `test_cuda_ops` decode-fast fast==shipped 0-ulp BIT-EXACT (assertion tightened from ≤1-ulp; 132/132), full 432/432; production default (unset = fast+cubin+async ON) `test_qwen27_paged_engine` 235/235 (token 6 = 198) + `test_qwen36_paged_engine` 315/315, both `=0` rollback arms 235/235 + 315/315; flag test inverted RED→GREEN default-ON 10/10. Perf win SURVIVES bit-identity: isolated 3.55 µs vs shipped 8.58 µs (2.41×), in-situ 27B engine-forward RmsNorm median 4.38 vs 16.13 µs (3.68×), total 48.3 vs 55.4 ms — strictly less GPU work, identical bits. The revert's sole cause (combination tokens) is fixed while the accepted c2 in-situ win (+1.446%) is retained; default flipped ON per the parity-enabler policy, binding grid re-measures. Closing record parity-ledger.md#L509; spec bit-safety addendum. Port anchors: RmsNormRowFastKernel + TryLaunchRmsNormDecodeFast, flag header, CPU flag test, CUDA parity test. 2026-07-17 (`CLAIM-EW-NORM-GATED-1`): GDN gated-RMSNorm decode-fast — BIT-IDENTICAL (0-ulp) port, DEFAULT ON. The c2/c8 attribution named the gated norm a batch-independent kernel-glue lever (+0.40 ms/step vs vLLM's fused gated norm, "RMSNorm-gated 0.403/fused/+0.40"). The shipped `RmsNormGatedRowKernel` (`cuda_gdn.cu:851`) is the same slow pattern the plain RMSNorm decode-fast already fixed: `kBlock=256` threads for a `Dv=128`-element row (upper half of every block idle) + `x` reloaded in the normalize pass. New `RmsNormGatedRowFastKernel` (behind `VT_RMSNORM_GATED_FAST`, default ON / `=0` rollback) uses `kGatedFastBlock=128` threads (one per element, no idle half) with `x` register-cached, and is BIT-IDENTICAL to shipped by construction: variance in shipped's exact per-element-square + shared-tree ORDER (for `d==128` each thread owns one element so the 128-thread tree == shipped's 256-thread tree; shipped's extra `s=128` step only adds the provably-zero `partials[128..255]`), `1.0f/sqrtf` (NOT `rsqrtf`), same `((xinv)w)act` multiply order + same silu/sigmoid act + same `__float2bfloat16` store + same padded-gate addressing. Unlike the plain RMSNorm (block-starved), the gated norm launches `THv` (~512-768 at c16) blocks so the win is idle-thread + redundant-load elimination, not thread-count-to-hide-latency. DGX (`dgx:~/work/vllm.cpp-ewnorm-gated`, clean `-Werror` CUDA build 0 warnings, CUTLASS sm120a NVFP4 + FA2 sm_121a hard-verified, one flock): Gate 2 `test_ops_gdn` gated decode-fast `fast==shipped` 0-ulp BIT-EXACT 140/140 (contiguous rank-2 + padded rank-3 strided gate, silu/sigmoid, c1-c16); full GDN 50/50 (2483/2483). Gate 3 isolated nsys pure-kernel avg fast/shipped: rows=768 (27B c16) 3.29/6.70 µs = 2.04×, rows=96 (c2) 1.55/2.14 = 1.38×, rows=48 (c1) 1.61/2.11 = 1.31× (all ≥1.3×; 4-shape aggregate 1.87× avg / 2.01× median). Gate 4 full production default set (async + GDN cubin + RMSNorm-fast + gated-fast all ON) `test_qwen27_paged_engine` 235/235 (token 6 = 198) + `test_qwen36_paged_engine` 315/315; both `VT_RMSNORM_GATED_FAST=0` rollback arms 235/235 + 315/315. Default flipped ON per the parity-enabler policy; binding grid re-measures. Port anchors: RmsNormGatedRowFastKernel + TryLaunchRmsNormGatedFast, flag header, CPU flag test, CUDA 0-ulp parity test; spec; ledger 2026-07-17 gated row. 2026-07-19 (`CLAIM-EW-NORM-GATED-2`): extended the fast gated-norm to the 35B MoE f32 path. WHY the 35B was on the SLOW kernel: `GdnOutDType` (`qwen3_5.cpp:2608`) is f32 for the MoE 35B (`num_experts=256`) vs bf16 for the dense 27B, so its core/z/weight are f32 and the former bf16-only launch guard fell it to the slow `RmsNormGatedRowKernel<float,bf16>` (the 77.6 ms / 3.3% of 35B prefill, task #57/#58). `RmsNormGatedRowFastKernel` is now templated `<Tin,Tout>` over the SAME `Load`/`Store` helpers as shipped `RmsNormGatedRowKernel<Tin,Tout>`, so every dispatched `(Tin,Tout)` is byte-for-byte identical to shipped; `TryLaunchRmsNormGatedFast` dispatches bf16→bf16 (27B), f32→bf16 (35B GlueFuse), f32→f32, bf16→f32 under `d==128` + shared-in-dtype∈{f32,bf16} + out∈{f32,bf16}; covers prefill AND decode; flag unchanged (still default ON). DGX (`dgx:~/work/vllm.cpp-gated-f32`, clean `-Werror` 0 warnings, one flock): bit-exact `test_ops_gdn` 260/260 0-ulp (adds f32→bf16, f32→f32 combos), full 53/53 (3201/3201); compute-sanitizer memcheck 0 errors; token-exact 35B 315/315 fast-ON + 315/315 rollback, 27B 235/235 + 235/235; isolated nsys 35B prefill `RmsNormGatedRowFastKernel<float,bf16>` 168.7 µs/call vs slow `RmsNormGatedRowKernel<float,bf16>` 261.8 µs = 1.55× (−2.79 ms/1024-tok prefill); in-situ TTFT 35B c1 −4.6% median / c2 within noise, 27B c1 −2.4% / c2 −3.5% (3/3, regression check). DEFAULT stays ON; ledger 2026-07-19 gated-f32 row.	gated-fast port; decode-fast port + flip verdict; reconciliation; ledger 2026-07-17 row; inventory	`DONE`	`5a53fb5`
`KERNEL-EW-NORM-QUANT`	Fused norm/add/activation plus FP8/FP4 quant	vLLM sources `CMakeLists.txt:394-397`; generated v0.25 trace families `triton_red_fused_fused_add_rms_norm_scaled_fp4_quant`; installed fusion pass `compilation/passes/fusion/rms_quant_fusion.py:97-98,634-674`; pass guard `compilation/pass_manager.py:162-167`; FlashInfer Add+RMSNorm+FP4 `flashinfer-ref/cute_dsl/add_rmsnorm_fp4quant.py:16-30,91-182` and RMSNorm+FP4 `rmsnorm_fp4quant.py:16-24,78-87`	FP8 fusion cuda_ops.cu:891; current FP4 model path remains separate norm and quant launches; historical `76e9047` shared-staging experiment is not present	FP8 fusion test; exact `3f256ab` trace records 127,040 long-named kernels, but the dumped body stores BF16 after add+RMSNorm and the wrapper separately invokes `scaled_fp4_quant.out`/`cvt_fp16_to_fp4`, matching ours. Oracle `fuse_norm_quant` is false. Generated graph/subgraph SHA `d58f81b8…9401` / `466e359a…9dd8`; status/vLLM-kernel SHA `9762c1e6…1d0c6` / `e4e916d1…565`. 2026-07-16 RECONFIRMED (`CLAIM-EW-NORM-QUANT-RECONCILE`): the fresh correct-state c16 trace `~/work/vllm.cpp-gdn-stateio-trace/20260716` shows COUNT PARITY — vLLM runs `cvt_fp16_to_fp4` at 144/win == ours' 144 `ScaledFp4Quant`/step, both 129 rmsnorm/step — so both arms run SEPARATE add+RMSNorm→bf16 then a separate FP4 quant; the `…scaled_fp4_quant…` substring in the `triton_red_` names is the Inductor graph-region label, not a fused quant. The 2026-07-16 `SUMMARY.json` "vLLM fuses add+rmsnorm+fp4quant" note is the misleading-name trap again and is CORRECTED. Residual delta (rmsnorm 391 vs 2006 µs) is cross-profiler-confounded (nsys node vs torch CUPTI; the 2026-07-14 rescan called the +1.81 ms residual a cross-profiler artifact) — the isolated ours RmsNorm is only 6-9 µs (microbench), not 15.5 µs. FUSION lever CLOSED (nothing to mirror); efficiency headroom (non-bit-exact ≤1.5×, ~0.3-0.5 ms c16 ceiling) reassigned to `KERNEL-EW-NORM-ACT`*	reconciliation; inventory; no FP4 spike is promoted from a misleading trace name	`PARTIAL`	-
`KERNEL-ROPE-QKNORM`	RoPE/MRoPE and fused QK-norm/RoPE/KV insertion	`CMakeLists.txt:391-406`; upstream core tests `test_fused_qk_norm_rope.py`, `test_mrope.py`	cuda_ops.cu:900	RoPE tests, CUDA RoPE	inventory	`PARTIAL`	-
`KERNEL-KV-CACHE`	KV reshape, insert, copy, and cache maintenance	`CMakeLists.txt:402-404`; upstream `tests/kernels/test_cache_kernels.py`; MLA cache write `csrc/libtorch_stable/cache_kernels.cu:401-442` (`concat_and_cache_mla_kernel`) + host wrapper `:842-905`, bound from `_custom_ops.py:2532` — VERIFIED to be vLLM's OWN csrc kernel, not a dependency's (no flashinfer/cutlass variant in the dense-bf16 path); upstream test `tests/kernels/attention/test_cache.py::test_concat_and_cache_mla`	cuda_cache.cu:97; `vt::ConcatAndCacheMla` (MLA campaign W3) CPU ref cpu_cache.cpp:74 + CUDA cuda_cache.cu:97, validation ops.cpp	cache tests; MLA cache write test_ops_mla_cache.cpp — CPU-vs-CUDA EXACT (a pure copy has no rounding) at the real 576-wide V2-Lite geometry, plus stride-driven cache-view + split-projection-source cases; fp8_ds_mla refused loudly	inventory	`ANCHOR-BACKFILL`	-
`KERNEL-KV-QUANT`	Quantized KV cache, including FP8 and NVFP4	NVFP4 KV sources `CMakeLists.txt:954-960,985-992`; Triton KV dtype dispatch `vllm/v1/attention/backends/triton_attn.py:250-383`	current paged path admits F32/BF16 only cuda_paged_attn.cu:2416	-	inventory	`INVENTORIED`	-
`KERNEL-ATTN-PAGED`	Dense reference plus paged prefill/decode fallback. HETEROGENEOUS per-layer head_dim (Gemma-4 G1b, 2026-07-28): `ReshapeAndCache`/`PagedAttention` already take `head_size` per CALL, so the runner's per-layer KV head_dim (`KVCacheConfig::per_layer_attn_specs`) drives them at Gemma-4's sliding 256 / global 512 with ZERO kernel change — no new kernel, the per-call head_size seam is what makes it additive	backend registry `vllm/v1/attention/backends/registry.py:34-120` and CUDA selector `vllm/platforms/cuda.py:360-493`	cuda_ops.cu:908, cuda_paged_attn.cu:2310	paged attention tests; heterogeneous per-layer head_dim exercised e2e by Gemma-4 E4B STRICT 32/32 test_gemma4_paged_engine.cpp	inventory	`ANCHOR-BACKFILL`	-
`KERNEL-ATTN-FA2`	FlashAttention-2 varlen prefill and GQA split-KV decode	exact dependency pin/build `cmake/external_projects/vllm_flash_attn.cmake:1-46`; v0.25 dispatch `flash_attn.py:674-721,935-977`; FA2 pure-decode swap/heuristic/dispatch `flash_api.cpp:262-327,587-714,754-779` at `2c839c33`	Prefill plus exact ratio-6 swap/heuristic/scratch adapter cuda_flash_attn_fa2.cu:429, queue cleanup cuda_backend.cu:58, operator dispatch cuda_paged_attn.cu:2501, cast-free model gate qwen3_5.cpp:2692, build CMakeLists.txt:559	Existing prefill plus ported tests tests:1354. Immutable `ae9e8ff` passes clean sm_121a build, 20/20 + 454,323 CUDA assertions, zero-error/zero-leak memcheck, both 27B arms and 35B correctness. Paired traces switch exactly between 240 main+combine / 0 old and 0 combine / 240 old, retain 3,536 FP4 GEMMs/producers and 64 plans, and show no capture allocation/free/sync or graph D2H. The completed c2/c16 component covers 12 legs/612 requests and reaches 1.017668×/1.006548× mean total throughput but strict-fails 35/40 timing + 5/8 memory. W3-G earns no speed credit and no exact grid follows. 35B FA2-PREFILL DEFAULT-ON (2026-07-18, `CLAIM-35B-FA2-FLIP-1`): `FuseAttnPreambleOn` (qwen3_5.cpp:1204) flipped default-ON all arches so the 35B ratio-8 layers take FA2 prefill (kernel `fa2_prefill` admits any ratio at head_dim 256, cuda_paged_attn.cu:2494); full default-set gate 35B 315/315 + 27B 235/235 + memcheck 0-err (`dgx:/tmp/fa2gates_u.log`, `/tmp/fa2_memcheck2.log`); the bf16-round-before-RoPE tighten was op-bit-identical but flipped the 27B tok6 near-tie in combination (RMSNorm-saga) ⇒ preamble ships untightened. 35B FA2-DECODE DEFAULT-ON (2026-07-19, `CLAIM-35B-FA2-DECODE-1`): the ratio-6-only split-KV DECODE was extended to the 35B ratio-8 (Hq/Hkv=16/2) hd-256 full-attn layers via new env `VT_FA2_DECODE_35B` (default ON, independent of the 27B `VT_FA2_DECODE`); the vendored `LaunchDecodeFA2Bf16` body was already generic in the GQA ratio, only the three gates (cuda_paged_attn.cu:2519, cuda_flash_attn_fa2.cu:446, qwen3_5.cpp:3282) were widened. nsys `--cuda-graph-trace=node` proves a clean 1:1 decode-kernel swap: OFF = `PagedAttentionDecodeGqaKernel<...(int)8...>` ×300 (grid (1,2,1)=2 blocks at single-req decode, no combine); ON = `flash_fwd_splitkv_kernel` ×300 + `flash_fwd_splitkv_combine_kernel` ×300 (split axis GridZ up to 16 fills the machine), old kernel absent. FULL default-set gate: 35B 315/315 + 27B 235/235 token-exact, operator 21/454,358 (adds ratio-8 parity ladder), memcheck 0 illegal-access; in-situ 35B A/B (input-1024, `VT_FA2_DECODE_35B=1` vs `=0`, 4 interleaved pairs) c1 TPOT −10.5%/+10.4% tput, c8 −3.2%/+2.8% tput, TTFT neutral (`dgx:~/work/vllm.cpp-35b-fa2-decode/{gpu_series,gates_engine,nsys_*}.log`)	W3-G split-KV decode spike	`ANCHOR-BACKFILL`	`CLAIM-SERVE-GATE-1` DENSE hd-64 NON-SPLIT arm added 2026-08-12 (`CLAIM-MM-SPEED-AUDIO-ENC-FA2`, spec S17, #432): `flash_fwd_hdim64_bf16_sm80.cu` instantiates upstream's `run_mha_fwd_hdim64` via the plain batch `run_mha_fwd_` entry (both head_dim 64 and the non-split entry are firsts here - every prior caller is paged, forcing split-KV), driven by `LaunchDenseFA2Bf16` for the b=1 dense non-causal Whisper encoder layout (null `cu_seqlens_q` selects the batch geometry in `BlockInfo`). Exposed as the additive op `vt::AttentionDenseFa2`. 115.8x on the encoder attention kernel (19,278 -> 166.5 us/layer, nsys 32 inst both arms), encoder forward 731.7 -> 133.0 ms (5.50x). OPT-IN, not default (`VT_WHISPER_ENC_FA2=1`): the oracle teacher-force PASSES the ratified 0.5-nat band (worst 0.125, over-band 0) but with 3 divergent positions vs the default's 0 - adoption is a developer decision. The MECHANISM is a HYPOTHESIS, not a fact (corrected 2026-08-12, review of #439): the earlier claim that FA-2's bf16 P-operand conversion (`flash_fwd_kernel.h:347`) causes the flips was REFUTED by mutation M4 - forcing exactly that rounding into the shipping scalar kernel left the token md5 UNCHANGED. Five differences remain candidates (QK^T reassociation; `exp2f` softmax `softmax.h:86,118`; `exp2f` online rescale `softmax.h:157`; bf16 P; PV reassociation) and none is isolated. Unit tests (added 2026-08-12, F4/F5): `tests/vt/test_ops_attention_dense_fa2.cpp` - fast-path agreement inside the bf16 envelope, reference-free FULL-key-range coverage (kills a `seqlen_k`-truncating filler), fall-through totality for causal / GQA / hd!=64 / f32, and the `VT_FA2_DENSE=0` A/B arm. `LaunchDenseFA2Bf16` now takes `bool causal` and THROWS rather than silently answering the non-causal question. Vendored template UNCHANGED => the 128/192/256 split-KV kernels are byte-identical.
`KERNEL-ATTN-FA3-FA4`	FlashAttention-3/4	FA3/FA4 build `setup.py:1113-1124`; runtime version selection `vllm/v1/attention/backends/fa_utils.py:132-250`	-	-	inventory	`INVENTORIED`	-
`KERNEL-ATTN-FLASHINFER-TRTLLM`	FlashInfer paged attention, XQA, and TensorRT-LLM generation kernels	availability/dispatch `vllm/utils/flashinfer.py:206-342,373-511`; SM90 decode and SM100+ prefill/decode gates	-	-	inventory	`INVENTORIED`	-
`KERNEL-ATTN-TRITON-FLEX-HPC`	Triton, FlexAttention, TurboQuant, and HPC attention families	registry `registry.py:34-120`; Triton contract `triton_attn.py:250-383`; HPC gate `hpc_attn.py:235-247`	-	-	inventory	`INVENTORIED`	-
`KERNEL-ATTN-MLA-SPARSE`	MLA and sparse attention: CUTLASS, FlashMLA, FlashInfer, FA, Triton, MSA W6: the MLA attention BLOCK + WEIGHT ABSORPTION — the layer that COMPOSES W3+W4+W5 — mla_attention.h + mla_attention.cpp <- `mla.py:119-181` (`MultiHeadLatentAttentionWrapper.forward`) over `mla_attention.py:553-620` (the cache-update-BEFORE-attention order), `:624-874` (`forward_impl`: the dispatch + the absorbed decode) and `:2344-2425` (`forward_mha`); `AbsorbKvBProjBf16` <- `:875-962 process_weights_after_loading` (split `:892-900`, permutes `:959-962`); `MakeMlaUpProjectFn` <- `:2141-2170` (the `kv_b_proj` callback W5 left open); `BuildDeepseekRopeCosSinCache` <- `deepseek_scaling_rope.py:76-118` over `rotary_embedding/common.py:34-70`; `MlaAttentionScale` <- `deepseek_v2.py:995,1067-1075` (the mscale^2 correction, kept SEPARATE from the rope's own rotation mscale). Absorption is a LOAD-TIME weight transform plus TWO batched GEMMs, not a fused kernel, so it needed only two new general primitives: `vt::BatchedMatmul` <- `torch.bmm` at `mla_attention.py:789` (q-side W_UK fold) and `:1034` (`_v_up_proj`), CUDA impl = cuBLASLt STRIDED-BATCHED cuda_matmul.cu (the cuBLASLt form of the cuBLAS `gemmStridedBatchedEx` torch.bmm resolves to; the only upstream alternatives are ROCm-only aiter fp8/fp4 bmm branches) + CPU ref cpu_ops.cpp, stride-driven because BOTH call sites pass `.transpose(0,1)` views; and `vt::ConcatMlaNopeRope` <- `ConcatMLAQKernel` (`csrc/libtorch_stable/concat_mla_q.cuh`) + wrapper `cache_kernels.cu:1555-1600`, GENERALIZED to arbitrary nope/rope widths and a head-BROADCAST rope operand so one op also serves `_concat_k_nope_k_pe` (`:2063-2092`) — CUDA cuda_mla_attn.cu, CPU cpu_ops.cpp. Two ADDITIVE relaxations of existing ops, integer-identical for contiguous tensors: `vt::RopeFromCache` stride-driven on q/k (DeepSeek rotates the TRAILING 64-dim slice and its `k_pe` is a column block of the fused kv_a projection) and `vt::MatmulBT` accepting a row-strided ACTIVATION (`kv_b_proj` applied to a 512-column slice of the 576-wide workspace, `:2160`)	CUDA priority `vllm/platforms/cuda.py:84-176` (`_get_backend_priorities`, both branches); MLA classes `vllm/v1/attention/backends/mla/.py`; MLA prefill selector `mla/prefill/selector.py:47-76`; capability filter `vllm/v1/attention/backend.py:307-360`; CUTLASS build `CMakeLists.txt:1037-1061` W6 test_mla_attention_block.cpp 10/10 cases / 2,372,644 assertions* and test_ops_mla_absorb.cpp 9/9 / 1,644,807 assertions on dgx sm_121 — ports of `tests/kernels/test_concat_mla_q.py` (BOTH arms incl. the NON-CONTIGUOUS transposed-nope case, compared bit-exactly since a concat is a pure copy), the MLA-geometry sweep of `tests/v1/attention/test_mla_backends.py`, and the two-pass-oracle discipline of `tests/kernels/attention/test_mla_decode_cpu.py`. THE ABSORBED-vs-UNABSORBED EQUIVALENCE IS PROVEN NUMERICALLY, THREE WAYS: an INDEPENDENT double-precision block oracle computing the attention BOTH ways agrees to < 1e-11 (the identity itself, at both query branches); our absorbed decode reproduces the UNABSORBED oracle to < 2e-4 (f32); and the SAME batch driven once through the ABSORBED MQA decode kernel and once through the UNABSORBED materialized-MHA prefill path agrees to < 3e-4 (CPU f32) / < 4e-2 (CUDA bf16) — two code paths sharing nothing but the weights. Real geometry throughout (V2-Lite 512/128/64/128/16-head, plus V3's 7168 / 128-head / `q_lora_rank=1536` for the lora branch, which has NO e2e coverage and says so). Decode-only / prefill-only-no-context / chunked-prefill-with-context / MIXED (decode packed FIRST) all gated; NaN-poisoned outputs; run-to-run BIT-exact; CUDA cases proven to EXECUTE (124,941 + 290,835 assertions when run alone). memcheck 0 errors, racecheck 0 hazards, synccheck 0 errors (the last requires `--num-cuda-barriers 65536`: the default table OVERFLOWS on a binary driving this many kernel families and the tool then emits a bogus `unspecified launch failure`). Clean CUDA build 0 warn/0 err; regression set UNCHANGED (27B 235/235, 35B 315/315, Coder 138/138, Qwen3-dense 664/664, OPT 36/36). No speed number — W9 owns tuning	SELECTION (W2) + the DECODE KERNEL (W4) + the PREFILL PATH and CHUNKED-CONTEXT LOOP (W5). What is still absent is the MLA attention BLOCK and MODEL (W6/W7). Priority TABLE cuda_attn_priority.h:49 (both branches, one row per upstream arch arm) + lookup cuda_attn_priority.h:86; `is_mla()`/`is_sparse()` filter registry.cpp:63; `TritonMLABackend` NAME + 3-D `get_kv_cache_shape` backend.h:519, backend.cpp:83, registration backend.cpp:108 — W4: `vt::MlaDecodeAttention` — `OpId::kMlaDecodeAttention` + args/validation ops.h, ops.cpp; CPU single-pass REFERENCE cpu_mla_attn.cpp (numerics from `csrc/cpu/mla_decode.cpp`); CUDA two-stage split-KV cuda_mla_attn.cu — `MlaDecodeStage1` <- `_fwd_grouped_kernel_stage1` (`triton_decode_attention.py:278-458`, IS_MLA `v = tl.trans(k)` branch `:424-431`), `MlaDecodeStage2` <- `_fwd_kernel_stage2` (`:575-639`), `ComputeNumKvSplits` <- `_compute_num_kv_splits` (`triton_mla.py:40-47`), split workspace via the house grow-only per-stream scratch (upstream's `_reserve_attn_logits_workspace` `:57-78`). Deterministic by construction: fixed ASCENDING split merge, NO atomicAdd. `TritonMLABackend::get_impl_cls()` now returns a real `TritonMLAImpl` backend.h, backend.cpp; PREFILL remains W5 and `TritonMLAImpl::forward` refuses a prefill-shaped batch by name. W5: `vt::MlaPrefillAttention` + `vt::GatherMlaCache` + `vt::MergeAttnStates` + the chunked-context driver — `vt::MlaPrefillAttention` cuda_mla_prefill.cu / CPU ref cpu_mla_prefill.cpp <- `mla/prefill/flash_attn.py:153-248` `FlashAttnPrefillBackend` (the ONLY MLA prefill backend reachable on sm_121 per `mla/prefill/selector.py:66-76`, and it HARD-RAISES with no fallback at `:191-194`), running over the vendored FA-2 through the NEW launcher entry `LaunchMlaPrefillFA2Bf16` cuda_flash_attn_fa2.cu plus two new explicit instantiations of the UNCHANGED generic template (`flash_fwd_split_hdim192_bf16{,_causal}_sm80.cu`). V is ZERO-PADDED 128->192 and the output sliced back, exactly as upstream's `requires_v_padding` path does (`flash_attn.py:88-99,164-168,196-197`) — which is WHY the asymmetric QK 192 / V 128 pair needs no asymmetric kernel. `vt::GatherMlaCache` <- `csrc/libtorch_stable/cache_kernels.cu:992-1064`; `vt::MergeAttnStates` <- `csrc/libtorch_stable/attention/merge_attn_states.cu:18-192` (BOTH `-inf` edge cases ported verbatim). The workspace-bounded loop is mla_chunked_context.h <- `mla_attention.py:1422-1451,1667-1745,2094-2199,2344-2425`. The paged launcher `LaunchPrefillFA2Bf16` that every non-MLA prefill calls is textually UNTOUCHED (211 insertions / 0 deletions in that TU; 2 new vendored files)	test_attn_backend_registry.cpp:146 (GB10 MLA list), :203 (`use_mla=true` -> `TRITON_MLA`, matching the W0 oracle observation), :230 (the DSA seam, proven both directions with a stand-in sparse backend) — ports of `test_attention_backends_selection.py` (MLA cases), `test_mla_prefill_selector.py`, `test_mla_prefill_registry.py`; W4 test_ops_mla_attn.cpp — port of `tests/kernels/attention/test_mla_decode_cpu.py` (`ref_mla` as a TWO-PASS oracle, its bs=4/mean_seq_len=256/h_q=16/d=576/dv=512/block=16 parametrization, BOTH varlen arms, and its NaN-padding out-of-bounds detector) plus the `test_mla_backends.py` shape sweep: ragged, multi-block, single-block/single-token, EVERY num_kv_splits in {1,2,3,4,5,8,16,17,64,300,512} (incl. splits > seq_len, the empty-split path both stages must skip), 128-head DeepSeek-V3 geometry, head counts 1/3/17 that do not fill a BLOCK_H tile, a 288/256 block-32 non-V2-Lite geometry, bf16 + f32, and run-to-run BIT-exactness over 5 runs. Gated on 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 warn/0 err; regression set UNCHANGED (27B 235/235, 35B 315/315, Coder 6/6, Qwen3-dense 16/16, OPT 6/6). NO speed number yet — decode perf is W9. W5 test_ops_mla_prefill.cpp 4/4 cases / 2,377,052 assertions and test_ops_mla_chunked_context.cpp 5/5 / 306,037 assertions on dgx sm_121 — ports of `tests/v1/attention/test_mla_backends.py` and `tests/v1/attention/test_mla_prefill_quant_output.py` (its fp8 arms NOT ported: they need device-capability family 100, unreachable on sm_121 — recorded, not dropped). REAL V2-Lite prefill geometry (QK 192 / V 128 / latent 576, block 16, mscale^2 scale) against an INDEPENDENT double-precision TWO-PASS oracle, plus — for the chunked loop — a SINGLE-SHOT whole-sequence oracle that never chunks: exact / +1 / -1 chunk boundaries, a request with NO context, a chunk in which a request contributes ZERO keys, ragged multi-chunk, 128-head V3, single-token queries, ADVERSARIAL reverse-interleaved block tables, NaN-poisoned outputs, run-to-run BIT-exact over 5 runs. memcheck 0 errors, racecheck 0 hazards, synccheck 0 errors on both binaries; clean CUDA build 0 warn/0 err; regression set UNCHANGED (27B 235/235, 35B 315/315, Coder 138/138, Qwen3-dense 664/664, OPT 36/36). Prefill perf is W9	MLA campaign spike	`PARTIAL`	`CLAIM-MLA-DEEPSEEK`
`KERNEL-ATTN-DFLASH-BLOCK`	DFlash in-block attention — the project's FIRST non-causal / bidirectional attention primitive (SPEC-DFLASH D2, DF-DRAFT-MODEL). Per-request uniform (1+k) query block attends within its own block: FULL-attention layers BIDIRECTIONAL (`causal=false`, no mask), SWA layers causal-within-window. f32 online softmax, GQA broadcast. A SEPARATE `vt::` op from the causal `kAttention`/`kPagedAttention` so every other model stays byte-identical	`vllm/model_executor/models/qwen3_dflash.py:86-146` (`_resolve_layer_attention`: full layers default non-causal, SWA causal) + `:149-263` (`DFlashQwen3Attention`); flashinfer non-causal path (vllm#48167 Blackwell non-causal attn, in-pin)	`OpId::kDFlashBlockAttention` + `DFlashBlockAttentionArgs` + decl ops.h:1713 + wrapper/validation ops.cpp:2069; CPU REFERENCE `DFlashBlockAttentionKernel` cpu_ops.cpp:1843 (three-pass block-local softmax, the authoritative impl); CUDA `DFlashBlockAttentionKernelCuda` cuda_ops.cu:1300 mirroring the causal `AttentionKernel` block-reduction recurrence with per-block bounds + the bidirectional/window mask; the draft model that consumes it qwen3_dflash.cpp:52	CPU GATE GREEN test_ops_dflash_block_attn.cpp:79 5 cases / 12 assertions — hand-checked non-causal (query 0 sees the future key), the RED causal-vs-non-causal separation (the mask is load-bearing), per-request cu_seqlens block isolation, SWA window bound, GQA; model forward test_qwen3_dflash_forward.cpp:116 5 cases / 95 assertions (RED full-layer-causal-flip); existing causal `test_ops_attention` 9/9·23 UNCHANGED. GPU GATE GREEN on dgx (2026-07-26, GB10 sm_121a): CUDA `-Werror=all-warnings` build clean (kernel compiles as-written, no change); CUDA==CPU parity test_ops_dflash_block_attn CUDA case 198412/198412 within the 1e-4 f32-softmax envelope over all 5 corners; `compute-sanitizer --tool memcheck` 0 errors; consumed by the draft-forward parity gate (test_qwen3_dflash_draft_parity, fc rel-L2 0.46% / hidden ≤1.3% vs the real vLLM draft). DONE 2026-07-27 with the DFlash block (`CLAIM-DFLASH-D14`): the D2 non-causal in-block primitive is the CPU/materialized reference the D12+ paged/warp kernels are gated against; closure ledger.	DFlash spec §1.3/§6 D2	`DONE`	`489a7544`
`KERNEL-ATTN-DFLASH-PAGED-BLOCK`	DFlash PAGED in-block attention — the CAPTURE-SAFE form of `KERNEL-ATTN-DFLASH-BLOCK` (SPEC-DFLASH D12 Part B, the CUDA-graph draft-attention primitive). The (1+k) block queries attend over `[PAGED context ; their own (1+k) block]`: the growing context enters as DATA (paged K/V cache `[pages,block_size,Hkv,D]` + per-request `seq_lens` + `block_table`, mirroring `PagedAttentionKernel`) instead of a variable-size materialized combined buffer, so the launch grid is STATIC over the fixed `Nq=(1+k)*num_reqs` rows and EVERY metadata input is a persistent DEVICE tensor read in place — NO `cudaMallocAsync`/`cudaMemcpyAsync` of a function-local host `cu_seqlens` (the [[cudagraph-capture-bakes-stack-addresses]] UAF class the eager `LaunchDFlashBlockAttention` had). Same f32 online softmax + D2 in-block mask over the COMBINED index; bit-identical to `DFlashBlockAttention` over the materialized `[context;block]` buffer	vLLM full CG `dflash/cudagraph.py` + `speculator.py:411-458` + `precompute_and_store_context_kv` (`qwen3_dflash.py:548-619`) @ `555967922`; paged read mirrors our `PagedAttentionKernel` cuda_paged_attn.cu:184	`OpId::kDFlashPagedBlockAttention` + `DFlashPagedBlockAttentionArgs` + decl ops.h + wrapper/validation ops.cpp; CPU REFERENCE `DFlashPagedBlockAttentionKernel` cpu_ops.cpp; CUDA `DFlashPagedBlockAttentionKernelCuda` cuda_ops.cu (static grid, persistent device metadata) + D14 WARP variant DFlashPagedBlockAttentionWarpKernel	GPU GATE GREEN on dgx (2026-07-27, GB10 sm_121a): CUDA `-Werror` clean (0 warnings); test_ops_dflash_paged_block_attn.cpp cross-checks CPU-paged == materialized `DFlashBlockAttention` across 6 corners (non-causal, causal-SWA, block isolation, GQA, multi-page, zero-context) + CUDA==CPU (f32+bf16) = 795648/795648 assertions; `compute-sanitizer --tool memcheck` 0 errors. D13 (2026-07-27, `CLAIM-DFLASH-D13`): WIRED INTO PRODUCTION — the single-request DFlash draft forward (`ForwardPagedBody`, `qwen3_dflash.cpp`) now runs the (1+k) block through this kernel reading a fixed-capacity paged `DflashDeviceKVStore`, and the whole draft step is captured into a per-request CUDA graph + replayed (the growing context enters only via the in-place `seq_lens`). Capture-correctness PROVEN: `test_qwen27_dflash_spec_decode` 27/27 with the graph BIT-IDENTICAL to eager (same tokens + acceptance 19/39/29/25); c1 throughput NEAR-PARITY with vLLM-DFlash-ON (ours 0.978×, ~2% below the tight 3-rep band; gap closed 0.917×→0.978× via the paged read, the CG is perf-neutral) — the kernel is landed + wired + gated; STAYS `ACTIVE` with the engine feature (the ~2% ≥vLLM residual is per-step compute for an nsys). D14 (2026-07-27, `CLAIM-DFLASH-D14`): the residual WAS this kernel → WARP-scoped variant added → SPEED GATE MET, `DONE`. An nsys (`--cuda-graph-trace=node`) attributed the D13 ~2% residual to THIS kernel: `DFlashPagedBlockAttentionKernel` = 242.9 ms = 1.8% of the graphed step's GPU time, median ~~460 us/call (grid `(nq=17,hq=32)` × kBlock=256 threads looping SERIALLY over C~~500-640 keys with a 256-wide shared-mem tree reduction + 2 `__syncthreads` PER key — the latency/sync storm the ViT tower fixed with `AttentionDenseFast`), vs vLLM's fused flash draft-attn ~0.15%. Added `DFlashPagedBlockAttentionWarpKernel` (cuda_ops.cu): ONE WARP per (block-query,head), `__shfl_xor` butterfly head_dim reduction, register accumulator, NO `__syncthreads`; SAME paged/block combined-index read + causal/SWA mask + GQA (copied verbatim from the block kernel), mirroring the shipped `AttentionWarpKernel`. Default ON; `VT_DFLASH_ATTN_BLOCK=1` keeps the bit-identical D12/D13 block kernel. Draft attn 242.9 → 77.9 ms (3.1×); our-ON c1 28.60 → 29.32 tok/s; FINAL 3-rep A/B our-ON 29.32 ≥ vLLM-ON 29.240 (non-overlapping bands, 1.003×) ⇒ ≥vLLM MET. Not bit-identical to the block kernel but same f32-online-softmax math within envelope; CUDA==CPU `test_ops_dflash_paged_block_attn` 795648/795648 (f32 1e-4/bf16 3e-2) + compute-sanitizer 0; e2e 27/27 graph==eager, acceptance 19/39/29/25 unchanged (1629 accepted identical warp-vs-block); SACRED 235/235 + MTP 9/9 inert; `-Werror` clean; closure ledger	DFlash spec §0 D12/D13/D14	`DONE`	`489a7544`
`KERNEL-DFLASH2-GROUPED-CONV`	DFlash2 grouped dynamic depthwise convolution — the project's FIRST dynamic (input-conditioned) convolution kernel (SPEC-DFLASH2 W2, #1314). `out[i,c] = sum_t (base[side,t,c] + delta[i,side,t,g(c)]) * x[i-t,c]`, with tap `t` contributing only where `(i mod block_size) >= t`, `g(c) = c / group_size`, and `block_size` the QUERY block `1 + k`. Three things separate it from the shipped `KERNEL-DEPTHWISE-CONV1D`: the kernel is DYNAMIC (a per-position `delta` projected from the sublayer input, added to a static per-channel `base`), it is GROUPED (one delta per group of channels, one base per channel), and its taps are ZEROED ACROSS THE BLOCK BOUNDARY rather than across the sequence — which is what lets a proposal position see the ones before it without another backbone pass. `base_kernel` dim 0 is the prepare/finish SIDE and not a tap; on the published 27B draft both axes are 2, so nothing but the port note and the shape assertion separates a correct load from a transposed one. Every intermediate rounds to the tensor dtype, mirroring upstream's bf16 chain, so the op is elementwise with NO reduction-order freedom and the CUDA arm is specified BIT-IDENTICAL to CPU rather than within an envelope	BEYOND-PIN — `vllm/model_executor/models/qwen3_dflash2.py` (`_grouped_conv`, `DFlashGroupedConv`, `DFlash2Qwen3DecoderLayer.forward`) @ vllm-project/vllm#52816 head `19c9351904df4c63042671bc67a866ca48dc7d6f`; the parity pin `555967922` does not carry the architecture and this row does NOT advance it	`OpId::kDFlashGroupedConv` + `DFlashGroupedConvArgs` + decl/wrapper `include/vt/ops.h::DFlashGroupedConv` and `src/vt/ops.cpp::DFlashGroupedConv`; CPU REFERENCE `src/vt/cpu/cpu_ops.cpp::DFlashGroupedConvKernel` (the authoritative impl); CUDA mirror `src/vt/cuda/cuda_ops.cu::DFlashGroupedConvKernelCuda` (one thread per (row, channel); `__fadd_rn`/`__fmul_rn` forbid the FMA contraction the CPU build pins off). Consumed by the draft through `src/vllm/model_executor/models/qwen3_dflash.cpp::DflashConvPrepare` and `src/vllm/model_executor/models/qwen3_dflash.cpp::DflashConvFinish`, called from all THREE layer bodies, with the uniform-block precondition in `src/vllm/model_executor/models/qwen3_dflash.cpp::CheckDflashConvBatch`; weights loaded by `src/vllm/model_executor/models/qwen3_dflash_weights.cpp::LoadQwen3DFlash`	CPU GATE GREEN 2026-08-19 (test_ops_dflash2_grouped_conv.cpp) 8 cases / 9930 assertions, `Status: SUCCESS!`, exit 0 (was 6 / 9410 on 2026-08-19, before the bf16 rounding cases below) — upstream's own sequential reference loop at block 5 (the `% block` arm), 8 and 16 (the two PUBLISHED checkpoints, `z-lab/Qwen3.8-27B-DFlash2` and `z-lab/Muse-Glimmer-30B-DFlash2`; upstream's parametrize covers 5 and 8 only), both published taps/group shapes on both sides, plus hand-computed corners for the block boundary, the group map and the side. MODEL GATE GREEN (test_qwen3_dflash2_draft.cpp) 16 cases / 108 assertions, `Status: SUCCESS!`, exit 0 — weights read off a REAL on-disk safetensors shard by the production loader, an IDENTITY conv proven BIT-IDENTICAL to no conv, and each conv driven ALONE through each of the three layer bodies. MUTATION-PROVEN 2026-08-19, each restored byte-for-byte and verified by sha256: deleting the call sites in `ForwardBlockLogits` (5 cases / 9 assertions red), in `ForwardWithCtxKVDev` (1/1) and in `ForwardPagedBody` (1/1); forcing `args.side` to 0 (op 2 cases/4353 assertions red, model 1/1); dropping the block mask (3/449); the wrong group map (3/7436); and dropping the uniform-block guard (1/1). The PER-STEP ROUNDING was added to that set on 2026-08-20, after the wave's second fresh review proved it had no executing assertion: replacing the bf16 branch of the `round` lambda in `src/vt/cpu/cpu_ops.cpp::DFlashGroupedConvKernel` with `return v;` compiled clean and left BOTH suites fully green, because every case in the op file ran in f32 where that rounding is the IDENTITY, and the model suite asserts only RELATIONALLY between two runs of the same kernel. Two CPU-only bf16 cases now pin it — one hand-computed against literals that differ from the round-once-at-the-end answer in six of eight outputs, one bit-exact at three shapes against a reference that rounds where UPSTREAM materializes — and under the same mutation the file is 8 cases / 2 failed, 9930 assertions / 225 failed, `Status: FAILURE!`. THREE gate repairs have now come out of this row's mutation passes and are recorded rather than hidden: activating both convs at once could not see one missing call site, the first side probe could not see a forced side, and no case at all could see the rounding policy. CUDA VERIFIED 2026-08-20 — the kernel and its registration are written and the CUDA==CPU bit-identity case exists over six shapes, and it has now COMPILED AND RUN on a device: an `rc` job on `dgx:gpu0` (GB10, sm_121a, `nvcc` 13.0) at W3 head `b29b6f886`, `BUILD_RC=0`, `COMPILE_ERRORS=0`, zero `no CUDA backend; skipping` lines, this suite among the five of six that passed (#1489). Spec `## Owed` O6 is DISCHARGED by that run. The AUTHORING HOST still has no `nvcc`, so the case still reports `no CUDA backend; skipping` there — a property of that box and no longer of the kernel	DFlash2 spec W2, #1314	`ACTIVE`	`CLAIM-SPEC-DFLASH2-W2`
`KERNEL-DFLASH2-SELECTOR-EDGES`	DFlash2 candidate-selector EDGE LATTICE — the transition scores the path walk chooses among (SPEC-DFLASH2 W3, #1314). `edge(b,l,p,c) = unary[b,l,c] + <pred_codebook[pid(b,l,p)] * project(h)[b,l], succ_codebook[cand[b,l,c]]>`, where `pid` is the request's verified ANCHOR token at step 0 (the same token for every predecessor slot, which is upstream's `anchor_token_ids[:, None, None].expand`) and the PREVIOUS step's candidate at every later step. This is what replaces the DFlash1 independent per-slot argmax: the draft no longer picks each slot alone, it scores adjacent transitions over the target head's top-K and walks the best path. Two bf16 materializations are reproduced at upstream's own two points -- the elementwise `predecessors * hidden` product and the einsum's output -- and the f32 `unary` is added after, which is torch's own promotion. UNLIKE `KERNEL-DFLASH2-GROUPED-CONV` this op is NOT specified bit-identical across backends: the rank contraction is a REDUCTION and the CUDA warp-shuffle tree sums in a different order, so it is gated within an f32 envelope and the difference is stated rather than inherited by analogy	BEYOND-PIN — `vllm/model_executor/models/qwen3_dflash2.py:208-276` (`_score_edges`, `CandidateSelector`) @ vllm-project/vllm#52816 head `66e5414c6d75a8529473d977f7458c140bbab8a0`, which superseded `19c93519` on 2026-08-19 (#1404) and is BYTE-IDENTICAL for this function; the parity pin `555967922` does not carry the architecture and this row does NOT advance it	`OpId::kDflash2SelectorEdges` + `Dflash2SelectorEdgesArgs` + decl/wrapper `include/vt/ops.h::Dflash2SelectorEdges` and `src/vt/ops.cpp::Dflash2SelectorEdges`; CPU REFERENCE `src/vt/cpu/cpu_ops.cpp::Dflash2SelectorEdgesKernel` (the authoritative impl); CUDA mirror `src/vt/cuda/cuda_ops.cu::Dflash2SelectorEdgesKernelCuda` (one WARP per (request, step, predecessor slot); `__shfl_xor_sync` contraction, dynamic shared for the gated product). Consumed by `src/vllm/model_executor/models/qwen3_dflash2.cpp::Qwen3DFlash2Model::SelectorEdgeScores`, called from `src/vllm/v1/worker/gpu/spec_decode/dflash2/speculator.cpp::Dflash2SelectCandidates`, which BOTH propose paths call -- `src/vllm/v1/worker/gpu/runner.cpp::GPUModelRunner::propose_drafts_block` (production) and `src/vllm/v1/worker/gpu/spec_decode/dflash/speculator.cpp::DflashProposeBlock`; weights loaded by `src/vllm/model_executor/models/qwen3_dflash_weights.cpp::LoadQwen3DFlash`	CPU GATE GREEN 2026-08-20 (test_ops_dflash2_selector_edges.cpp) 7 cases / 203 assertions (199 at W3; W4's review added 4), `Status: SUCCESS!`, exit 0 — upstream's OWN sequential reference loop at upstream's own parameters (batch 2, steps 4, top_k 3, rank 5, vocab 17), bit-exact in f32; the ANCHOR arm (all K predecessor rows of step 0 identical, later steps not, and changing the anchor moves step 0 alone); the PREDECESSOR indexing (swapping two of step l-1's candidate ids swaps exactly those two predecessor rows of step l); the UNARY broadcast (a per-CHILD bias, not per-edge); and the bf16 ROUNDING PLACEMENT with hand-written literals chosen so the three candidate placements answer differently (ours 7.71875, round-once-at-the-end 7.6875, unrounded 7.699830055236816) -- which is the executing assertion W2's second review proved the convolution's evidence lacked. PRODUCTION-REACHED and mutation-proven at the RUNNER: test_dflash2_runner_reach.cpp 3 cases / 14 assertions drives a real `LoadedEngine` over a synthetic Qwen3.5-dense target plus an in-memory DFlash2 draft, and the walk refusal names this op's own output (`scored-transitions=27 requests=1 steps=3 top_k=3`). MUTATION-PROVEN 2026-08-20, each restored byte-for-byte and verified by sha256 — see the row spec's `## Now`. W4's fresh review found this row's own `[B,L,K,K]` OUTPUT-lattice guard ungated and repaired it here: the refusal case's two bare `CHECK_THROWS` would each have been answered by a neighbouring guard ONCE THE NAMED ONE WAS DELETED, so deleting the lattice check alone left the suite green. Same #1518 correction as the sibling row: with the check PRESENT the case throws `scores must be [B,L,K,K]` (`src/vt/ops.cpp:3284`); with it deleted the throw falls through to `contiguous tensors required` (`:3299`). The two trailing axes are now driven by genuinely contiguous wrong-extent output views and every refusal is matched on its message; deleting the check reddens 1 case / 2 assertions. CUDA VERIFIED 2026-08-20 — the kernel and its registration are written and the CUDA-vs-CPU envelope case exists over two shapes (upstream's small one and the PUBLISHED rank 256 / K 16), and it has now COMPILED AND RUN on a device: `dgx:gpu0` (GB10, sm_121a, `nvcc` 13.0) at W3 head `b29b6f886`, `BUILD_RC=0`, `COMPILE_ERRORS=0`, zero `no CUDA backend; skipping` lines, this suite among the five of six that passed, so the warp-shuffle contraction lands inside the asserted envelope at rank 256 (#1489). Spec `## Owed` O10; the AUTHORING HOST still has no `nvcc` and still skips there	DFlash2 spec W3, #1314	`ACTIVE`	`CLAIM-SPEC-DFLASH2-W3`
`KERNEL-TOPK-PAIRS`	Top-k that EMITS the surviving (id, value) pairs — the vocabulary selection the DFlash2 candidate selector consumes (SPEC-DFLASH2 W3 / `## Risks/decisions` D2, #1314). The shipped `KERNEL-SAMPLE-*` threshold search MASKS below the k-th largest and returns no indices; this one COMPACTS the survivors and orders exactly k of them. THE TIE-BREAK IS THE CONTRACT and not an implementation detail: the pivot-bracket search converges to an exact array VALUE, so `{x >= thr}` keeps whole tie groups atomically and can hold more than k members, and a backend that chose among equals differently would reorder the selector's candidate slots and move acceptance without raising anything. The order is DESCENDING value, ties by ASCENDING index -- `torch.topk`'s CPU order and what FlashInfer's `deterministic=True` exists to provide. `num_org_vocab_padding` mirrors upstream's `logits[..., -num_pad:] = -inf` and is 0 on every path this engine ships (the DFlash lane's lm_head is the raw unpadded tensor and there is no vocab-parallel sharding), so it is gated SYNTHETICALLY rather than claimed as checkpoint coverage	BEYOND-PIN — `vllm/model_executor/models/qwen3_dflash2.py:60-64` (`_topk`) @ vllm-project/vllm#52816 head `66e5414c6d75a8529473d977f7458c140bbab8a0`: `torch.topk(scores, k, dim=-1)` off CUDA, FlashInfer's radix `top_k(..., sorted=True, deterministic=True)` on it. FlashInfer's `topk.cuh` is NOT ported — `## Risks/decisions` D2 records why: 3380 lines of general kernel (multi-CTA, deterministic mode, three tie-break modes, dynamic shared sizing) for a shape that is fixed and small here, K=16 over 248320 for about 224 rows	`OpId::kTopKValuesIndices` + `TopKValuesIndicesArgs` + decl/wrapper `include/vt/ops.h::TopKValuesIndices` and `src/vt/ops.cpp::TopKValuesIndices`; CPU REFERENCE `src/vt/cpu/cpu_ops.cpp::TopKValuesIndicesKernel` (the authoritative impl; `std::partial_sort` under the explicit comparator, because `partial_sort` is not stable and leaving ties to the algorithm's swaps would make the order an implementation accident); CUDA `src/vt/cuda/cuda_sample.cu::TopKValuesIndicesRowKernel`, which EXTENDS the shipped sort-free block-cooperative pivot-bracket search beside `ApplyTopKTopPRowKernel` in the same file rather than adding a second one. Consumed by `src/vllm/model_executor/models/qwen3_dflash2.cpp::Qwen3DFlash2Model::ComputeCandidates`	CPU GATE GREEN 2026-08-20 (test_ops_topk_values_indices.cpp) 10 cases / 196 assertions, `Status: SUCCESS!`, exit 0 — three HAND-WRITTEN tie cases (ties inside the kept set, a tie group STRADDLING the k-th boundary, a tie group larger than k), the org-vocab padding mask asserted AGAINST its own unmasked answer, that the logits are not mutated in place (unlike `vt::ApplyTopKTopP`), a -inf-saturated row, and bulk shapes against a full sort. The bulk reference is independent in ALGORITHM and not in the tie RULE, which it restates rather than derives, so the hand-written literals are the correctness anchor for the ties and the full sort is a consistency check on the selection. PRODUCTION-REACHED through `Dflash2SelectCandidates` and mutation-proven at the runner with `KERNEL-DFLASH2-SELECTOR-EDGES`. CUDA RUN 2026-08-20, and ONE ROW DISAGREES — written, registered, and now compiled and executed on `dgx:gpu0` (GB10, sm_121a, `nvcc` 13.0) at W3 head `b29b6f886`: `BUILD_RC=0`, `COMPILE_ERRORS=0`, 562 device assertions against 202 CPU-only, and the TIE rows AGREE across the two algorithms. The NaN row does not: 562 assertions, 550 passed, 12 FAILED, `Status: FAILURE!`, all twelve on `"NaN sorts first, as torch.topk does"`, because `TopKValuesIndicesRowKernel`'s `fmaxf`/`fminf` bracket returns the non-NaN operand and its `r[j] > thr` survivor test is false for a NaN, so the kernel cannot select one (#1489). The device cases now SKIP that row by name and `include/vt/ops.h` states the asymmetry; #1489 owns reconciling the kernel. Spec `## Owed` O10; the AUTHORING HOST still has no `nvcc` and still skips there	DFlash2 spec W3, #1314	`ACTIVE`	`CLAIM-SPEC-DFLASH2-W3`
`KERNEL-DFLASH2-PATH-WALK`	DFlash2 candidate-selector PATH WALK — the choosing that turns the transition lattice into k draft tokens (SPEC-DFLASH2 W4, #1314). Start at the verified anchor, take the best child of the current predecessor row, then read the NEXT step's block at the row just chosen: `previous = 0; for l in steps: index = argmax(scores[b,l,previous,:]); tokens[b,l] = cand[b,l,index]; previous = index`. The slot-to-slot dependency is the shape of the problem, so the STEP LOOP LIVES INSIDE ONE PROGRAM PER REQUEST rather than one launch per step -- upstream's own `(num_reqs,)` grid with `num_warps=1`, and what spec `## Risks/decisions` D3 requires from the first landing, because the identical sequential walk shipped host-side in DSpark and measured 28% of the 27B draft step (#436) before `SampleSequentialDevice` moved it. TWO CONTRACT POINTS decide tokens and are pinned by literal cases: a tie resolves to the LOWEST slot (`tl.max(..., return_indices=True)`'s own rule, and it picks the next step's PREDECESSOR row too, so it moves the rest of the path rather than one token), and an all `-inf` row -- upstream's fully masked lane -- resolves to slot 0 rather than to "no index". Both arms reach those answers identically: seed at `-inf` with slot index `top_k`, keep only what STRICTLY exceeds the running best (so a NaN never wins on either side), then collapse a `top_k` seed to 0 -- and the CUDA arm's per-lane scan is STRICT ONLY since W4's review, because its `(v == best && j < slot)` disjunct was unreachable once a lane had claimed anything and its one effect was letting a lane holding `-inf` claim on the `-inf` seed, which is a slot the CPU arm refuses (`[NaN,-inf]` read cpu 0 and cuda 1). The lower-slot tie rule stays in the cross-lane butterfly, where lanes combine out of slot order. GREEDY IS THE ONLY ARM, and it is upstream's greedy arm exactly: at the moved head the walk's two hand-written branches collapse into one `gumbel_noised_argmax` whose temperature is `temperature if SAMPLE_PROBABILISTIC else 0.0`, and `SAMPLE_PROBABILISTIC` is `draft_logits is not None`, which is set only for `draft_sample_method == "probabilistic"` -- a value `vllm::ParseSpeculativeConfigJson` refuses BY NAME here. UNLIKE `KERNEL-DFLASH2-SELECTOR-EDGES` this op IS specified bit-exact across backends: it performs no arithmetic, only comparisons and one gather, so there is no reduction order to differ in	BEYOND-PIN — `vllm/v1/worker/gpu/spec_decode/dflash2/speculator.py:16-79` (`_selector_walk_kernel`) and `:148-172` (`DFlash2Speculator._sample_path`) @ vllm-project/vllm#52816 head `66e5414c6d75a8529473d977f7458c140bbab8a0`. The head MOVED under this row and it moved HERE (#1404): at `19c93519` this file is +37/-34 different -- the walk carried a hand-written `temperature == 0.0` branch beside a hand-written Gumbel branch, wrote into a private `_selector_tokens` buffer that a trailing `copy_` moved into `draft_tokens`, and the speculator ALWAYS allocated a proposal distribution. At `66e5414c` the branches are one `gumbel_noised_argmax` call, `_selector_tokens` and its `copy_` are DELETED, and `draft_logits` is optional and `None` for greedy behind a new `draft_logits_spec` on the speculator base. The greedy ANSWER is identical at both heads; the SHAPE ported here is the new one. The parity pin `555967922` does not carry the architecture and this row does NOT advance it	`OpId::kDflash2PathWalk` + `Dflash2PathWalkArgs` + decl/wrapper `include/vt/ops.h::Dflash2PathWalk` and `src/vt/ops.cpp::Dflash2PathWalk`; CPU REFERENCE `src/vt/cpu/cpu_ops.cpp::Dflash2PathWalkKernel` (the authoritative impl); CUDA mirror `src/vt/cuda/cuda_ops.cu::Dflash2PathWalkKernel` / `Dflash2PathWalkKernelCuda` (ONE BLOCK per request, one warp, `__shfl_xor_sync` argmax, the step loop inside the kernel and `previous` a per-lane register every lane derives from the same butterfly). Consumed by `src/vllm/v1/worker/gpu/spec_decode/dflash2/speculator.cpp::Dflash2WalkPath`, which BOTH propose paths call -- `src/vllm/v1/worker/gpu/runner.cpp::GPUModelRunner::propose_drafts_block` (production) and `src/vllm/v1/worker/gpu/spec_decode/dflash/speculator.cpp::DflashProposeBlock` -- and whose output IS the draft a DFlash2 checkpoint proposes	CPU GATE GREEN 2026-08-20 (test_ops_dflash2_path_walk.cpp) 7 cases / 49 assertions, `Status: SUCCESS!`, exit 0 — upstream's own `_selector_walk_kernel` control flow at `SAMPLE_PROBABILISTIC=False` over three shapes including the PUBLISHED one (top_k 16, 7 steps); the PREDECESSOR CARRY with hand-written literals built so the carrying answer and the row-0 answer differ at step 1 and the off-by-one answer differs from both; the TIE resolving to the lowest slot AND then choosing the next step's predecessor row; the all `-inf` row resolving to slot 0; the GATHER, with ids that are not the identity permutation of the slots; and (W4 REVIEW REPAIR) that a NaN never wins a slot. A `>=` reduction fails THREE of those cases, not one (#1518 corrects the earlier claim that the tie rows and the -inf row are answered correctly under `>=`): turning `>` into `>=` reddens `a tie resolves to the LOWEST slot` (2 assertions, got 13 for 11 and 22 for 20), `an all -inf row resolves to slot 0` (1, got 33 for 31) and `a NaN never wins a slot` (2), for 3 cases / 5 assertions, `Status: FAILURE!` -- an ascending scan under `>=` keeps the LAST maximum, so both of the other rows answer slot K-1. What the NaN case adds is the NaN CLASS and the row on which the two backends actually diverged. RED-BEFORE: `rc=2`, 10 `error:` lines, `‘Dflash2PathWalk’ is not a member of ‘vt’`. W4's FRESH REVIEW returned FAIL on this row's shape guard and the repair is in the same wave. The `[B,L,K,K]` case mutated ONE `shape` field of a tensor `Contig` had produced, so dropping either trailing conjunct or deleting the whole `VT_CHECK` left all four suites green. The MECHANISM recorded for that was backwards and #1518 corrects it: the mutation does desynchronise the strides (`probe IsContiguous=0`), but the walk checks shape at `src/vt/ops.cpp:3331` and contiguity at `:3339`, so with the guard PRESENT the shape check still answered (`probe threw: vt: dflash2-path-walk: scores must be [B,L,K,K] matching candidate_ids at src/vt/ops.cpp:3331`) and the case did reach the guard it named. The real hazard is the DELETED state: remove the shape check and the throw falls through to `contiguous tensors required` (`:3336` once removed), so a bare `CHECK_THROWS` still passes and the deletion is invisible -- measured at 7 cases / 47 assertions, `Status: SUCCESS!`, rc 0. Both trailing axes are now driven by GENUINELY CONTIGUOUS wrong-extent lattices ([B,L,K,K-1] and [B,L,K-1,K]) matched on the message, and the three mutations redden 1 case / 1 assertion, 1/1 and 1/2. PRODUCTION-REACHED and mutation-proven at the RUNNER: test_dflash2_runner_reach.cpp 3 cases / 83 assertions drives a real `LoadedEngine` through `propose_drafts_block`, reads the drafts off the production `VT_SPEC_TRACE` line at real fd 2, and requires them to MOVE when D9's output scalars move -- which the DFlash1 per-slot argmax cannot do, because those scalars touch only the candidate values the selector reads. MUTATION-PROVEN 2026-08-20, each restored byte-for-byte and verified by sha256, each with `compile_rc` printed: dropping the `previous` carry (op 2 cases / 17 assertions red), `>` to `>=` (2/3), the `-inf` collapse to `K-1` (1/1), the gather ignoring the winner (4/26), the walk's per-request row indexing (guard suite 1/2), the `[B,L,K,K]` guard's child conjunct (1/1), its predecessor conjunct (1/1) and the whole `VT_CHECK` (1/2), `Dflash2WalkPath`'s candidate-set check and its i32-range check (guard suite 1 case / 2 assertions each, 1/4 together), DELETING the runner's walk call site (reach 2 cases / 5 red), REPLACING the walk with the DFlash1 argmax (reach 1/1, the D9-scalars case, which is the guard-independent reachability proof), and deleting `DflashProposeBlock`'s walk call (draft suite 1 case red). CUDA VERIFIED 2026-08-20 (#1518, correcting a `CUDA UNVERIFIED — owed to the operator's GPU lease` state this cell carried at the W4 merge). The kernel and its registration are written and the CUDA==CPU BIT-EXACT case exists over three shapes (a sub-warp K, the published K 16, and K 40 so the strided load and the shuffle both run past one warp) with a forced tie group and a forced all `-inf` row; and the fixture now CHAINS a third forced row -- a NaN row at step 2 predecessor 0, reached because the tie group answers slot 0 and the all -inf row answers slot 0 -- which is the row the two arms disagreed on before the lane comparator was reconciled. It COMPILES pre-merge -- `src/vt/cuda/cuda_ops.cu` is in the CUDA source list in `CMakeLists.txt` and CI's `build-cuda-fat` job builds it for ten architectures (`80;86;87;89;90a;100a;103a;110;120a;121a`) on every non-closed `pull_request` event -- and it has RUN: the operator executed this suite on `dgx:gpu0` (GB10, sm_121a) at the W4 merge commit, reporting 83 assertions on device against 49 on CPU, `Status: SUCCESS!`, zero `no CUDA backend; skipping` lines, `CUDA_OBJECTS_BUILT=34`, the chained NaN row among the increment -- so the lane-comparator deletion is verified on device too. The AUTHORING HOST still has no `nvcc` and still reports `no CUDA backend; skipping CUDA dflash2-path-walk parity` locally. Spec `## Owed` O11	DFlash2 spec W4, #1314	`ACTIVE`	`CLAIM-SPEC-DFLASH2-W4`
`KERNEL-ATTN-DSA-SPARSE-INDEX`	DeepSeek-V4 DSA "Lightning Indexer" sparse-attention SELECTION — the project's FIRST sparse candidate-selection primitive (DeepSeek-V4-Flash W3). Two ops: (1) the weighted-MQA INDEXER LOGIT `logit[t,s] = Σ_h w[t,h]·ReLU(q[t,h]·k[s])` over the causal candidate window (the per-head ReLU is load-bearing — it is what makes the indexer a learned sparse SELECTOR, not a plain attention score), where `w[t,h] = weights_proj[t,h]·index_head_dim^-0.5·index_n_heads^-0.5`; and (2) the per-row causal top-k that keeps the `index_topk=512` highest-logit keys (short-context: every candidate, ascending; else top-k with -1 padding). Distinct from every dense/paged/MLA family, which score ALL keys — this one PICKS a sparse key subset the downstream MLA then attends over. W3 also lands the two 512-wide-MLA OUTPUT seams V2/V3 lack (per-head attention-sink softmax + grouped output-LoRA `wo_a` bmm→`wo_b`) as portable host references beside it	MQA logit `vllm/v1/attention/ops/triton_fp8_mqa_logits.py:120-156` (dot→×kv_scale→ReLU→×weights→Σheads); weight fold `vllm/model_executor/layers/sparse_attn_indexer.py:203-207`; top-k `sparse_attn_indexer.py:488-497` + short-context `vllm/models/deepseek_v4/attention.py:70-86,:813-831`; sinks `deepseek_v4/nvidia/flashinfer_sparse.py:777,:896`; grouped output-LoRA `deepseek_v4/nvidia/ops/o_proj.py:58-73` @ `555967922`	Portable host reference (device kernel is a W7 residual) deepseek_v4_dsa.cpp + deepseek_v4_dsa.h: `DsaIndexerWeightFold` / `DsaIndexerLogits` / `DsaTopkSelect` / `SoftmaxWithSink` / `GroupedOutputLora`	CPU UNIT GATE GREEN (2026-07-28, `-Wall -Werror -Wextra` 0-warn): test_deepseek_v4_dsa.cpp 13/13 cases · 38 assertions — hand-derived literal cases (the ReLU clip, the weight fold, short-context all-select, full top-k, tie→smaller-index, causal-window offset, sink probability mass, sink numerical stability, grouped-LoRA) + from-first-principles double-precision references on randomized shapes (indexer logits + grouped output-LoRA rel-L2 < 1e-6). Full-model gate is multi-Spark-blocked (156.7 GiB); MHC (W5) + sqrtsoftplus/hash MoE (W6) + device kernel + forward integration (W7) are named residuals	deepseek-v4-flash spike §W3	`SPIKE`	`CLAIM-DEEPSEEK-V4-W3`
`KERNEL-ATTN-DSA-COMPRESSOR`	DeepSeek-V4 DSA COMPRESSOR + fp8_ds_mla KV-cache state — the second half of the sparse-attention stack (DeepSeek-V4-Flash W4). Where `KERNEL-ATTN-DSA-SPARSE-INDEX` SELECTS keys, this POOLS + QUANTIZES them into the compressed latent the MLA reads and defines how it is cached across steps. Three ops: (1) the softmax-weighted window POOL — at a compress boundary the compressor gathers `(1+overlap)·compress_ratio` KV-state rows and computes, PER head-dim column, `softmax(score, dim=0)·kv` (each channel pools the window with its OWN weights — the load-bearing nuance), then RMSNorm; (2) the fused save-time APE add `score_state = score + ape[position % compress_ratio]`; (3) the fp8_ds_mla KV-cache STATE layout — the 512-wide latent split into a 448-wide NoPE part quantized to FP8 e4m3 with per-64 UE8M0 power-of-two block scales (exponent `= ceil(log2(absmax/448))`, byte `= exp+127`) and a 64-wide RoPE part stored bf16, at a 576-byte token stride with a padded 7+1 scale region — plus the dequant READ (`nope = e4m3·2^(byte-127)`, `rope = bf16`)	pool+RMSNorm `vllm/models/deepseek_v4/common/ops/fused_compress_quant_cache.py:198-218`; save-time APE `common/ops/save_partial_states.py:92-101`; fp8_ds_mla store `fused_compress_quant_cache.py:220-297`; layout `deepseek_v4/compressor.py:307-309`; dequant READ cross-checked to SGLang `v0.5.15` `dsv4/dequant_k_cache.py:12-18,:122-136` @ `555967922`	Portable host reference (device kernel is a W7 residual) deepseek_v4_compressor.cpp + deepseek_v4_compressor.h: `CompressorSaveScoreApe` / `CompressorPoolNorm` / `MakeFp8DsMlaLayout` / `Fp8DsMlaEncodeToken` / `Fp8DsMlaDecodeToken`	CPU UNIT GATE GREEN (2026-07-29, Debug full-library build, 0-warn on the new TUs): test_deepseek_v4_compressor.cpp 12/12 cases · 164 assertions — hand-derived literal cases (APE modulo wrap; per-column softmax pool proven load-bearing via the column-ratio-survives-RMSNorm case; window masking; V4 layout 448/64/576/7+1; all-ones→UE8M0 byte 119 exact round-trip; value-3→byte 120; bf16 rope verbatim) + from-first-principles double-precision references (pool+norm rel-L2 < 1e-6; independent UE8M0 scale-byte recompute; encode→decode round-trip < 0.05 fp8 granularity). RED-first PROVEN: perturbing the scale bias `+127→+126` fails 4 cases / 135 assertions; revert restores 12/12. Honest gate form: hand-case + structural review vs vLLM+SGLang `file:line` (fixed-config 167B not constructible at a tiny shape ⇒ NOT a dumped-oracle rel-L2). Full-model gate multi-Spark-blocked (156.7 GiB); MHC (W5) + sqrtsoftplus/hash MoE (W6) + the fused device kernel + forward integration (W7) are named residuals	deepseek-v4-flash spike §W4	`SPIKE`	`CLAIM-DEEPSEEK-V4-W4`
`KERNEL-MHC-SINKHORN`	DeepSeek-V4 Manifold/Markov Hyper-Connections (MHC) — the Sinkhorn-normalized hc_mult-stream residual manifold (DeepSeek-V4-Flash W5, the hardest V4 brick). V4 replaces the plain `residual + RMSNorm` stream with a `[tokens, hc_mult=4, hidden]` MANIFOLD of parallel residual streams, mixed at every attn/ffn boundary by a doubly-stochastic matrix and collapsed by a learned head. Four ops: (1) the `hc_sinkhorn_iters=20` Sinkhorn normalization of the hc_mult×hc_mult mixing matrix — a row-softmax seed (`+eps`), a col-norm, then $(\text{iters}-1) \times [\text{row}-\text{norm}, \text{col}-\text{norm}]$ toward a doubly-stochastic matrix (the AXIS ALTERNATION and the ITERATION COUNT are load-bearing at non-converged counts — RED-first proven); (2) the mHC pre mix — flatten the streams, project through `hc__fn` with a FOLDED weight-free RMSNorm `rsqrt(sqrsum/(hc·H)+rms_eps)`, split into pre/post/comb gates (`pre=σ+hc_eps`, `post=σ·hc_post_alpha(2.0)`, `comb=Sinkhorn`), collapse to the single `layer_input`, and optionally FOLD the model's attn_norm/ffn_norm RMSNorm; (3)* the mHC post mix — fold the block output back into the manifold via the comb matrix (`Σ_i comb[i,j]·res[i,h]`) + the post gate; (4) the hc_head collapse — weight-free RMSNorm → `hc_head_fn` → sigmoid gate → weighted stream sum → one hidden vector. EAGER-REF FINDING: corrects the W0 "ZERO eager reference upstream" premise — the pinned vLLM DOES ship an eager PyTorch reference (`mhc/torch.py` `mhc_pre_torch`/`mhc_post_torch`, `triton.py` head collapse); four upstream impls (torch.py, tilelang_kernels.py `_sinkhorn_fwd`, tilelang.py, SGLang mhc.py) agree byte-for-byte on the Sinkhorn	mHC pre/post + Sinkhorn `vllm/model_executor/kernels/mhc/torch.py:56-106` (byte-identical `tilelang_kernels.py:126-153` `_sinkhorn_fwd`, `tilelang.py` `mhc_pre_big_fuse_with_norm`); head collapse `triton.py:108-140` + `tilelang.py:720-748`; constants `hc_post_alpha=2.0`/`hc_pre_eps=hc_sinkhorn_eps=hc_eps` `vllm/models/deepseek_v4/nvidia/model.py:818-821,:886-894,:1023-1041`; cross-checked SGLang `v0.5.15` `python/sglang/srt/layers/mhc.py:110-126` @ `555967922`	Portable host reference (device kernel + `DeepseekV4Model::Forward` assembly are W7 residuals) deepseek_v4_mhc.cpp + deepseek_v4_mhc.h: `MhcSinkhorn` / `MhcPre` / `MhcPost` / `HcHeadCollapse`	CPU UNIT GATE GREEN (2026-07-29, Debug full-library build, `-Wall -Werror -Wextra` 0-warn on the new TUs): test_deepseek_v4_mhc.cpp 14/14 cases · 125 assertions — hand-derived literal cases (all-zero Sinkhorn → uniform doubly-stochastic 1/hc; symmetric-2×2 fixed point `[[.75,.25],[.25,.75]]`; iteration-count load-bearing; MhcPre fn=0 gate midpoints; RMSNorm fold `[1,3]→[1,3]/√5`; MhcPost identity-comb + post-add; mix sums over the first comb index; hc_head fn=0 stream mean) + from-first-principles DOUBLE-PRECISION references (Sinkhorn/MhcPre/MhcPost/HcHead f32==f64 rel-L2 < 1e-5..1e-4; doubly-stochastic convergence to row/col sums=1). RED-first PROVEN both levers: perturb the Sinkhorn iteration count (`iters-1→iters-2`) fails 1 case/9 assertions AND swap a normalization axis fails 2 cases/12 assertions (caught by a dedicated SMALL-iteration-count gate, since at 20 iters the Sinkhorn has converged and ±1 is within tolerance); revert restores 14/14·125. Honest gate form: DERIVED-eager-reference + hand-case + structural review vs vLLM+SGLang `file:line` (fixed-config 167B not constructible at a tiny shape ⇒ NOT a dumped-oracle rel-L2). OPEN QUESTION: end-to-end bf16 residual/layer_input rounding between steps is a W7 device concern, not folded into these f32/f64 refs. Full-model gate multi-Spark-blocked (156.7 GiB); sqrtsoftplus/hash MoE (W6) + device kernel + forward assembly (W7) are named residuals	deepseek-v4-flash spike §W5	`SPIKE`	`CLAIM-DEEPSEEK-V4-W5`
`KERNEL-MOE-SQRTSOFTPLUS-HASH`	DeepSeek-V4 MoE router + clamped-SwiGLU deltas — the three genuinely-new-vs-V2/V3 MoE pieces (DeepSeek-V4-Flash W6). V4 keeps the DeepSeek grouped-GEMM / 256-expert w13/w2 / shared-expert / NVFP4 machinery (REUSED, not re-ported) but replaces three primitives. Three ops: (1) the router SCORE function `sqrt(softplus(x))` (`softplus(x)=log(1+exp(x))`, then sqrt) — distinct from V2/V3's sigmoid/softmax `noaux_tc`; the sqrt∘softplus COMPOSITION is load-bearing (RED-first proven); (2) the router: score all experts, add `e_score_correction_bias` for SELECTION ONLY, pick top-k OR — for the first `num_hash_layers` HASH layers — look experts up directly in the `tid2eid` [vocab, topk] token-id→expert table (BYPASSING top-k), GATHER weights from the UNBIASED scores, renormalize, ×`routed_scaling_factor` (the bias-affects-selection-not-weights split + the hash bypass are load-bearing, both RED-first proven); (3) the clamped SwiGLU expert activation `SiluAndMulWithClamp` — `gate` clamped `max=limit` (max ONLY), `up` clamped `[-limit,+limit]` (BOTH sides), then `gate·sigmoid(alpha·gate)·(up+beta)`; the ASYMMETRIC clamp is load-bearing (RED-first proven). MegaMoE (SM100-only) is NOT the GB10 target — this mirrors the FusedMoE-fallback router GB10 runs	score `vllm/model_executor/layers/fused_moe/router/fused_topk_bias_router.py:88`; router `fused_topk_bias_router.py:75-118` (`_topk_softplus_sqrt_torch`) + hash branch `:100-106` + dispatch `:254-265`; hash-table wiring `vllm/models/deepseek_v4/nvidia/model.py:562-578,:686,:696-717`; FusedMoE fallback `nvidia/model.py:647-691`; clamped SwiGLU `vllm/model_executor/layers/activation.py:197-201` (`SiluAndMulWithClamp.forward_native`), used by DeepseekV4MLP `nvidia/model.py:126-133`; cross-checked SGLang `v0.5.15` `python/sglang/srt/layers/moe/{topk.py:1013-1014, hash_topk.py:137-180}` @ `555967922`	Portable host reference (device kernels landed W7-device — see `KERNEL-DSV4-W7-DEVICE`) deepseek_v4_moe.cpp + deepseek_v4_moe.h: `SqrtSoftplus` / `SqrtSoftplusRouteTopk` / `ClampedSwiGLU`	CPU UNIT GATE GREEN (2026-07-29, Debug full-library build, `-Wall -Werror -Wextra` 0-warn on the new TUs): test_deepseek_v4_moe.cpp 12/12 cases · 716 assertions — hand-derived literal cases (sqrt∘softplus composition `softplus(x)=4 ⇒ score=2`; bias flips selection but weight stays the UNBIASED 1.0 not 3.0; renormalize by the unbiased sum; routed_scaling_factor; hash `tid2eid` picks {3,1} where top-k would pick {2,0}; asymmetric clamp gate=-5 kept vs up clamped to -2; gate/up clamp boundaries; alpha/beta) + from-first-principles double-precision references (router f32==f64 rel-L2 < 1e-5 + exact ids; SqrtSoftplus f64 + monotonicity; ClampedSwiGLU rel-L2 < 1e-6). RED-first PROVEN all three levers: drop the sqrt → 8 cases/493 assertions fail; gather weights from the BIASED scores → 2 cases/181 fail; symmetric-clamp the gate → 2 cases/6 fail; revert restores 12/12·716. Honest gate form: host-reference + hand-case + structural review vs vLLM+SGLang `file:line` (fixed-config 167B not constructible at a tiny shape ⇒ NOT a dumped-oracle rel-L2). Full-model gate multi-Spark-blocked (156.7 GiB); the device kernels reuse the existing grouped-GEMM + `DeepseekV4Model::Forward` assembly (W7) + the strict/near-tie engine gate (W8) are named residuals	deepseek-v4-flash spike §W6	`SPIKE`	`CLAIM-DEEPSEEK-V4-W6`
`KERNEL-DSV4-W7-DEVICE`	DeepSeek-V4-Flash W7-DEVICE — the four NEW V4 op families' CUDA kernels, each a 1:1 DEVICE port of the landed portable HOST reference (the oracle the four SPIKE rows above pinned), registered through the vt OpProvider seam (`kDeepseekV4{Mhc,Dsa,Compressor,Moe}`) so `DeepseekV4Model::ForwardDevice` can dispatch them: (MHC) Sinkhorn + mHC pre/post + hc_head collapse; (DSA) indexer weight-fold + weighted-MQA ReLU logits + causal top-k + per-head attention-sink softmax + grouped output-LoRA; (Compressor) softmax-window pool + RMSNorm + save-time APE + fp8_ds_mla KV encode (UE8M0 block scale + e4m3, bf16 rope) / decode; (MoE) sqrtsoftplus/hash router + clamped SwiGLU. The 512-wide MLA attention + expert grouped-GEMM REUSE the existing NVFP4/FP8 kernels (`kMlaDecodeAttention`/`kMoeGroupedGemmNvfp4`, `cuda_mla_attn.cu`/`cuda_moe*.cu`) and are NOT re-ported	the SAME `file:line` the host refs cite (the `KERNEL-{MHC-SINKHORN,ATTN-DSA-SPARSE-INDEX,ATTN-DSA-COMPRESSOR,MOE-SQRTSOFTPLUS-HASH}` upstream columns), @ `555967922`	cuda_deepseek_v4.cu (kernels + host-vector launchers + OpProvider registration) + deepseek_v4_device.h / deepseek_v4_device.cpp (seam resolvers); `DeepseekV4Model::ForwardDevice` composes them (deepseek_v4.cpp)	DGX GB10 (sm_121a) UNIT GATE GREEN + RUNTIME-VERIFIED (2026-07-29): test_cuda_deepseek_v4.cpp 11/11 cases · 153 assertions — each device kernel vs its host-ref oracle at small shape: BIT-EXACT ids (DSA causal top-k, sqrtsoftplus/hash router selection), `-inf` mask exact (indexer out-of-window), near-tie rel-L2 < 1e-4 for the fp reductions (Sinkhorn, pool/softmax, sqrtsoftplus — device `expf`/`sqrtf`/`rsqrt` vs host), fp8_ds_mla encode→decode within the e4m3 granularity bound, bf16 rope bit-exact; PLUS the ForwardDevice composition gate (device forward == host forward, rel-L2 < 2e-3 over the 4-family tiny-config interleave). compute-sanitizer memcheck 0 errors. RED-first PROVEN: dropping the sqrt in the device sqrtsoftplus fails 3 cases / 6 assertions (sqrtsoftplus + router weights + ForwardDevice); revert restores 11/11·153. Build: CUDA `-Werror` clean (the #155 voxtral GCC-13 `-O2` array-bounds/stringop false positive neutralized locally). Honest 3-state: RUNTIME-VERIFIED at small shape on real GB10; the real-checkpoint paged-engine e2e stays W8 (156.7 GiB does not fit ONE GB10). DECODE GLUE-FOLD LANDED (2026-08-03, GB10 sm_121a): the resident-decode `norm_rope_rows` + MHC-pre kernels folded FP64→FP32 — `VT_V4_ROPE_FLOAT` (fused norm+RoPE 4.58→0.46 ms/step ~10×, decode +6.1%) + `VT_V4_MHC_LEAN` (finish block 256→1024 + sqrsum-fold, +0.7%, floored by 86 sequential single-block launches/step), both default-ON + BYTE-EXACT (decode ids `=1`/`=0` token-identical via the resident-decode path `--gpu --kv-cache`); `test_cuda_deepseek_v4` Brick-7 + Brick-B 20/20·67073 PASS on GB10; net decode 14.02→14.96 tok/s → 90.7% of ds4 ~16.5 (`CLAIM-DSV4-ROPE-FLOAT` / `CLAIM-DSV4-MHC-LEAN`; ds4 bar corrected from the unreproduced 17.13 anchor to the fair same-session ~16.5, the later MHC-SINK4 reached ~96%)	deepseek-v4-flash spike §W7	`SPIKE`	`CLAIM-DEEPSEEK-V4-W7-DEVICE`
`KERNEL-KDA-DELTA`	Kimi Delta Attention (KDA) gated-linear-attention delta vs plain GDN — a genuinely new gated-linear-attention family, the shared unblocker for Kimi-Linear-48B and Kimi-K3 (W4). `KimiGatedDeltaNetAttention` SUBCLASSES `GatedDeltaNetAttention`, so its conv-state/cache layout, `GDNAttentionMetadata`, chunked-delta recurrence and WY solve are REUSED from our landed GDN — this row owns ONLY the four KDA-specific deltas plain GDN lacks: (1) a per-channel `[H,D]` low-rank decay via an `f_a_proj→f_b_proj` bottleneck (GDN has only a per-HEAD scalar decay from `A_log`); (2) the decay GATE `g = -exp(A_log[h])·softplus_β(g1+dt_bias)` per channel (β=1, thr=20; `kda_gate_fwd_kernel` decode) + its chunk-local cumulative-sum prefill variant (`kda_gate_cumsum_fwd_kernel`, folds `RCP_LN2`); (3) the sigmoid-gated output norm `FusedRMSNormGated(head_dim, activation="sigmoid")` = `rmsnorm(x)·w·σ(g)` (the gated-linear-attention output norm GDN lacks); (4) three separate q/k/v short causal convs (`conv_size=4`, silu) + the q/k L2-norm preprocessing (`x/sqrt(Σx²+eps)`, SUM not mean). ADDITIVE — does NOT touch `cuda_gdn.cu`/`gdn_attn.cpp`, so the Qwen3.6-27B/35B GDN gate is structurally untouched (like DSA kept shared-MLA untouched)	decay bottleneck `vllm/model_executor/layers/mamba/gdn/kimi_gdn_linear_attn.py:142-156,:245`; decay gate `vllm/third_party/flash_linear_attention/ops/kda.py:1541-1600,:1603-1646`; chunk-cumsum `kda.py:1182-1254,:1257-1303`; gated norm `kda.py:463-487` (`:436` eps=1e-5); short conv `kimi_gdn_linear_attn.py:171-198,:324-356`; q/k L2-norm `kda.py:1511-1513` + `ops/l2norm.py:42-43,:96` @ `555967922`	Portable host reference (device kernel is a named residual) kimi_kda.cpp + kimi_kda.h: `KdaLowRankDecay` / `KdaDecayGate` / `KdaDecayGateChunkCumsum` / `FusedRMSNormGated` / `KdaShortConv` / `L2NormRows`	CPU UNIT GATE GREEN (2026-07-28, `-Wall -Werror -Wextra` 0-warn): test_kimi_kda.cpp 14/14 cases · 36 assertions — hand-derived literal cases (f_b∘f_a bottleneck; `-exp(A_log)·softplus` with the >thr linearisation; per-head A_log + per-channel dt_bias; chunk-cumsum reset+`RCP_LN2` fold; sigmoid-gated norm; swish-vs-sigmoid branch; per-head-dim normalisation; causal-depthwise+silu conv; zero-init-state edge; L2-norm SUM-not-mean) + from-first-principles double-precision references on randomized shapes (decay gate, gated norm, short conv rel-L2 < 1e-6). Honest gate form: host-reference + structural review, NOT a dumped-oracle rel-L2 — the REAL e2e gate is the Kimi-Linear-48B-A3B proxy vs the pinned oracle (DGX-blocked; K3 2.8T does not fit one GB10). Named residuals: the KDA CUDA device kernel + the Kimi-Linear-48B proxy gate — anchor `tests/vllm/models/test_kimi_kda.cpp:41`	kda-kernel-delta spike	`SPIKE`	`CLAIM-KDA-KERNEL`
`KERNEL-ATTN-DENSE-FLASH`	Flash-TILED dense non-causal attention — the SHARED-MEMORY-TILED form of `AttentionDenseFast` for long non-causal contexts (multimodal-speed §14, the Whisper AUDIO encoder — hd-64, non-causal, 1500 frames × 32 layers). A block of `kFlashBr=16` query-warps (512 threads) SHARES each streamed `kFlashBc=64`-column K/V tile out of shared memory (classic FlashAttention K/V tiling): the CTA cooperatively loads a K/V tile into shared memory, then each warp runs its online-softmax update reading K/V from shared memory, killing `AttentionWarpKernel`'s O(t²) redundant global K/V re-reads (one full K/V sweep per (query,head)). One q-head per CTA (all warps share the GQA kv-head). BIT-IDENTICAL to `AttentionDenseFast`: the per-warp arithmetic (per-lane head_dim grouping `lane+32k`, butterfly `__shfl_xor`, sequential j-order, f32 online-softmax `m`/`l`/`acc`) is copied verbatim, only K/V bytes come from shared memory instead of global ⇒ token-identical by construction. Head_dim-generic (`npl=(d+31)/32`; the register blocking allows d≤256 but the K/V tile's dynamic shared memory is what BINDS — d≤192 bf16 / d≤96 f32 under CUDA's default 48 KiB cap, see the 2026-08-21 entry): since 2026-07-28 (multimodal-speed §16, `CLAIM-MM-SPEED-QWEN-IMAGE`) ALSO the default for the Qwen3-VL / Qwen3.6-27B VISION tower per-frame self-attention (hd-72, non-causal, 784 patches) — byte-identical to the warp `AttentionDenseFast` it replaced there (bench 0/1,003,520 mismatch; STRICT image/video e2e 32/32)	STRUCTURE ported 1:1 from vendored FlashAttention-2 `compute_attn_1rowblock` flash_fwd_kernel.h:52 (sK/sV shared tiles :163-165 + the `for(int n_block…)` K/V-tile stream + online rescale); non-causal encoder dispatch cross-checked to vLLM `WhisperEncoderAttention` whisper.py:255	`OpId::kAttentionDenseFlash` + decl ops.h + wrapper/validation ops.cpp; CUDA `AttentionDenseFlashKernel`/`AttentionDenseFlashKernelCuda` cuda_ops.cu; CPU maps to `AttentionKernel` (byte-identical) cpu_ops.cpp; wired default in whisper_audio.cpp (`VT_WHISPER_ENC_WARP=1`/`VT_WHISPER_ENC_EAGER=1` A/B) + qwen3_vl_vision.cpp (vision tower default since §16; `VT_QWEN3VL_ATTN_WARP=1`/`VT_QWEN3VL_ATTN_EAGER=1` A/B)	GPU GATE GREEN on dgx (2026-07-28, GB10 sm_121a, base `af1ed76b`): CUDA `-Werror` 0-warn (cutlass-ON + FA2-ENABLED banners). `test_voxtral_e2e` 16/16 default-flash; flash/warp/eager token dumps md5-IDENTICAL (`89923566…`) ⇒ ZERO token flips; goldens md5 UNCHANGED (`voxtral_golden.json 8ab87b7e…`, `voxtral_neartie.json 937b9ad3…`, before==after). Proof-of-run nsys `AttentionDenseFlashKernel` 32 inst, ZERO `AttentionWarpKernel`/naive on encoder; RED confirmed (corrupt kernel → gate FAILS → restore → 16/16); `compute-sanitizer --tool memcheck` 0 errors; 3 runs byte-identical. A/B (same binary, `flock`, rep0 dropped): attention 35.11 → 19.29 ms/layer (1.82×, NON-OVERLAPPING); encoder forward ~1834 → ~1375 ms (1.33×). NOT at parity: ~1.37 s vs vLLM ~43 ms TTFT (~32×, was ~44×) — the scalar warp-per-query recurrence is now serial-latency-bound over 1500 keys (L2 already served much of the redundant reads ⇒ 1.8× not 16×); gap-closer is a tensor-core MMA hd-64 non-causal FA2 instantiation (LARGE) + resident encoder weights (MEDIUM). Vision tower (§16, 2026-07-28): extended to the Qwen3-VL/27B tower (hd-72, 784 patches) — STRICT image/video e2e 32/32, bench flash-vs-warp 0/1,003,520 mismatch, nsys default 4B e2e `AttentionDenseFlashKernel` 24 inst/zero warp, RED 30/46→46/46, sanitizer 0; A/B warp 148.3→flash 142.3 ms = 1.04× (small — the vision attention at t=784 is serial-latency-bound not bandwidth-bound; the tower already BEATS vLLM at 0.57× eager) 2026-08-21 (`CLAIM-ATTN-RUNG-VISIBLE`, issue #1544, spec attention-rung-visibility.md): the advertised head_dim contract is now the LAUNCHABLE one, and the naive rung stops being a silent default. The op stated `d <= 256` (cuda_ops.cu `LaunchAttentionDenseFlash`) while requesting `2kFlashBcdsizeof(Tin)` bytes of DYNAMIC shared memory with no `cudaFuncSetAttribute` anywhere in `src/vt/cuda/`, so the driver's default 48 KiB cap made the real ceiling 192 bf16 / 96 f32* — Kimi (192 f32, 96 KB) and Qwen3.5 (256) would have taken a bare launch error from the `cudaGetLastError` at the bottom of the launcher, naming nothing they could do instead. The bound now lives in ops.h as `AttentionDenseFlashSmemBytes` / `AttentionDenseFlashMaxHeadDim`, PURE host arithmetic so a box with no GPU can execute it, tied to the kernel by two `static_assert`s on `kFlashBc` and the register blocking; the launcher refuses above it naming `vt::AttentionDenseFast`, which uses NO shared memory and does serve those widths. NARROWING was chosen over `cudaFuncSetAttribute(cudaFuncAttributeMaxDynamicSharedMemorySize)`: d=256 f32 wants 128 KiB, above the opt-in per-block cap of the consumer Blackwell parts gated here, so the opt-in would still leave the widest advertised width a lie AND cannot be verified without a device. Strictly additive for callers — the bound is INCLUSIVE, so d=192 bf16 lands exactly on 49152 and still launches. Mirrors `supports_head_size` / `get_supported_head_sizes` backend.py:155-163, consulted BEFORE dispatch rather than discovered by launching. Same change adds check-attention-rung-consistency.py (preflight + CI): a model TU naming `vt::Attention` needs a `// VT-ATTN-NAIVE:` reason beside the call, so the six deliberate sites now say why and a new author gets a red instead of a silent ~500x. `kAttention` and every existing caller's numerics are UNTOUCHED by construction — the checker executes no model code and the head_dim guard only fires where the launch already failed. CPU-GATED: checker green, with unit + mutation cases in test_check_attention_rung_consistency.py. NEITHER the site triple this cell used to quote (9 sites / 6 marked / 3 excused) NOR that suite's case count is recorded here, for one reason: a count of another file stored in this one reds on every row that legitimately adds or removes a `vt::Attention` call, which is the drift lock AGENTS.md `## Records` forbids and the same defect #1629 names inside that suite. Run the checker for the live triple and the suite for the live count, new head_dim contract cases in test_ops_attention.cpp. OWED #1573: the on-device refusal case and its reachability mutation are PENDING a lease — `dgx:gpu0` was held by the developer, and the CPU cases pin the arithmetic, never that the launcher calls it.	multimodal-speed §14 + §16	`ACTIVE`	`CLAIM-MM-SPEED-AUDIO-ENC-KERNEL` + `CLAIM-MM-SPEED-QWEN-IMAGE` + `CLAIM-ATTN-RUNG-VISIBLE`
`KERNEL-MOE-ROUTING`	Router top-k, align, permute/unpermute, combine, activation	core MoE sources `CMakeLists.txt:1135-1157`; M=1 decode parallelization mirrors `topk_softmax_kernels.cu:192-242,494-537` (moeTopK/topkGating) + `moe_align_sum_kernels.cu:147-185,295-324`; grouped-topk (`noaux_tc`) `fused_moe/router/grouped_topk_router.py:106-161` (`forward_native`; the fused `ops.grouped_topk` at `:28-70` is the same formula), upstream tests `tests/kernels/moe/test_grouped_topk.py`, `test_routing.py`	cuda_moe.cu:349; parallel router argmax cuda_moe.cu:61; parallel moe_align BlockScan cuda_marlin_repack.cu:224; grouped-topk (MLA campaign W3) — additive `MoeRouterTopKArgs` fields + optional `e_score_correction_bias` arg ops.h, CPU ref `MoeRouterGroupedTopKKernel` cpu_ops.cpp + CUDA `MoeRouterGroupedTopKKernel` cuda_moe.cu. A SEPARATE kernel: `num_expert_group == 0` still dispatches the original, so the existing router is byte-identical by construction	routing tests; byte-exact router+align parity test_ops_moe_grouped.cpp:451; 35B 315/315 gate; grouped-topk test_ops_moe_router_grouped.cpp at REAL DeepSeek-V3 dims (256 experts, n_group=8, topk_group=4, top_k=8, sigmoid, routed_scaling 2.5, WITH `e_score_correction_bias`) vs an INDEPENDENT sort-based transcription of the upstream formula, plus isolated cases for bias-selects/unbiased-weights, top-2-sum-vs-max group scoring, the group mask excluding the global argmax, and renorm-before-scaling; CPU-vs-CUDA ids EXACT + run-to-run bit-reproducible	inventory	`ANCHOR-BACKFILL`	`CLAIM-MOE-DECODE-PARALLEL-1`
`KERNEL-MOE-UNQUANTIZED`	Unquantized grouped/batched MoE GEMM	core MoE sources `CMakeLists.txt:1135-1157`; upstream `tests/kernels/moe/test_unquantized_backend_selection.py`	activation/combine subset cuda_moe.cu:349; grouped execution remains NVFP4-specialized	MoE tests, grouped tests	inventory	`PARTIAL`	-
`KERNEL-MOE-QUANTIZED`	FP8/INT8/NVFP4/MXFP4 grouped MoE	CUTLASS/FP4 builds `CMakeLists.txt:865-1002`; NVFP4 oracle `fused_moe/oracle/nvfp4.py:38-276`	NVFP4 fallback cuda_matmul_nvfp4.cu:761, Marlin cuda_moe_marlin.cu:156	NVFP4 grouped tests; 35B gate	inventory	`PARTIAL`	-
`KERNEL-MOE-SPECIAL`	Cooperative top-k, DSV3 router, Hadacore, DeepEP/AllToAll integrations	target/build sources `CMakeLists.txt:408-424,1063-1072,1282-1299`; FlashInfer comm probes `vllm/utils/flashinfer.py:171-203`	-	-	inventory	`INVENTORIED`	-
`FIX-FP8-PLAN-CAPTURE`	The fp8 cuBLASLt lane's plan cache (`VT_FP8_PLAN_CACHE`) defaults ON, because on CUDA 13.3 `cublasLtMatmulAlgoGetHeuristic` fails inside CUDA-graph capture and the uncached fp8 lane queries it per call -- the fp8 half of #1732, whose bf16/f32 half is PR #1741	vLLM is structurally immune: torch caches the selected cuBLASLt algo per shape and vLLM warms eagerly before capture (the argument recorded in #1741's spec); the cache is the vt-runtime mirror of that reuse	flag fp8_plan_cache.h, the two flagged lanes cuda_matmul.cu	red-first polarity table in test_fp8_plan_cache.cpp; default/rollback ctest arm pair for `test_ops_fp8_cutlass`; GPU default-env graphed 35B gate after #1741 lands (measured 2026-08-24 with the env forced: token-exact on all three arms)	fix-fp8-plan-capture.md	`DONE`	#1843
`KERNEL-GDN-PACKED-DECODE`	Pure non-spec packed GDN recurrent decode for FP16/BF16/F32	default and dispatch `vllm/envs.py:117,1123-1125`, `qwen_gdn_linear_attn.py:1286-1298,1644-1695`; packed body/launch `fla/ops/fused_recurrent.py:255-478`; upstream test `tests/kernels/test_fused_recurrent_packed_decode.py:13-98`	public API, validation, portable CPU recurrence, registered CUDA kernel, register-resident perf kernel GdnPackedDecodeRegTileKernel + rollback flag + CPU flag test, exact model dispatch (`ShouldUsePackedGdnDecode`, `GdnBlockPaged`), request-identity slot pool, trace finalizer, production component driver, and component finalizer	Clean `f18ca23`/`9ad8fb7` close G0/G1; `f344dec` closes W1D2/G2; `7ff713e` + `24cea4f` close W1D3 structure; the identity-keyed slot fix at `c172336` proved the c16 repair (model gates 235/235, `--diagnostic-c16` 3/3). W1D3 CLOSES on EQUIVALENCE over eight sealed component roots + the 8-pair locked c16 A/B (−0.205% ± 0.30, <1σ) + the 24-window trace attribution (packed GPU-cheaper, kernel compute −1.30..−1.58%/step, no attributable packed-side cost): no STABLE regression on any axis; every failing axis is a sign-flipping band-edge statistic of a true-zero effect. The eighth (first 22-leg: cold-discard pair + 5 reps) seal `complete-failed` at 38/40 + 8/8 memory, stability clean, `validation_error=None`, paired-consistency PASS at BOTH c2/c16 (the two fails are c2 `median_tpot_ms` 0.9899 and c2 pooled `p99_ttft_ms` 0.8464). Packed stays the default (exact-upstream semantics, 48-launch reduction); `VT_GDN_PACKED_DECODE=0` is the rollback; no `complete-pass` marker exists and NO speed credit is claimed. Component tests 79/79, all tools 162/162; runner test 8/8. Correctness blast radius in BENCHMARKS; closure ledger. qkvz (`KERNEL-GEMM-BF16` W2) is UNBLOCKED and the exact grid is authorized. 2026-07-16 PERF LEVER (test-first, CPU-gated, DGX-pending): the named +2.06 ms/step recurrence-tiling lever is ported into `GdnPackedDecodeRegTileKernel` — one warp per `[BV=32,BK]` tile, state block REGISTER-RESIDENT (no shared-state round-trip, no cross-warp `__shfl` reduction, no `__syncthreads`; mirrors vLLM FLA `num_warps=1`/`num_stages=3` `fused_recurrent.py:256-336`), default-on behind `VT_GDN_PACKED_REG_TILE` (=0 → legacy bit-for-bit, same binary), selected for `bv==32 && dk in {32,128}`. Boundary-fixture bit-exactness preserved (sequential per-row Dk reduction). PROOF FAILED (dgx 54f0541): oracle boundary FAIL + c16 700.5 vs 793.6 tok/s → default flipped OFF (`309c218`). 2026-07-16 RESOLVED via the sanctioned vendored Triton cubin (`CLAIM-GDN-DECODE-TRITON`): Phase-1 cuobjdump MEASURED the codegen cause — vLLM FLA decode cubin REG:205/0-spill vs the naive hand port REG:255+STACK:48 (spills); DECISION = vendored cubin (`gdn_decode_h48`, 27B-only, `VT_GDN_PACKED_DECODE_TRITON` default OFF, hand kernel stays default). DGX gates: AOT op test 28/28, full GDN 49/49, oracle boundary 12/12, 27B model gate 235/235 token-exact with the Triton path ON, memcheck 0/0. c16 A/B: triton [817.51, 821.06, 822.55] vs legacy [813.77, 815.62, 815.30] tok/s — paired mean +5.48 tok/s (+0.67%), monotone (+3.74/+5.44/+7.25), 3/3 pairs positive; mean TPOT triton [161.04, 160.49, 160.35] vs legacy [162.09, 161.65, 161.93] = -1.26 ms (-0.78%) (median TPOT -1.13 ms); w0 cold-discard (triton 821.48/160.44) excluded. Anchors: shim, vendored cubin, launcher `TryTritonPackedDecode`, default-ON flag header + CPU flag test, AOT test. ACCEPTANCE MET (oracle PASS + consistent c16 TPOT improvement + no throughput regression). 2026-07-16 DEFAULT FLIP ON (`CLAIM-GDN-DECODE-TRITON-FLIP`): per MIRROR policy (the vendored kernel IS vLLM's exact token-identical FLA kernel, run by default upstream) `VT_GDN_PACKED_DECODE_TRITON` flipped OFF→ON, joining the sibling GDN Triton kernels; `=0` is the same-binary rollback. Test-first default-ON pure-header predicate + CPU flag test (RED→GREEN 10/10). 35B: no specialization — historically excluded at the model level by a dense-only term in `ShouldUsePackedGdnDecode` and by a launcher guard that rejected `Hv=32`; BOTH are gone — `GDN-MOE-BF16-OUT` (`5ae2c100f`) removed the `dense_model` term and `TryTritonPackedDecode` accepts `Hv=32` and dispatches `gdn_decode_h32_default` (`cuda_gdn.cu:5207`, `:5239`). The MoE loader never building `in_proj_ba` (#1169) was ONE of the terms still keeping the 35B off the packed leg; `GDN-MOE-PACKED-BA` owns that and builds the owner on every MoE safetensors checkpoint. The NVFP4 35B default arm is still excluded by the #365 fp8-tower term (`VT_GDN_PACKED_DECODE_FP8_TOWER`, default OFF; with `VT_GDN_FP8_IN_BF16` default OFF the predicted `mixed_qkv` stays F32 too), so it reaches the packed leg only under those two levers; the bf16 35B (`Qwen/Qwen3.6-35B-A3B`) default arm is the one that reaches it. Flip gates ALL EIGHT PASS exit 0 (root `~/work/vllm.cpp-gdn-decode-triton-flip` `gates.verdict`, `-DVLLM_CPP_TRITON=ON` + CUTLASS-4.5.0/nvcc-13.0, CUTLASS/FA2 configure-log lines verified, one flock): 27B DEFAULT (Triton) 235/235 + `=0` rollback 235/235; 35B DEFAULT 315/315 + `=0` rollback 315/315 (inert); AOT op test 28/28 (default fires cubin, `=0` fires legacy); full GDN 49/49 (2,343/2,343); oracle boundary 12/12; memcheck 28/28, 0 errors. No new A/B (9dd7d3f's +5.48 tok/s / −1.26 ms TPOT stands). The next binding grid runs the Triton decode path by default; no separate flip speed credit is claimed.	packed decode	`DONE`	`e47b4d6`
`KERNEL-GDN-AOT-BF16`	GDN prefill/decode and BF16-output Triton/CuTe fast-path coverage	backend resolution `qwen_gdn_linear_attn.py:150-411`; FLA kernels `fla/ops/chunk_delta_h.py:43-371`; FlashInfer GDN entry `flashinfer-ref/gdn_prefill.py:37-87`	canonical AOT contract, BF16 manifest bases, typed dispatch, `GdnOutDType()` — BF16 default on EVERY arm, `VT_GDN_OUT_BF16=0` the f32 rollback	Existing AOT/safety/native gates are green. The BF16/f32 component is 1.007989×, 16/20 timing and 2/4 memory. Immutable `3f256ab` is the exact v0.25 denominator: c16 total throughput passes at 1.027889×, but mean TPOT/ITL remains 0.987450× normalized; c1-c8 total and decode-shaped axes fail. Diagnostic cross-profiler totals are local `GdnDecodeFusedKernel` 19.101 s/73,578 calls and vLLM fused recurrent 28.659 s/70,848 calls; they are not a valid speed ratio. The independently gateable pure-decode mismatch is now owned by `KERNEL-GDN-PACKED-DECODE`; this row retains prefill AOT, BF16 default and scratch/teardown scope. The 35B GDN recurrence output and z gate are no longer f32 — `GDN-MOE-BF16-OUT` (#1168) removed the model-shape branch from `GdnOutDType`, so `VT_GDN_OUT_BF16` is the whole decision on both arms; this row's own prefill AOT and scratch scope is unchanged by that. Strict component/pool/teardown remain open	inventory	`READY`	-
`GDN-MOE-BF16-OUT`	The GDN recurrence output `dcore` `[T,Hv,Dv]`, the `z` output gate and the gated-RMSNorm weight that has to match them follow the ONE model dtype on every arm, dense and MoE alike, instead of resolving bf16 from a dense checkpoint and f32 from a MoE one	vLLM branches on no model shape here: `core_attn_out = torch.zeros(..., dtype=hidden_states.dtype)` `qwen_gdn_linear_attn.py:870-873`, `z` a split of the bf16 `in_proj_qkvz` `:843`, `:859-860`, `RMSNormGated` with no dtype override `:459-465`, one shared `Qwen3_5ForCausalLMBase` `models/qwen3_5.py:280-297` @ `5559679`; `VLLM_ENABLE_FLA_PACKED_RECURRENT_DECODE` defaults True with no shape term `envs.py:124`. SGLang agrees (`qwen3_5.py:522-536`, `gdn_triton.py:43` @ `f63458b5be`)	resolver + parser qwen3_5.cpp, contract qwen3_5_internal.h; `detail::ShouldUsePackedGdnDecode` lost its `e.dense_model` term in the same change	CPU tier only. `test_qwen35_paged_forward` enters through `ModelRegistry::Forward` on a MoE config and reads the dtypes off `dcore` and the projected gate, and runs a second time under `VT_GDN_OUT_BF16=0` so the documented rollback is gated too; `test_qwen27_paged_forward` pins the parser truth table and the shape-free eligibility. Nothing is measured. The 35B correctness gate, `test_qwen36_paged_engine` 315/315, `test_qwen27_paged_engine` 235/235 inertness, the `VT_GDN_OUT_BF16` 0-versus-1 same-binary A/B and the `nsys` memory-format confirmation are all `PENDING` on a GPU host and are listed in the spec's `## Now`	gdn-moe-bf16-out.md	`GATING`	#1168
`GDN-MOE-PACKED-BA`	The MoE safetensors loader builds the ONE merged `in_proj_ba` owner (`[2*Hv,H]` nk, rows `[b; a]`) that only the dense loader built, so `has_packed_ba` is true on the 35B and packed GDN decode — the vendored FLA `gdn_decode_h32` cubin — becomes reachable on a MoE checkpoint for the first time: on the bf16 35B (`Qwen/Qwen3.6-35B-A3B`) at the default, and on the NVFP4 35B only under `VT_GDN_PACKED_DECODE_FP8_TOWER=1 VT_GDN_FP8_IN_BF16=1`, because the #365 fp8-tower term (default OFF) still excludes its native-FP8 default arm. Loader-only: the consumer `ProjectGdnBA` and the eligibility predicate are already general	one `MergedColumnParallelLinear` per GDN layer on BOTH arms: `packed_modules_mapping["in_proj_ba"] = ["in_proj_b", "in_proj_a"]` on the shared `Qwen3_5ForCausalLMBase` `models/qwen3_5.py:281-297`, shard order `b=0, a=1` `:217-218`, `create_ba_proj` `output_sizes=[num_v_heads]*2` `gdn/qwen_gdn_linear_attn.py:516-530`, invoked once per forward arm `:802,:844,:905,:945` @ `5559679`	`LoadGdn` qwen3_5_weights.cpp (the one edit), dense precedent qwen3_5_dense_weights.cpp, consumer `ProjectGdnBA` + `detail::ShouldUsePackedGdnDecode` qwen3_5.cpp	red-first hermetic loader test through `LoadQwen3_5MoeLayer`; CPU split-vs-merged bit-identity in `test_qwen35_paged_forward`; GPU `test_qwen36_paged_engine` 315/315 on the NVFP4 35B (counters `0/0` at the default, `packed_launches > 0` under the two levers), bf16 35B default-arm counters and arm-vs-`VT_GDN_PACKED_DECODE=0` greedy identity, and the same-binary TPOT A/Bs — see the spec's `## Gates`	gdn-moe-packed-ba.md	`ACTIVE`	#1169
`KERNEL-GDN-SCRATCH`	Stream-ordered GDN AOT scratch allocation, reuse, and capture safety	FLA/FlashInfer runtime path `qwen_gdn_linear_attn.py:1073-1161`; CuTe Blackwell scheduler `flashinfer-ref/gdn_kernels/blackwell/gated_delta_net_chunked.py:177-182,379-846`	stream-owned pool, queue cleanup, test poison hook	same-stream poison/reuse/growth/pool-off assertions; all 11 buffers are filled with `0xff` before reuse and dispatch counters reject hand fallback; integrated current-main GB10 graph gate, sanitizer, trace and pool A/B pending	inventory	`READY`	-
`KERNEL-GDN-REPLAYSSM`	ReplaySSM buffered output-only decode for the GDN path: a per-slot ring of the last `L` steps' rank-1 factors `(d, k, g)`, with the full `[HV,V,K]` state written back only every `L` steps	Algorithm (vLLM, MIRROR SOURCE) @ `555967922`: kernel `layers/mamba/ops/selective_state_update_replayssm_output_only.py` (precompute `:22-129`, main `:131-472`, non-flush route `:275-279`, flush route `:358-362`, entry `:474-523`), launch heuristic `ops/replayssm_config.py`, ring dtypes/shapes `layers/mamba/mamba_utils.py:84-93,202-221`, layer wiring `mamba_mixer2.py:1105-1140`, cursor `v1/attention/backends/mamba_attn.py:77-82,575-620`, ring origin `v1/attention/backend.py:483-486`, `replayssm_buffer_len` default 16 `config/cache.py:148`, `use_replayssm` default `False` `:152`, refusals `config/vllm.py:2318-2337` (spec decode `:2329`), introduced `866fea2b` (vllm#48018). Mamba2 ONLY — four walls to GDN: `config/vllm.py:2318-2322` refuses a model without `supports_replayssm`; only `NemotronHForCausalLM` sets it (`models/nemotron_h.py:711,755,792`); `GDNAttentionMetadataBuilder` (`v1/attention/backends/gdn_attn.py:82`) does not subclass the builder that derives the cursor; the kernel asserts scalar-per-head `A` (`:540-542`) + Mamba2 `(B,C)` groups (`:528-530`). `gated_delta_net_state_dtype` (`mamba_utils.py:120-128`) appends no ring. Still true 877 commits past the pin. GDN application (SGLang, SECONDARY ORACLE) @ `f63458b5`: `layers/attention/fla/fused_recurrent_linear_replayssm.py` (`:50` credits vLLM, `:45-47` the `L=1` algebraic identity), `--enable-linear-replayssm` default `False` + `--linear-replayssm-cache-len` default 16 `server_args.py:1972-1987`, rings `mem_cache/memory_pool.py:465-487`, cursor `:641-645`, dispatch `linear/gdn_backend.py:317-327` + `linear/kernels/gdn_triton.py:84-119`, introduced `a10a24e9` (sglang#28451)	NOT IMPLEMENTED. `grep -rniE 'replayssm' src include tests` is empty. The path this would join: legacy packed decode cuda_gdn.cu:2324-2438 (state read `:2405`, write `:2437`), reg-tiled `:2516-2639` (default OFF), fused `:2805`, scan `:2221`, dispatch `:2680-2772`, vendored Triton AOT bridge `src/vt/cuda/cuda_gdn.cu::TryTritonPackedDecode` (fwd decl `:2673`, definition `:5201-5258`); state spec qwen3_5_common.cpp:37-106 (shape/dtype vectors `:84-87`), SSM dtype resolution `src/vllm/model_executor/models/qwen3_5.cpp::ResolveMambaSsmCacheDType` (`:544-554`), speculative arm `include/vt/ops.h::GdnSpecDecode` (`:2519-2522`)	NO MEASUREMENT EXISTS AND NONE IS CLAIMED. Source-verified anchors only, read at both pins. §5 of the spec DERIVES from the 27B shape (`HV=32, V=128, K=128, H_k=16`, fixture `tests/vllm/test_hf_config.cpp:80-82`): honest state-traffic ratio 0.61-0.63x at `L=16`, NOT SGLang's published 0.53x (`fla/bench_gdn_replayssm_decode.py:58-63`), which models `(1+1/L)/2` and omits both the flush re-read and every ring access; ring costs +18.9% mamba page (395,264 B against a 2,097,152 B state) so ~19% fewer GDN slots, against vLLM's ~7% on Nemotron shapes, because GDN's state is `VK` while the ring is `L(V+K)`; `L* = sqrt(4HVVK/(HVV + aH_kK))` = exactly 16 at this shape, matching both upstream defaults, with `L=8` and `L=32` both worse at 0.658x. Neither upstream is bit-exact and neither claims to be (vLLM fp32 `rtol=1e-4, atol=1e-3` / bf16 `6e-2, 2e-1`, `tests/kernels/mamba/test_replayssm_standard_decode_mamba2.py:47-57`; SGLang fp32 `atol=1e-4, rtol=1e-3`, `atol=2e-6, rtol=1e-5` at `L=1` alone, `test_linear_replayssm_decode.py:83-88,226-231` — atol/rtol roles are SWAPPED between the two). REGISTER PRESSURE (R1) IS RETIRED BY MEASUREMENT, 2026-08-18. Six kernels in one translation unit, `nvcc`/`ptxas` 13.0.88 for `sm_121a` on `orin:gpu0` under an `rc run` lease, `COMPILE_RC=0`, re-emitted identically across three `rc run` jobs — all on the same worker and toolkit, so this is compiler determinism on one host and NOT host or toolkit independence (evidence): in the shape the pin specifies — the state slice streamed over `nf_dstate_tile = 32` / `fl_dstate_tile = 64` with disjoint per-branch locals (`selective_state_update_replayssm_output_only.py:309-311,361,391-393,395`, `replayssm_config.py:47-52`) — the decode step compiles at REG:42 / STACK:0 (non-flush route, 15 of every 16 steps), REG:94 / STACK:0 (flush route) and REG:95 / STACK:0 (both fused), against a same-toolkit control transcribed from the shipped `GdnPackedDecodeRegTileKernel` (`cuda_gdn.cu:2516-2612`) at REG:255 + STACK:56. A first probe that held the `[BK=128]` row resident spilled at REG:255 + STACK:96, and that reading was WITHDRAWN: the mirror source deliberately does not hold the slice whole, and `nvdisasm -g` puts all 24 of that probe's spill LOADS inside the 1-in-16 flush branch. The remaining walls are R2 and the traffic model, not the register file: the 27B decode default is the vendored FLA cubin, which hard-guards `state.dtype == kF32` (`cuda_gdn.cu:5224`), and the hand register-tiled kernel built on the resident tile lost c16 700 vs 794 tok/s (`gdn_packed_decode_triton.h:9-14`, `cuda_gdn.cu:2722`) — whether a STREAMED kernel's exposed load latency costs more than the traffic it saves is UNMEASURED. The state stays `READY`: W0 is done and its answer is positive, but nothing was implemented, so no lifecycle transition happened and W1 is simply the next claimable item. A register count is not a throughput, latency or memory measurement, so no `docs/BENCHMARKS.md` entry is owed either. Motivation is the OPEN Qwen3.8-27B bf16 decode gap (c4 total 0.918x, output 0.963x, BENCHMARKS `:192-205`). Payoff is UNMEASURED; the spec §6 names the A/B (same binary, default-OFF lever, the SHIPPED cubin path as denominator, interleaved arms, clocks pinned under the GPU lock, `failed == 0` asserted, capacity as an axis)	gdn-replayssm	`READY`	-
`KERNEL-SSM-MAMBA`	General Mamba selective scan, causal convolution, SSD, and linear attention kernels	selective scan source `CMakeLists.txt:401`, `csrc/libtorch_stable/mamba/selective_scan_fwd.cu:394-820`; upstream mamba tests; decode conv-update FLA Triton `causal_conv1d.py:15-192` @ `e24d1b24`	only Qwen GDN-specific convolution/state path, including indexed cache gather/scatter cuda_gdn.cu:156 and GDN dispatch cuda_gdn.cu:3695; decode conv-update decode-fast CausalConv1dUpdateFastKernel + TryLaunchConvUpdateFast + flag header; prefill conv-fwd register-window CausalConv1dFwdRegKernel + LaunchConvFwdReg + fused post-conv per-V-head split GdnPostConvSplitKernel + fused post-conv fast (128-thread + 128-bit V copy) GdnPostConvFastKernel + GdnVecCopy8 + prefill flag header + CPU flag test	Indexed BF16/F32 cache-boundary + i8-mask CUDA tests test_ops_gdn.cpp:994, focused 7/7 memcheck and W1 structural trace pass; these GDN tests do not prove generic Mamba support. 2026-07-18 (`CLAIM-CONV-UPDATE-FAST-1`): decode conv-update kernel-efficiency sub-lever landed, DEFAULT ON (the c16-trace scan lever #5; row stays `INVENTORIED` — broad Mamba support unchanged). `CausalConv1dUpdateFastKernel<Tin,Tout,TState,WIDTH>` behind `VT_CONV_UPDATE_FAST` (default ON / `=0` rollback) is BIT-IDENTICAL (0-ulp) to the shipped `CausalConv1dUpdateKernel`: same float op order, only a 2D grid (removes two int64 div/mod per thread) + a WIDTH-templated register-cached state row reused for the conv accumulation AND the roll (mirrors FLA's per-`KERNEL_WIDTH` `col0..col3` register specialization). DGX-PROVEN (flock, clean `-Werror`, CUTLASS+FA2 hard-verified): `test_ops_gdn` conv-update decode-fast `fast==shipped` BYTE-EXACT (0-ulp) on `out`+`conv_state` 330/330 (k∈{3,4,5}, bf16+f32 state, ±bias, silu/identity, compact+scattered cache incl. NULL-block); full GDN 51/51; CPU flag test 10/10. Isolated nsys pure-kernel 27B c16 shape (batch=16 conv_dim=10240 k=4 bf16): shipped 7,072 ns vs fast 3,680 ns median = 1.92× (clears ≥1.3× bar). Full default set (all fast kernels ON) 27B 235/235 + 35B 315/315; `=0` rollback arms 235/235 + 315/315. `benchmark_binding=false`; binding grid re-measures. 2026-07-18 (`CLAIM-GDN-PREFILL-CONV-1`): PREFILL conv-fwd + fused post-conv kernel-efficiency sub-lever landed (row stays `INVENTORIED`). `CausalConv1dFwdRegKernel` (`VT_CONV_REG` DEFAULT ON / `=0`→tiled) is the register-resident sliding-window mirror of vLLM FLA `_causal_conv1d_fwd_kernel` (`causal_conv1d.py:397-452`): per-channel weights preloaded to registers, `(k-1)`-tap register sliding window (each x loaded once, coalesced), token-axis chunked over grid.z for `n<=4` low-batch prefill. `GdnPostConvSplitKernel` (`VT_GDN_POSTCONV_SPLIT` OPT-IN) mirrors `_fused_post_conv_kernel` grid `(cdiv(L,BLOCK_T), H+HV)` (per-V-head). BIT-IDENTICAL (0-ulp) to the shipped tiled/megablock kernels — same tap-order f32 sum + q/k L2-norm byte-for-byte. DGX-PROVEN: byte-exact reg==tiled + split==megablock 268 GPU assertions + full GDN 3081/3081; memcheck 0; 27B 235/235 + 35B 315/315 (final defaults reg ON/split OFF). nsys 35B (`~/work/prefill-attr-conv-35b`): conv c1 −4.7% / c6 −7.3% (bandwidth-bound) ⇒ DEFAULT ON; post-conv split c1 −3.8% / c6 +4.7% (near-neutral — GdnPostConv is q/k-L2-norm-dominated) ⇒ OPT-IN. TTFT c1 −0.39% within run-noise (conv ~2.5% of GPU). Residual vLLM conv gap = bf16 traffic (VT_GDN_IN_BF16, task #40 sibling), not structure. `benchmark_binding=false`. 2026-07-19 (`CLAIM-GDN-POSTCONV-FAST-1`): fused post-conv FAST sub-lever landed, DEFAULT ON (row stays `INVENTORIED`). A fresh production-path nsys (`--cuda-graph-trace=node`, VLLM_CPP_TRITON=ON, input-1024) confirmed the GDN chunk compute runs the vendored FLA Triton AOT cubins by default (`chunk_gated_delta_rule_fwd_kernel_h_blockdim64` / `chunk_fwd_kernel_o` / `recompute_w_u_fwd_kernel` / `chunk_scaled_dot_kkt_fwd_kernel` = at FLA parity), so the #1 remaining NON-AOT GDN kernel on BOTH models is the fused post-conv prep. The split (above) measured neutral/slower, so `GdnPostConvFastKernel` (`VT_GDN_POSTCONV_FAST` DEFAULT ON / `=0`→megablock) keeps the megablock grid `(T,Hk+1)` but makes two BYTE-IDENTICAL changes for Dk==Dv==128: 128 threads/block (the 128-wide L2-norm tree is the 256-wide tree minus a leading `+0` step; better reduction occupancy) + a 128-bit-staged V copy (`GdnVecCopy8`: raw `int4` when dtypes match, else the SAME `__bfloat162float`/`__float2bfloat16` converts). No arithmetic reordered; grounded in FLA `_fused_post_conv_kernel` (`fused_gdn_prefill_post_conv.py:57-149`). Isolated nsys per-call: 27B 175.6→133.0 ms (−24.3%), 35B 93.5→70.3 ms (−24.8%); in-situ TTFT (input-1024, 3 reps): 27B c1 −1.14% / c2 −1.31%, 35B c1 −0.72% / c2 −0.99% (all reps positive). BIT-EXACT: 27B 235/235 + 35B 315/315 on default(fast)+`=0`(megablock) arms; CPU flag test 28/28; clean `-Werror`. Evidence `dgx:~/work/vllm.cpp-gdn-chunk/prof`. `benchmark_binding=false`. 2026-08-12 (#496): the SSD half of this row's title is named and specced. Everything above is the GDN arm; the generic Mamba2 selective scan has no local symbol at all — `include/vt/ops.h` has `kCausalConv1dFwd/Update/SpecUpdate`, `kRmsNormGated`, `kGdnPrefill/Decode`, `kKdaGatedDeltaRule` and no chunked state-space scan and no selective state update. Gated delta rule is not Mamba2: GDN carries the `(I − β kkᵀ)` removal term and a per-head scalar decay, Mamba2's SSD is a diagonally-decayed gated linear recurrence driven by `A_log`/`dt`/`D` with `B`/`C` shared across `n_groups`. Owed: `mamba_chunk_scan_combined_varlen` (`mamba/ops/ssd_combined.py:157-235` @ `555967922`), `selective_state_update` (`mamba/ops/mamba_ssm.py:497`), and the silu-gated GROUP RMS norm (`mamba_mixer2.py:100-149`, a sibling of our sigmoid `kRmsNormGated`, not a parameter of it). Blocks 10 `INVENTORIED` model rows; immediate driver is `NemotronHForCausalLM` (Nemotron-3.5-Lightning-30B-A3B-NVFP4: 23 mamba / 6 attention / 23 MoE, `n_groups=8`, `ssm_state_size=128`, `chunk_size=128`). 2026-08-12 (#496): W1 LANDED at `47960a009` — CPU host references `vt::Mamba2ChunkScan` (varlen chunked SSD prefill), `vt::Mamba2StateUpdate` (single-token selective update) and `vt::RmsNormGatedGroup` (silu-gated GROUP RMS norm), gated against a sequential per-token recurrence transcribed independently in `double` from upstream's OWN non-chunked CPU kernel (`csrc/cpu/mamba_kernels.hpp`), plus chunk-size invariance over {8,16,32,64,128}. Two fresh reviews: round 1 FAIL (the gated-norm `weight` was read through `Ptr<float>()` while the validator accepted bf16/f16 — a 2x heap over-read on the dtype upstream actually passes, since `Mixer2RMSNormGated.weight` is `torch.ones(...)` at model dtype; plus four unpinned guarantees and a refuted "equivalent mutant"), round 2 PASS (13 mutations, F1 swept repo-wide for a second instance — none). Row stays `INVENTORIED`: a host reference is not generic Mamba support. Owed: the CUDA arm (W2), the `MambaSpec` producer for Mamba2 shapes (W3), and one decode refusal `SUBCASE` (spec §8.2). No performance claim is made	Mamba2 SSD spec; prefill conv spec; conv-update decode-fast spec; inventory	`INVENTORIED`	-
`KERNEL-SAMPLING`	Sampler reductions, top-k/top-p, probabilities, penalties, masks, RNG	stable sampler/top-k `CMakeLists.txt:399-400`; upstream `tests/v1/sample/test_sampler.py`, `test_topk_topp_sampler.py`	cuda_sample.cu:638	sampling tests	inventory	`ANCHOR-BACKFILL`	-
`KERNEL-COLLECTIVES`	Custom/quick/FlashInfer all-reduce and fused collective kernels	custom all-reduce `CMakeLists.txt:405`; runtime `device_communicators/custom_all_reduce.py:109-264`; upstream distributed tests	-	-	inventory	`INVENTORIED`	-
`KERNEL-SPEC-DECODE`	Speculative-decode expand/copy, rejection, acceptance, and draft sampling kernels	Triton setup kernels `vllm/v1/spec_decode/utils.py:306-565`; proposer path `llm_base_proposer.py:67-1315`	-	-	inventory	`INVENTORIED`	-
`KERNEL-GEMM-CPU-ELEM-X86WIDE`	AVX2 + AVX-512 elementwise CPU GEMM tiers. Our x86 tier is SSE2 128-bit `MR=2` (cpu_matmul_elem.cpp:226-231) while the hardware offers `avx2`/`avx512f`/`avx512_bf16`/`fma`. Measured 2026-08-06: that width gap alone is 3.5x (vt 166.6 vs ggml-no-llamafile 587.0 GFLOP/s at `131,2048,512`, f16 weight + f32 acts). Sibling of quant row G5: G5 owns the QUANTIZED x86 kernels, this owns the ELEMENTWISE ones, which nothing covered before	llama.cpp `ggml/src/CMakeLists.txt:371-401` (`GGML_CPU_ALL_VARIANTS` per-ISA build), `ggml-cpu/arch/x86/quants.c` @ `237ad9b96`	planned: new `src/vt/cpu/cpu_matmul_elem_avx2.cpp` + `_avx512.cpp`, per-file `COMPILE_OPTIONS` next to the i8mm block (CMakeLists.txt:744-760), runtime `__builtin_cpu_supports` probe in `BuildTier()`	planned: widen the `test_ops_matmul_elem.cpp` tier sweep (memcmp byte-identity, unchanged bar)	specs/cpu-elem-gemm-wide-isa-and-tiling.md	`SPIKE`	`CLAIM-KERNEL-CPU-ELEM-WIDE-1`
`KERNEL-GEMM-CPU-TILED`	tinyBLAS-style M x N register-tiled sgemm, all architectures. Isolated 2026-08-06 by rebuilding ggml with `GGML_LLAMAFILE=OFF`: our NEON kernel is at parity with ggml's stock kernel (216-242 vs 208-215 GFLOP/s, ahead on 4 of 6 shapes), so the entire Arm 16-bit deficit is llamafile's tiled GEMM, worth ~1.9x on f16 and ~1.2x on f32, and ~2.4x on x86. Not a defect in `KERNEL-GEMM-CPU-ELEM`, an absent capability. That ggml denominator is SUPERSEDED: it came from our own local-only fork `237ad9b96`, built from the same uncommitted working tree as the Vulkan arm, so the tree is in no repository. None of the 65 fork commits touch `llamafile/sgemm.cpp`, so the compared kernel matches `b9827`, but the parity verdict is still owed a re-take against stock `b10451` (#1003 row 11), because it is what attributes the Arm 16-bit deficit to an absent capability rather than to this kernel	llama.cpp `ggml/src/ggml-cpu/llamafile/sgemm.cpp` @ `237ad9b96`, a superseded location owed re-anchoring under #1003 (tiles M and N, loops K, which is what makes it compatible with our byte-identity rule)	planned: raise `MR` and add an `NR` dimension to the `btm` family per tier; no new op, no new seam	planned: existing `test_ops_matmul_elem.cpp` battery, which already gates memcmp byte-identity	specs/cpu-elem-gemm-wide-isa-and-tiling.md	`SPIKE` (RESCOPED 2026-08-06 after a deep dive, byte-identity RETAINED by decision. The earlier REFUTED verdict was too strong: it held for llamafile's STRUCTURE, not for the gap. Proven byte-identical and reachable: (a) the `nk` [K,N] family needs NO transpose and is exactly byte-identical to `bt` at every shape, worth ~1.15x; (b) `cpu_ops.cpp:135` gives M blocking ONLY to `bt`, so `nk` runs mr=1 and reloads the weight per activation row, and a prototype shows adding it is worth 1.44x-1.78x. Plan: M-blocked `nk` kernel, then repack [N,K]->[K,N] at load like the q8_0 G7 tier. Genuinely unreachable and NOT pursued: FMA and K-vectorised hsum, the only two real costs of the contract. Superseded verdict, see `.agents/benchmark-record.md`: llamafile's 1.9x comes from vectorising along K + hsum + FMA, all order-changing, so it is unreachable while byte-identity holds; raising `kMrNeon` 4->6 measured a WASH, +5% on three shapes and -7 to -8% on two. Needs a PRODUCT call: rescope, or add an opt-in non-byte-identical fast path with its own NMSE gate)	`CLAIM-KERNEL-CPU-ELEM-WIDE-1`
`KERNEL-CPU-CONV2D-SUBSAMPLE`	torch `nn.Conv2d` as a real `vt::` op (groups/stride/padding/dilation), the conformer audio-encoder front end. Before this the ONLY Conv2d in the tree was a host `std::vector<float>` loop, gemma4_audio.cpp:92 `Conv2dK3S2P1`, which stays as an independent correctness reference. One op serves all three `ParakeetEncoderSubsamplingConv2D` stages (dense 1->C, depthwise groups=C, pointwise 1x1) and vLLM's native `Conv2dSubsampling`	transformers 5.3.0 `transformers/models/parakeet/modeling_parakeet.py:357` (:369-371, :377-386, :388), which is what vLLM RUNS (`vllm/model_executor/models/parakeet.py:37,62`); vLLM-native sibling `vllm/model_executor/models/conformer_encoder.py:18`	`vt::Conv2d` / `OpId::kConv2d` (declaration + contract include/vt/ops.h:2077-2099, validating wrapper src/vt/ops.cpp:2226-2260, registration CMakeLists.txt:773-778); CPU kernel src/vt/cpu/cpu_conv2d.cpp:65-107	tests/vt/test_ops_conv2d.cpp:1-364 (wired tests/CMakeLists.txt:891-901) — 4 cases / 1631 assertions, memcmp byte-identity vs an INDEPENDENT in-test scalar reference across dtype x shape x thread count; mutation-verified (group map, stride/dilation axis, bias placement all killed)	specs/parakeet-conformer-encoder.md (row P1)	`ACTIVE` (CPU tier LANDED + byte-identity gated; the CUDA provider is the remaining work)	`CLAIM-PARAKEET-KERNELS-P1P3`
`KERNEL-DEPTHWISE-CONV1D`	NON-CAUSAL depthwise `nn.Conv1d(C,C,K,groups=C)`, the conformer convolution module's temporal mixer: centre-padded, stateless, activation-free, strided/dilatable. A deliberate SIBLING of `vt::CausalConv1dFwd` (Mamba/GDN: causal, persistent `conv_state`, folded SiLU), which is NOT modified — widening the causal op would have put a branch in a hot decode kernel and risked its byte-exactness	`modeling_parakeet.py:116` (padding :136, ctor :138-146, applied :180); vLLM-native sibling `conformer_encoder.py:229`	`vt::DepthwiseConv1d` / `OpId::kDepthwiseConv1d` (include/vt/ops.h:2101-2121, wrapper src/vt/ops.cpp:2262-2296); CPU kernel src/vt/cpu/cpu_conv1d_depthwise.cpp:63-95	tests/vt/test_ops_conv1d_depthwise.cpp:1-338 — 5 cases / 1184 assertions, same byte-identity bar, both accepted weight layouts, plus a left-only-padding cross-check that the causal op reproduces element for element	specs/parakeet-conformer-encoder.md (row P2)	`ACTIVE` (CPU tier LANDED + byte-identity gated; the CUDA provider is the remaining work)	`CLAIM-PARAKEET-KERNELS-P1P3`
`KERNEL-ATTN-RELPOS`	Transformer-XL relative-position ENCODER self-attention — no KV cache, no paging, no RoPE, non-causal. Every other attention path in `vt` is a decoder path (RoPE + paged/flash KV, `ops.h` kAttention/kPagedAttention/kMla*), so none could express it. The upstream `_rel_shift` pad/reshape/slice is carried as the CLOSED FORM `raw(i, T-1-i+j)`, so no `[T,2T-1]` scratch is materialised; the unit test's reference does the literal reshape, which is what proves it. The two upstreams differ only in where the scale lands, exposed as `AttentionRelPosArgs::scale_after_sum` rather than chosen	`modeling_parakeet.py:259` (forward :302-347, `_rel_shift` :349-355, `eager_attention_forward` :225-255); vLLM-native sibling `conformer_encoder.py:170` (:188-217, :179-186)	`vt::AttentionRelPos` / `OpId::kAttentionRelPos` (include/vt/ops.h:2123-2165, wrapper src/vt/ops.cpp:2298-2345); CPU kernel src/vt/cpu/cpu_attn_relpos.cpp:77-158	tests/vt/test_ops_attn_relpos.cpp:1-450 — 7 cases / 368 assertions, byte-identity vs a LITERAL-`_rel_shift` reference, GQA ratios, padding mask, both scale placements, thread counts	specs/parakeet-conformer-encoder.md (row P3)	`ACTIVE` (CPU tier LANDED + byte-identity gated; the CUDA provider is the remaining work)	`CLAIM-PARAKEET-KERNELS-P1P3`
`KERNEL-CONV3D`	General 3-D convolution as a real `vt::` op — the primitive `vt` had on NO device, and the reason the LTX-2.5 video VAE decode ran entirely on the host while every oracle runs it GPU-resident. Not a parameter of `kConv2d`: the ACCUMULATION ORDER differs and is part of the contract (one f32 partial per INPUT CHANNEL, bias seeded FIRST, vs `kConv2d`'s single flat accumulator with the bias last), the same sibling relationship `kConv1d` has with `kDepthwiseConv1d`. Batch 1 and a merged weight leading axis, because `vt::Tensor` caps rank at 4 (tensor.h:12) — both refused BY NAME rather than folded silently. Four other host conv3d loops (`minimax_h3_vae_cnn.cpp`, `minimax_h3_video_vae.cpp`, `ltx2_upsampler.cpp`, `clip_mmproj_gguf.cpp`) are owed it and are deliberately NOT rewired	torch `nn.Conv3d` as `CausalConv3d` instantiates it — Lightricks/LTX-2 @ `fd4ded7f2` `packages/ltx-core/src/ltx_core/model/video_vae/convolution.py:292-302`, called once at `:312`; placement is a BUILD-time decision upstream (`loader/single_gpu_model_builder.py:267-288`, CUDA by default at `:273`; `devices.py:29-39` resolves CUDA -> MPS -> CPU)	`vt::Conv3d` / `OpId::kConv3d` (contract include/vt/ops.h, validating wrapper src/vt/ops.cpp); CPU kernel src/vt/cpu/cpu_conv3d.cpp; CUDA kernel src/vt/cuda/cuda_conv3d.cu; reached from the render path through `Ltx2ConvVideoDecode` on the queue `Ltx2VideoEngine::Load` resolved	tests/vt/test_ops_conv3d.cpp — 4 cases / 2035 assertions, `memcmp` byte-identity vs an INDEPENDENT in-test scalar reference across dtype x shape x thread count, plus the case that asserts the per-input-channel partial is a DIFFERENT number from the flat accumulator. Reachability: tests/vllm/models/test_ltx2_vae.cpp counts the dispatches through `GetOpProviderStats` entered at `Ltx2VideoDecodeStreaming` (44 cases / 3131 assertions, no golden re-baselined); deleting the dispatch reds 12 of them. Device arm: tests/vllm/multimodal/test_diffusion_device_seam.cpp runs upload + a non-`kCPU` dispatch + download against the fake unified-memory XPU backend and requires `memcmp`-identical pixels. The CUDA arm has never been compiled or executed (#1452)	specs/ltx25-device-residency.md `## W5`	`ACTIVE` (CPU tier LANDED + byte-identity gated and REACHED from the render path; the CUDA arm is registered and unexecuted, #1452, and the decode's non-convolution stages are still host-side, #1451)	`CLAIM-LTX25-DEVICE-RESIDENCY-W5`

2026-08-07 KERNEL-SSM-MAMBA checkpoint. Exact (sequence, 8-token chunk) prefill-conv metadata/dispatch is default ON on sm_120 after byte-exact gates. Rebased-main same-binary causal-conv is 718.704→233.955 ms (3.072x) and the profiled enclosing workload improves 2.272%; pinned vLLM remains 145.421 ms (1.609x open). Lifecycle stays INVENTORIED because generic Mamba coverage and the 27B/35B release gates are unchanged. Spec and evidence.

2026-08-08 KERNEL-SSM-MAMBA checkpoint. The opt-in VT_GDN_POSTCONV_TOKEN_TILE path ports vLLM/FLA's 16-token, per-head, four-warp post-conv schedule for 128-wide Q/K/V heads. It preserves the current 128-lane reduction tree exactly while retaining Q/K values in registers. Qwen3.5-4B sm_120 graph-node time is 227.887→122.587 ms (1.859x) with byte-identical production tokens and every enclosing axis positive; the same-tool vLLM residual is 1.135x. Portable 6/6, CUDA GDN 67/67·4384 and cached 4B 3/3·1672 pass. It remains opt-in pending repeated and 27B/35B gates; generic Mamba lifecycle stays INVENTORIED. Spike and result.

2026-08-08 KERNEL-SSM-MAMBA causal-conv residual checkpoint. The opt-in width-four, one-channel arm removes runtime-width work at the unchanged 64-block feature grid and is byte-exact. Qwen3.5-4B sm_120 graph-node time is 234.605→219.506 ms (-6.44%); dominant waves improve about 149.5→140.1 us and every observed enclosing axis is slightly positive. CUPTI registers rise 43→52, so the win is not occupancy. The separately measured 256-channel arm halves the grid but rises to 58 registers and is 4.05% slower than the K4 one-channel arm; that hypothesis is falsified. Portable 9/9·88, CUDA GDN 67/67·4631 and paged-forward 4/4·8 pass. Arm 1 remains opt-in; generic Mamba lifecycle stays INVENTORIED. Spike and result. 2026-08-09 KERNEL-SSM-MAMBA sm_120 campaign anchor. #206 tracks the RTX 5070 Ti Qwen3.5-4B Pareto campaign. Reviewed atomic pretoken admission removed the frontend/batching confound; the corrected exact three-repetition comparison is 6831.71 vs 6643.40 tok/s (1.0283x), while TTFT, TPOT/ITL and E2E remain 1.0853x / 1.0165x / 1.0288x slower and VRAM is +118.7 MiB. The exact default-OFF GDN decode BV16+swizzle+REGSTATE stack improves local throughput, TTFT, TPOT and E2E, and slack-only memset adds a smaller further local win. Two deeper candidates are closed: geometric argmax scratch merely moved wait from cudaFree to stream synchronization and regressed TPOT ~1%; BF16 vector writeback improved y800 only 0.143% with one losing raw leg. Both products/tests were removed while their specs and same-tool traces remain. Lifecycle stays INVENTORIED; generic Mamba and 27B/35B coverage are unchanged. Campaign spec.

Count invariants

This table has exactly 38 practical kernel-family rows.
Baseline lifecycle counts are 8 ANCHOR-BACKFILL, 0 READY, 4 PARTIAL, 12 ACTIVE, 0 GATING, 1 DONE, and 13 INVENTORIED.
KERNEL-DFLASH2-GROUPED-CONV was added on 2026-08-19 (SPEC-DFLASH2 W2, #1314) and is the tenth ACTIVE row. It is a separate family from KERNEL-DEPTHWISE-CONV1D rather than a variant of it: that kernel's weights are static per channel and its mask is causal over the SEQUENCE, while this one's are projected per position from the sublayer input, grouped, and masked over the QUERY BLOCK. It was recorded as staying ACTIVE rather than DONE because its CUDA arm had never compiled (spec ## Owed O6). That reason expired on 2026-08-20: the arm compiled and ran on dgx:gpu0 and O6 is discharged (#1489). The state is left unchanged here because a lifecycle move owes docs/STATUS.md, docs/BENCHMARKS.md and the moved row spec's ## Now, and because whether this row is correctly ACTIVE at all is already #1469's open question.
KERNEL-DFLASH2-SELECTOR-EDGES and KERNEL-TOPK-PAIRS were added on 2026-08-20 (SPEC-DFLASH2 W3, #1314) and are the eleventh and twelfth ACTIVE rows. They are TWO rows and not one because they are two kernels with different shapes and different gates: one is a small dense contraction over two codebooks whose difficulty is the predecessor indexing and the bf16 rounding placement, the other is a sort-free selection over a 248320 vocabulary whose difficulty is the TIE-BREAK. KERNEL-TOPK-PAIRS is a separate family from the shipped sampling threshold search rather than a variant of it: that one masks in place and returns no indices, this one compacts, orders and emits pairs, and its CUDA arm extends that search in the same file rather than duplicating it. Both stay ACTIVE rather than DONE because neither CUDA arm has ever compiled (spec ## Owed O10).
KERNEL-DFLASH2-PATH-WALK was added on 2026-08-20 (SPEC-DFLASH2 W4, #1314) and is the thirteenth ACTIVE row. It is a separate family from KERNEL-DFLASH2-SELECTOR-EDGES rather than a second entry point into it, on the axis that decides kernel families here: the lattice op is a DENSE CONTRACTION with a reduction and an f32 envelope, and the walk is a SEQUENTIAL SELECTION with no arithmetic at all and a bit-exact cross-backend contract. Their grids differ for the same reason — one block per (request, step, predecessor slot) against one block per REQUEST with the step loop inside it, which is the whole point of spec ## Risks/decisions D3. It stays ACTIVE rather than DONE because its CUDA arm has never compiled on the authoring host (spec ## Owed O11).
The two counts above are the W2 baseline plus W3's two rows, and they now UNDERCOUNT BY ONE: KERNEL-CONV3D landed on main concurrently with W3 (#1007) and did not bump them, though it did bump the gated total in scripts/check-agent-record.py. That constant is the authority and is correct at 56; these two prose numbers are not gated by anything and re-deriving the whole file's lifecycle census is not W3's scope. Named here rather than silently corrected or silently left, because a reader comparing the two surfaces will otherwise think one of them is wrong about W3.
The distinct KERNEL-CPU-A76-Q8-DOT compiler/assembly family (GATING, PR #79) was added on 2026-08-06; scripts/check-agent-record.py pins the total row count.
The prefill/scratch GDN ACTIVE rows are valid only while CLAIM-PR3 remains active; the FA2 row is valid only while CLAIM-SERVE-GATE-1 claims it in .agents/coordination.md. KERNEL-GDN-PACKED-DECODE closed DONE (owner e47b4d6, W1D3 EQUIVALENCE), so its CLAIM-GDN-BA-ROUNDING-1 lead moved to the KERNEL-GEMM-BF16 qkvz row. Integration must reconcile a row if its claim closes first.
Adding a family requires an upstream/dependency anchor, tests-to-port entry, spike order, and updated counts in the same commit. Splitting a block preserves its original ID as an umbrella or records an explicit supersession mapping.