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.
| ID | Item | Upstream | Our code | Tests/evidence | Spike/spec | State | Owner |
|---|---|---|---|---|---|---|---|
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 |
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 OpIds 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. |
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) |
portable-fusion-framework.md | SPIKE (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 | ACTIVE |
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+KernelTmaWarpSpecialized*FP8FastAccum → 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 |
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 attn*sigmoid(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 |
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 | ACTIVE |
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 |
cuda_marlin_repack.cu:129, cuda_moe_marlin.cu:156 | dense tests, MoE tests; 35B gate | inventory | ANCHOR-BACKFILL |
- |
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) |
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 ((x*inv)*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 T*Hv (~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 | ACTIVE |
CLAIM-SERVE-GATE-1 |
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:271, 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 (nq=17,hq=32) × kBlock=256 threads looping SERIALLY over C__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-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 (iters-1)×[row-norm, col-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) |
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 |
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, d≤256): 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) |
multimodal-speed §14 + §16 | ACTIVE |
CLAIM-MM-SPEED-AUDIO-ENC-KERNEL + CLAIM-MM-SPEED-QWEN-IMAGE |
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 | ACTIVE |
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 |
- |
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, 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 — excluded at the model level by the dense-only ShouldUsePackedGdnDecode (qwen3_5.cpp:49) and the launcher guard rejects its Hv=32 shape (cmake H=32=35B) anyway (clean fallback; a 35B cubin would be dead code). 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, 27B-only default/f32 fallback | 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. Every 35B path stays f32; strict component/pool/teardown remain open |
inventory | ACTIVE |
CLAIM-PR3 |
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 | ACTIVE |
CLAIM-PR3 |
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 |
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 |
- |
- This table has exactly 32 practical kernel-family rows.
- Baseline lifecycle counts are 8
ANCHOR-BACKFILL, 0READY, 4PARTIAL, 6ACTIVE, 0GATING, 1DONE, and 13INVENTORIED. - The prefill/scratch GDN
ACTIVErows are valid only whileCLAIM-PR3remains active; the FA2 row is valid only whileCLAIM-SERVE-GATE-1claims it in.agents/coordination.md.KERNEL-GDN-PACKED-DECODEclosedDONE(ownere47b4d6, W1D3 EQUIVALENCE), so itsCLAIM-GDN-BA-ROUNDING-1lead moved to theKERNEL-GEMM-BF16qkvz 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.