docs: M2.7 tensor-core NVFP4 GEMM — vLLM gap 16.8x->7.5x; GDN scan now dominates

Free-box measured vs vLLM: prefill TTFT 14.4->6.1s (2.36x), 8x1024x128 total
70->157 tok/s (2.24x), batched TTFT 100->32.8s (3.05x). Paged gate 16/16. nsys:
fp4 MoE GEMM 70.7%->11.6% of prefill; bottleneck moved to GdnScanKernel (63.8%).
Next levers: GDN chunk-parallel scan (prefill), decode fp4-GEMV + fast-argmax.

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

Author: mudler

.agents/parity-ledger.md

@@ -47,3 +47,4 @@ Columns:
 | 2026-07-04 `69eab0d` (GB10, apples-to-apples vs vLLM per the new protocol directive) | **device-resident weights** (killed the redundant per-op weight RE-UPLOADS — the ~600MB embed table was re-copied to device EVERY op, + all q/k/v/o/GDN/router/norm weights per layer; now a lazy per-OwnedTensor device cache, upload-once). This is what the 67.5%-memcpy profile was largely made of | numerically identical (same bytes/upcast); CPU path unchanged | **BOTH sides measured on the SAME workload 8×128×64 conc4 on GB10 (per the always-compare-vs-vLLM directive):** OURS **14.33 output tok/s / 43.05 total** (device-resident, +38% vs the pre-change 10.36/31.13 on this config), vs **vLLM 0.24.0 ENFORCE-EAGER 131.3 output / 1181.8 total** (gpu-mem capped 0.7 to coexist with the box's LocalAI). **Gate-#1 gap (eager-vs-eager, fair near-term): ~9.2× decode output, ~27× total (prefill-dominated).** vs FULL vLLM (WITH cuda graphs, ~153/1377): larger — the graph delta is exactly our next unlock. test_qwen36_paged_engine 16/16 correct. NOTE (directive honored): the prior step re-based the bench config without a matching vLLM run — this row fixes that; the vLLM baseline needed gpu-mem capping + retries around LocalAI memory fluctuation + a transient profiling race. NEXT: CUDA graphs (decode host-API kill → the ~9× decode gap) + GEMM/prefill fusions (the ~27× total gap) |
 | 2026-07-04 `f95b131` (GB10, device-resident forward) | **device-ify the mid-forward host glue** (Phase 1 of the CUDA-graph unlock): the forward's per-op `Download()`/`Synchronize()` + the host-loop glue (attention gate-split + sigmoid output-gate, GDN g/beta + conv q/k/v split, shared-expert sigmoid gate) → 6 new device elementwise vt ops (CastBf16, AttnGateSplit, SigmoidGateBf16, GdnGBeta, GdnConvSplit, SharedExpertGate; CUDA+CPU+unit tests) + device-resident matmul helpers returning DBuf (no Download); blocks now thread hidden/res as device DBuf. The fp4+CUDA DECODE body is now PURE async-on-stream (only Embedding + the final-logits Download remain host) — the prerequisite for CUDA-graph capture | 6 glue ops == CPU reference (unit tests); paged==dense bit-exact anchor holds | **VALIDATED + MEASURED ON GB10:** test_qwen36_paged_engine 16/16 token-for-token; test_ops_glue PASS; CPU ctest 82/82. **Measured +46% output/total, +76% per-stream decode** (fair A/B: 3.97→5.79 output tok/s, both runs under heavy LocalAI GPU contention that depressed absolutes vs the free-box 14.33 — a free-box re-measure vs vLLM 131 is the open item [directive: needs both sides on a free box; LocalAI was pinning the GPU ~89%]). NEXT: **Phase 2 CUDA-graph capture** (now unblocked — the decode body is sync-free): hoist per-token inputs (positions/slot_mapping/block_table/seq_lens/qsl) to persistent device buffers, pre-warm the pool (no cudaMalloc in capture), keep Embedding outside, wire BeginCapture/Replay into GPUModelRunner (re-capture on batch-shape change) → collapse the ~88% host-API to one cudaGraphLaunch/token |
 | 2026-07-04 `2999431` (GB10 FREE-BOX, decode CUDA-graph capture) | **Phase 2: wire CUDA-graph capture/replay into the decode path** (`Qwen3_5DecodeGraph`: Embed outside the graph, capture the residual+layers+norm+lm_head region, persistent host input vectors mutated in place — probe confirmed pageable H2D memcpyAsync IS capturable on GB10 so no device-input-buffer threading needed; MoE expert ptr/scale arrays made resident; gated to num_reqs==1 fp4/CUDA, `VLLM_CPP_CUDAGRAPH=0` disables). Embedding op kept outside capture (it cudaMalloc/syncs a bounds flag — illegal in capture). | paged greedy gate 16/16 token-for-token WITH graph active; CPU 79/80 (the 1 red `test_qwen36_weights` is pre-existing: asserts bf16 experts, checkpoint is NVFP4) | **VALIDATED + FREE-BOX MEASURED (LocalAI worker stopped, 0% contention).** Graph win is MARGINAL: num_reqs=1 TPOT 67.21→65.48ms (**+2.6%**); num_reqs=8 **−7% (regresses)** → gated to num_reqs==1, zero batched regression. **HONEST NEGATIVE — CUDA graphs are NOT the gate-#1 unlock:** Phase 1's async-on-stream conversion already hid host launch overhead behind the GPU, so decode is now **GPU-compute-bound**. **FIRST CLEAN FREE-BOX vLLM COMPARISON (always-compare directive, 8×1024×128 enforce-eager):** vLLM **1124 total / 124.9 output tok/s** vs ours **70.2 total / 7.8 output** → **~16× slower**, dominated by **prefill TTFT ~100s** + slow reference kernels (naive NVFP4 dequant-GEMM, GDN scan, attention), NOT launch overhead. **PROFILE RE-ORIENTED: the remaining gate-#1 gap is GPU-KERNEL/PREFILL COMPUTE SPEED** (tensor-core NVFP4 MMA GEMM vs cutlass, flashinfer-class attention), not host overhead. Graph infra kept (correct, gated, mirrors vLLM for upstream porting). |
+| 2026-07-04 `bc0a8d7` (GB10 FREE-BOX, tensor-core NVFP4 W4A16 GEMM) | **M2.7: move the prefill-dominant fp4 W4A16 GEMMs onto Blackwell tensor cores** (`nvcuda::wmma` m16n16k16 bf16×bf16→f32, dequant-into-shared-bf16 — dequant stays CUDA-core, MMA on tensor cores). Dense `MatmulNvfp4Wmma` (64×64 tile, BK=32, 2×2 warps, col-major B-load = no transpose). Grouped `MoeGroupedGemmNvfp4Wmma` (the 70% lever): added **device-side expert grouping** (counting sort: histogram/prefix/atomic-scatter into pair-row + source-act-row, ragged per-BM-tile expert map, tail-masked) → dense bf16 WMMA GEMM per expert, buying BOTH weight-reuse AND tensor cores. Self-contained in the op (identical signature, no engine/model changes); large-P prefill only (small-P decode keeps the naive kernel byte-for-byte); `VT_NVFP4_WMMA=0` toggle for same-binary A/B. Mirrors killgate 0035 (Marlin grouped) shape, adapted to modelopt layout + added the expert-grouping (0035's input was pre-sorted). | unit `test_ops_nvfp4_matmul` 4/4 (37/37 assert: single-tile, multi-tile P/E>BM ragged, odd N/K tails, WMMA==per-expert dequant-matmul ref); paged gate **16/16 token-for-token**; CPU 82/82; -Werror clean rebuild green | **VALIDATED + FREE-BOX MEASURED vs vLLM (0% contention, same-binary A/B).** Prefill TTFT **14.40→6.10s (2.36×)**; 8×1024×128 total **70.16→157.27 tok/s (2.24×)**, output 7.79→17.47, batched TTFT **100.06→32.76s (3.05×)**. **vLLM oracle (8×1024×128 enforce-eager, same box) = 1181.78 total / 131.31 output → total gap CLOSED 16.8×→7.5×.** nsys: fp4 MoE GEMM **70.7%→11.6%** of prefill (now wmma); **bottleneck MOVED to `GdnScanKernel` 63.8%** (the recurrence — next lever, out of scope here). HONEST: kernel ~3–4 TFLOP/s useful (below bf16 peak — small N=512 MoE tiles, ragged BM waste ~32 rows/expert, CUDA-core dequant); the 2.36× is driven as much by expert-grouping (weight reuse) as tensor cores. Native-fp4-MMA stretch NOT attempted (killgate 0034/0035: W4A4 FP4-MMA regressed on GB10 — low marginal leverage now GDN dominates). Tuning headroom: larger tiles, cp.async pipeline, drop BM waste. |

.agents/state.md

@@ -976,3 +976,17 @@
   breakdown by kernel, then attack the dominant cost (almost certainly the NVFP4 GEMM → a
   tensor-core MMA kernel, the killgate prior art). GPU is free — keep it serialized to one
   workstream. Restart `local-ai-worker` when the measurement campaign pauses.
+- **2026-07-04 (M2.7 tensor-core NVFP4 GEMM — gap 16.8×→7.5×; prefill bottleneck now GDN scan)**
+  — `bc0a8d7`. Moved the prefill-dominant fp4 W4A16 GEMMs (MoE grouped + dense projections)
+  onto Blackwell tensor cores (bf16 WMMA, dequant-into-shared) + device-side expert grouping
+  for the MoE (counting sort → dense per-expert GEMM, weight-reuse + tensor cores). Correctness
+  preserved (paged gate 16/16, unit 4/4, CPU 82/82, -Werror clean). **Free-box measured vs
+  vLLM:** prefill TTFT 14.4→6.1s (2.36×), 8×1024×128 total 70→157 tok/s (2.24×), batched TTFT
+  100→32.8s (3.05×). **Gap to vLLM CLOSED 16.8×→7.5×** (vLLM 1181 total / 131 output). nsys:
+  fp4 MoE GEMM 70.7%→11.6% of prefill; **the prefill bottleneck MOVED to `GdnScanKernel`
+  (63.8%)** — the gated-delta-net recurrence, a sequential-scan problem (needs chunk-parallel
+  scan, NOT MMA). HONEST: kernel below bf16 peak (~3-4 TFLOP/s useful; win is part tensor-core
+  part expert-grouping) — tuning headroom remains. NEXT LEVERS (measured): (1) **GDN chunked/
+  parallel scan** — now 63.8% of prefill, THE next prefill lever; (2) M2.8 decode fp4-GEMV +
+  fast-argmax (the ~100× argmax, the naive M=1 fp4 GEMV, route the cublas bf16 gemvx to fp4).
+  GPU stays free+serialized (LocalAI worker down, restart disabled per user directive).