docs: record M2.3 GDN chunk-parallel scan (gap 7.5x->5.84x) + 27B bring-up findings

35B M2.3: chunk-parallel GDN prefill scan (2ce938f, mirrors FLA chunk.py) — GDN
scan 3.0x faster, batched TTFT 33.2->19.4s, total 155.6->202.5 tok/s, gap to vLLM
7.5x->5.84x, paged gate 16/16. Bottleneck moved to PagedAttention (28.6%).
27B: shares the 35B backbone wholesale; W4A4 weight encoding == modelopt NVFP4, so
the M2.7 tensor-core GEMM carries it with a one-line reciprocal (fast path, no new
kernel); 27B is VL-multimodal (text-first, ViT deferred). Also lands the GDN test
cosmetic cleanup. Note: 2ce938f bundled the GDN kernel under a docs message.

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

Author: mudler

.agents/parity-ledger.md

@@ -49,3 +49,4 @@ Columns:
 | 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. |
 | 2026-07-04 (27B W4A4 CPU scaffolding) | **27B dense NVFP4 W4A4 — CPU-first correctness scaffolding** (the 2nd MVP gate, ZERO prior bring-up; CPU-only, no GPU touched). Surveyed the real checkpoint (`unsloth/Qwen3.6-27B-NVFP4`, read-only on dgx) + pinned upstream → `.agents/qwen27b-w4a4-notes.md` (arch dims/layer-pattern, SHARED-vs-NEW vs the 35B, the compressed-tensors W4A4 format from the actual manifest, the bring-up plan). **TWO findings that change the plan:** (1) the 27B is `Qwen3_5ForConditionalGeneration` — a **VL multimodal** model (vision_config/ViT + image/video tokens), not plain text; text path first, ViT deferred. (2) quant is compressed-tensors **W4A4** (weights AND activations fp4, dynamic per-token) — BUT the on-disk weight encoding == modelopt NVFP4 (only the tensor names + a reciprocal-global-scale differ), so the existing M2.7 W4A16 tensor-core GEMM can carry the 27B weights with a 1-line reciprocal fix (recommended fast path; true W4A4 fp4-MMA is a measured follow-up — killgate 0034/0035 found it regressed on GB10). Delivered: CPU W4A4 dequant + activation-quant emulation reference `nvfp4_emulation.h/.cpp` (mirrors `nvfp4_emulation_utils.py` + `compressed_tensors_w4a4_nvfp4.py` + `kernels/linear/nvfp4/emulation.py` @ e24d1b24) + skipping greedy-parity gate `test_qwen27_paged_engine.cpp` | mirrors pinned upstream file:line (weight dequant `dequantize_to_dtype` swizzle=False + the CT `1/weight_global_scale`; activation `ref_nvfp4_quant` using on-disk `input_global_scale` directly; `cast_to_fp4` buckets; `run_nvfp4_emulations`). NO GPU/oracle run — the greedy + throughput oracle (pip-vLLM on the same ckpt / `vllm bench throughput`) is set up but GPU-gated (notes §5-§6) | **CPU-GREEN, clean rebuild (-Werror).** `test_ct_nvfp4_emulation` 6 cases / 81 assertions PASS (hand-computed vs the emulation math: F32↔F8E4M3 round-trip, CastToFp4 buckets, CT weight dequant, activation round-trip identity + rounding, full emulated W4A4 matmul). Full suite **84/84** (was 82; +2 new, no regression to the 35B). NEXT (notes §5): config+loader plumbing → dense forward assembly (reuse 35B GDN/attn, swap MoE→dense SwiGLU) → [GPU] oracle golden + W4A4 GEMM (fast path 6a reuses M2.7) → flip the gate |
+| 2026-07-04 `2ce938f` (GB10 FREE-BOX, GDN chunk-parallel prefill scan) [NOTE: kernel bundled into the `2ce938f` "docs" commit by a concurrent `git add -A` — content is correct+committed, message is mislabeled] | **M2.3: chunk-parallel GDN prefill scan** mirroring vLLM FLA `chunk.py` (chunk=64). 4 kernels: `GdnChunkCumsumKernel` (←cumsum), `GdnChunkWUKernel` (←scaled_dot_kkt+solve_tril+wy_fast fused — forward-substitution for the 2 WY cols, no explicit triangular inverse), `GdnChunkDeltaHKernel` (←chunk_delta_h, the only cross-chunk-sequential kernel, `[Dv,Dk]` state in 64KB opt-in shmem), `GdnChunkOKernel` (←chunk_o). Intra-chunk parallel over the whole grid; only cross-chunk recurrence sequential (depth ⌈T/64⌉≈16 vs T=1024). Numerically stable via the bounded `exp(G_i−G_j)`≤1 form (avoids f32 overflow). Prefill only; decode 1-step untouched; `VT_GDN_CHUNKED=0` toggle. Found+fixed a multi-head `hstate` indexing bug (`(gc*hv_n+hv)` was missing the Hv dim → heads collided). | unit `test_ops_gdn` 23/23 (310 assert: 1/2/tail chunks, multi-head Hk2/Hv4 + Hk4/Hv8 GQA-2, varlen; chunked==sequential f32 max|diff|~3e-8, bf16~6e-5); paged gate **16/16 token-for-token** with chunked as default; full suite green | **VALIDATED + FREE-BOX MEASURED vs vLLM (same-binary A/B).** GDN scan itself **3.0× faster** (nsys 63.7%→37% of prefill). Single-prompt TTFT 6086→4705ms (1.29×); **8×1024×128 batched TTFT 33.2→19.4s (1.71×)**, total **155.6→202.5 tok/s (1.30×)**, output 17.3→22.5. **vLLM 1181.78/131.31 → gap CLOSED 7.5×→5.84×.** BOTTLENECK MOVED: `PagedAttentionKernel` now the largest at **28.6%** (full-attn-layer prefill → FlashInfer-class GQA 16/2 is the next prefill lever), `GdnChunkDeltaH` 18.5% next (tensor-core headroom on the state rank-update). |

