mudler · localai-bot · Aug 6, 2026
diff --git a/.agents/NOW.md b/.agents/NOW.md
@@ -16,15 +16,15 @@ checkpoint on `upstream/main` at `59674cf1d`.
 |---|---|---|
 | Laguna NVFP4 decode speed | **Closed: PARITY+ 1.03x** (byte-exact, default; `VT_LAGUNA_RESIDENT_BF16W` bf16-residency). Benchmark record | vLLM K-run set when convenient |
 | DeepSeek-V4-Flash decode | **Closed: BEATS ds4 1.144x** (`VT_V4_RESIDENT_W`, byte-exact). Phase-2 routed-expert residency NEGATIVE (−3.4%), default-OFF | — |
-| f32-out GEMV audit | Only laguna + deepseek_v4 bf16 tower affected; gate models & on-framework dense unaffected (bf16-out, e2e-verified) | Re-verify deepseek_v4 bf16 tower same-tool |
+| f32-out GEMV audit | Only laguna + deepseek_v4 bf16 tower affected; gate/on-framework dense unaffected | Re-verify deepseek_v4 tower same-tool |
 | Invocation-parity prevention | CI guard (`check-gemv-invocation-consistency.py`) + AGENTS.md checklist landing | Review + merge; CUDA build-verify `kGemvHeuristicAlgos` on dgx |
 | MiniMax-H3 lane | Portable path complete; e2e prompt-conditioned video on real weights (Thor). Speed = NVFP4 FP4 device path, sm_121-gated | PR #26 rebase + supports-audit synthesis |
 | Kimi-Linear-48B (KDA+NoPE-MLA+MoE) | **Full-model GB10 e2e RUNS** (bf16-resident §13): CPU+CUDA 13/13·656, no OOM. **Token gate NEAR-TIE 106/128** (6/8 token-exact) | device GDN/MLA islands + bf16 stream; 1.59 tok/s; default OFF |
 | 35B fresh grid | **BOUND** @`1ea26427`: tput 0.93-1.03x, c16 0.93x. INTAKE + Option A both **RESOLVED NEGATIVE** (H2D-out-of-capture tput WASH) | Real lever left: prefill glue (task #61) |
 | Qwen3.5-4B revalidation | 0.9971x @`59674cf1` (#35); TTFT/PSS pass, TPOT/ITL open | `docs/bench-evidence/` |
-| MXFP4 parity (Qwen3-8B) | **`MARLIN-STRUCT`: decode-graph + gate_up FUSION default-ON (marlin 180→144 GEMM/step = vLLM-structural); #44 3/3, 0.6B/4B 184/184, 32B-NVFP4A16 142/142** | residual = marlin CTA + flash |
+| MXFP4 parity (Qwen3-8B) | **`MARLIN-STRUCT`: decode-graph + gate_up FUSION default-ON (180→144 GEMM/step); #44 3/3, 0.6B/4B 184/184, 32B 142/142** | residual = marlin CTA + flash |
 | ROW-SERVE-ASYNC-DENSE-MIRROR | **LANDED+dgx-VERIFIED** (`f9c969ae`): #31 async mirror on classic dense Qwen3; gate RED→GREEN, SACRED 184/184 | Residual: sibling scope one-liner |
-| MXFP4 parity goal | graph+fuse default-ON. c8 residual: marlin CTA 144 vs 48 = DOMINANT +1,177us (`VT_MARLIN_E1_PAR1` opt-in → near-parity, but flips strict 32B token → default-OFF), flash +784, glue +195 | NEXT: dense-template marlin port + full binding (oracle) |
+| MXFP4 parity goal | graph+fuse default-ON; c8 dominant residual = marlin CTA 144 vs 48. **`KERNEL-MARLIN-DENSE-PORT` gated-OFF** (`VT_MARLIN_DENSE`): vLLM's own dense marlin = byte-preserving E=1; 3 dense `.cu` compile-clean dgx; unit WRITTEN | NEXT (dgx): strict dense-ON vs oracle + nsys + binding (state) |
 
 In-flight branches (default-OFF, not pushed): `laguna-fp4proj-prod` (fp4),
 laguna bf16/legacy/pipeline-gemv, `ds4-hc-expand-fuse`. Records:

diff --git a/.agents/kernel-matrix.md b/.agents/kernel-matrix.md
@@ -118,7 +118,7 @@ host/sched. Detail: state `KERNEL-FA2-GQA-SWAP-FLIP`.
 | `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](../src/vllm/model_executor/models/qwen3_5_dense_weights.cpp#L165), merged/fallback qkvz dispatch [forward:2372](../src/vllm/model_executor/models/qwen3_5.cpp#L2372), eligibility seam [internal:52](../src/vllm/model_executor/models/qwen3_5_internal.h#L52), stride-aware [consumers:930](../src/vt/ops.cpp#L930), unchanged [cuBLASLt path:188](../src/vt/cuda/cuda_matmul.cu#L188), exact [oracle generator:1](../tools/bench/gdn_ba_projection_oracle.py#L1), and explicit [trace contracts:113](../tools/bench/online_gate.py#L113) | 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](../tests/vt/test_ops_gdn.cpp#L2367), [loader/CPU-exactness:456](../tests/vllm/models/test_qwen27_dense_forward.cpp#L456), [eligibility:451](../tests/vllm/models/test_qwen27_paged_forward.cpp#L451), 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](../tests/vllm/models/test_qwen27_paged_forward.cpp#L491), 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](specs/gdn-merged-input-projections.md); [packed decode](specs/gdn-packed-decode.md) | `ANCHOR-BACKFILL` | CLAIM-GDN-BA-ROUNDING-1 |
 | `KERNEL-GEMM-FP8` | FP8/INT8 scaled-mm plus static activation quant | C2x/C3x dispatch `CMakeLists.txt:737-863`; stable quant sources `:383-388`; vLLM cuBLASLt fp8 reuses an in-graph plan (nvjet_sm121_qqtst_* kernels) — no per-call heuristic | [cuda_matmul_fp8_cutlass.cu:312](../src/vt/cuda/cuda_matmul_fp8_cutlass.cu#L312), [cuda_matmul.cu:345](../src/vt/cuda/cuda_matmul.cu#L345); **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](../src/vt/cuda/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](../src/vt/cuda/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](../src/vt/cuda/fp8_plan_cache.h) + [GetOrBuildCachedFp8Plan/BuildFp8Plan](../src/vt/cuda/cuda_matmul.cu#L358) (`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](../tests/vt/test_ops_fp8_cutlass.cpp#L188); **byte-exact cached==fresh** [test_ops_fp8_cutlass.cpp#L387](../tests/vt/test_ops_fp8_cutlass.cpp#L387) + `VT_FP8_PLAN_CACHE=1` on-arm ctest; **CPU key/flag** [test_fp8_plan_cache.cpp](../tests/vt/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](../include/vt/ops.h), [cpu_ops.cpp](../src/vt/cpu/cpu_ops.cpp), [cuda_glue.cu](../src/vt/cuda/cuda_glue.cu)); model glue `ResidentFp8Qkv`/`MergedFp8QkvD`/`MergedFp8QkvEligible` + `ProjectFullAttnQkv` branch ([qwen3_5.cpp](../src/vllm/model_executor/models/qwen3_5.cpp)), resident fields ([qwen3_5_weights.h](../include/vllm/model_executor/models/qwen3_5_weights.h)); byte-exact CPU tests [test_ops_glue.cpp](../tests/vt/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](specs/fp8-merged-qkv-projection.md) | [inventory](specs/kernel-family-inventory.md); [fp8-merged-qkv](specs/fp8-merged-qkv-projection.md) | `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](../src/vt/cuda/cuda_matmul_nvfp4.cu#L1207), public zero-lifecycle [contract](../include/vt/ops.h#L518), and candidate/fallback/graph/alignment [tests](../tests/vt/test_ops_nvfp4_fp4.cpp#L959). **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](../src/vt/cuda/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](../include/vt/cuda/cuda_profiler_control.h#L13), [driver](../scripts/dgx-online-serving.sh#L14), batch-keyed [validator](../tools/bench/online_gate.py#L99), fail-closed [c2 finalizer](../tools/bench/finalize_low_batch_trace.py#L1), and [finalizer tests](../tests/tools/test_low_batch_trace_summary.py#L1) 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](../src/vt/cuda/cuda_matmul_nvfp4.cu)) + `VT_SILU_FP4_FAST` ([SiluAndMulFp4QuantFastKernel](../src/vt/cuda/cuda_matmul_nvfp4.cu)): 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](../tests/vt/test_ops_nvfp4_fp4.cpp) + full suite 24/24 (26,976), [flag header](../src/vt/cuda/fp4_quant_fast.h) + [CPU test](../tests/vt/test_fp4_quant_fast.cpp) 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](../src/vt/cuda/fp4_quant_fast.h) inverted to default-ON `=0`-rollback; [CPU flag test](../tests/vt/test_fp4_quant_fast.cpp) RED→GREEN 20/20; [CUDA byte-exact test](../tests/vt/test_ops_nvfp4_fp4.cpp) 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`](../src/vt/cuda/cuda_matmul_nvfp4.cu), [cpu composite](../src/vt/cpu/cpu_ops.cpp), [ops decl](../include/vt/ops.h)) 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](../src/vllm/model_executor/models/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](../tests/vt/test_ops_nvfp4_fp4.cpp) **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](specs/glue-fusion-2026-07-19.md) — anchor `tests/vt/test_ops_nvfp4_fp4.cpp:113` | [small-M spike](specs/nvfp4-small-m-dispatch.md); [W3-E spike](specs/nvfp4-direct-swizzled-scales.md); [W3-C spike](specs/nvfp4-persistent-plan-cache.md); [W3-F device-alpha spike](specs/nvfp4-device-alpha.md); [W3-G spike](specs/fa2-gqa-split-kv-decode.md); [W3-H normal-producer spike](specs/nvfp4-bf16-producer-vectorization.md); [W3-I fused-producer spike](specs/nvfp4-fused-silu-producer.md) | `ANCHOR-BACKFILL` | CLAIM-SERVE-GATE-1 |
-| `KERNEL-GEMM-MARLIN-W4A16` | FP4 W4A16 Marlin dense/grouped GEMM and repack | Marlin generation/targets `CMakeLists.txt:548-679,1168-1274`; capability floor `marlin_utils_fp4.py:29-35` | [cuda_marlin_repack.cu:129](../src/vt/cuda/cuda_marlin_repack.cu#L129), [cuda_moe_marlin.cu:156](../src/vt/cuda/cuda_moe_marlin.cu#L156) | [dense tests](../tests/vt/test_ops_nvfp4_matmul.cpp#L201), [MoE tests](../tests/vt/test_ops_moe_grouped.cpp#L453); 35B gate | [inventory](specs/kernel-family-inventory.md) | `ANCHOR-BACKFILL` | - |
+| `KERNEL-GEMM-MARLIN-W4A16` | FP4 W4A16 Marlin dense/grouped GEMM and repack | Marlin generation/targets `CMakeLists.txt:548-679,1168-1274`; capability floor `marlin_utils_fp4.py:29-35`; **DENSE marlin_gemm** `csrc/libtorch_stable/quantization/marlin/marlin.cu:326-541,545` | [cuda_marlin_repack.cu:129](../src/vt/cuda/cuda_marlin_repack.cu#L129), [cuda_moe_marlin.cu:156](../src/vt/cuda/cuda_moe_marlin.cu#L156); **`KERNEL-MARLIN-DENSE-PORT` (gated OFF `VT_MARLIN_DENSE`): vLLM's OWN dense W4A16 GEMM** [cuda_marlin_dense.cu](../src/vt/cuda/cuda_marlin_dense.cu), lifted dispatcher [marlin_mm_dense.cu](../src/vt/cuda/marlin/libtorch_stable/quantization/marlin/marlin_mm_dense.cu), op [ops.h `kMarlinDenseGemm`](../include/vt/ops.h), routing [dense_nvfp4_gemm.h](../include/vllm/model_executor/models/dense_nvfp4_gemm.h) — byte-preserving E=1 replacement for the single-expert MoE route (no par-regroup ULP) | [dense tests](../tests/vt/test_ops_nvfp4_matmul.cpp#L201), [MoE + DENSE tests](../tests/vt/test_ops_moe_grouped.cpp#L453); 35B gate; **dense-port unit RED-first battery (dense-vs-ref-vs-grouped, M=1..8, NVFP4+MXFP4, wrong-stride RED)**; 3 new dense `.cu` compile-clean dgx sm_121a; GPU exec PENDING | [inventory](specs/kernel-family-inventory.md) | `ANCHOR-BACKFILL` | `KERNEL-MARLIN-DENSE-PORT` |
 | `KERNEL-GEMM-INT-LOWBIT` | AWQ/GPTQ/integer Marlin, Machete, and AllSpark | `CMakeLists.txt:468-731`; Marlin types/capability `marlin_utils.py:43-149` | - | - | [inventory](specs/kernel-family-inventory.md) | `INVENTORIED` | - |
 | `KERNEL-GEMM-W4A8` | CUTLASS W4A8 dense/MoE | `CMakeLists.txt:1004-1035`; upstream test `tests/kernels/quantization/test_cutlass_w4a8_moe.py` | - | - | [inventory](specs/kernel-family-inventory.md) | `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](specs/kernel-family-inventory.md) | `INVENTORIED` | - |