.agents/state.md

@@ -1049,3 +1049,24 @@ dgx inspection of the checkpoint).
   dense forward assembly (reuse 35B GDN/attn, swap MoE→dense SwiGLU) → [GPU]
   capture the pip-vLLM greedy golden + wire the W4A4 GEMM (fast path 6a) → flip
   `kW4A4ForwardReady`.
+- **2026-07-04 (TWO MVP tracks landed — 35B gap 7.5×→5.84×; 27B bring-up: shares 35B backbone)**
+  35B: **M2.3 GDN chunk-parallel prefill scan** (`2ce938f`, mirrors FLA chunk.py; GDN scan 3.0×
+  faster, batched TTFT 33.2→19.4s, total 155.6→202.5 tok/s, **gap 7.5×→5.84×**, paged gate 16/16,
+  found+fixed a multi-head hstate bug). Prefill bottleneck MOVED to `PagedAttentionKernel` 28.6%
+  (FlashInfer-class GQA 16/2 = next prefill lever) then GdnChunkDeltaH 18.5%.
+  27B: **CPU-first W4A4 bring-up** (`559e5cc`). TWO PLAN-CHANGING FINDINGS: (1) the 27B is a
+  **VL-multimodal** model (`Qwen3_5ForConditionalGeneration` + vision_config) — do the TEXT path
+  first, defer the ViT (T1/T2). (2) It **shares the 35B hybrid backbone wholesale** (GDN, gated
+  attn, RoPE, Gemma norm — GQA ratio 3 vs 2 is a dims-only change); the ONLY new structure is the
+  **dense SwiGLU MLP** (vs MoE) + the W4A4 quant. CRUCIALLY the **W4A4 weight encoding == modelopt
+  NVFP4 byte-identical** (CT just stores globals as divisors) → the existing **M2.7 tensor-core
+  GEMM can carry the 27B with a one-line `1/weight_global_scale` reciprocal + name remap (bf16
+  activations, W4A16-style)** — the FAST PATH to 27B correctness+throughput, NO new kernel. True
+  fp4×fp4 MMA (6b) is a risky optional (killgate 0034/0035 saw W4A4 fp4-MMA regress on GB10).
+  Landed: CPU W4A4 dequant/activation reference + `test_ct_nvfp4_emulation` 6/81 assert, skipping
+  27B greedy gate scaffold, `.agents/qwen27b-w4a4-notes.md`. Full suite **84/84**.
+  27B REMAINING (ordered): [CPU] loader plumbing + dense forward assembly (reuse 35B, MoE→SwiGLU);
+  [GPU] capture pip-vLLM greedy oracle golden → wire the M2.7 GEMM (reciprocal fast path) → close
+  greedy + throughput gates vs oracle. CONCURRENCY LESSON: two code-writing subagents on the same
+  working tree caused a commit-bundling mislabel (`2ce938f`) — use worktree isolation or explicit-
+  path staging for parallel code-writers, never `git add -A` while a subagent is mid-write.

tests/vt/test_ops_gdn.cpp

@@ -850,26 +850,6 @@ void RunGdnChunkedVsSequential(const std::vector<int32_t>& qsl, int64_t hk, int6
   run("0", out_seq, st_seq);
   run("1", out_chunk, st_chunk);
 
-  {
-    auto os = Unpack(out_seq, cb.out), oc = Unpack(out_chunk, cb.out);
-    float mad = 0.0f, mrd = 0.0f;
-    size_t argmax = 0;
-    for (size_t i = 0; i < os.size(); ++i) {
-      const float ad = std::fabs(oc[i] - os[i]);
-      if (ad > mad) {
-        mad = ad;
-        argmax = i;
-      }
-      const float rd = ad / (std::fabs(os[i]) + 1e-6f);
-      if (rd > mrd) mrd = rd;
-    }
-    MESSAGE("out max|diff|=" << mad << " at i=" << argmax << " seq=" << os[argmax]
-                            << " chunk=" << oc[argmax] << " maxrel=" << mrd);
-    float smad = 0.0f;
-    for (size_t i = 0; i < st_seq.size(); ++i)
-      smad = std::max(smad, std::fabs(st_chunk[i] - st_seq[i]));
-    MESSAGE("state max|diff|=" << smad);
-  }
   CheckClose(Unpack(out_chunk, cb.out), Unpack(out_seq, cb.out), atol, rtol);
   CheckClose(st_chunk, st_seq, atol, rtol);
 }