docs(deepseek-v4): ds4 Q8_0 GEMV — ds4 SIDE ncu captured + DIFFed, mechanism = register-resident MLP

Closes the residual of ds4-q8-ncu-2026-07-30 (ds4 side never captured; 8 prior
agents stalled). sudo ncu --replay-mode application --graph-profiling node on ds4's
matmul_q8_0_preq_warp8_kernel over the byte-identical 6.15 GiB Q8_0 tower, GB10 sm_121a.

DIFF (ds4 vs ours): long-scoreboard 17.2 vs 54.4, regs 56 vs 39, occupancy 60.3% vs
71.9%, L2-hit 2.6% vs 10.7%, bytes/sector 1.70 vs 2.0. long-scoreboard inversely tracks
register count across 3 kernels; occupancy is non-monotone -> mechanism = register-resident
memory-level parallelism (deeper in-flight load pipeline), NOT L2 residency (REFUTED: ds4
L2 LOWER) and NOT coalescing (REFUTED: ds4 bytes/sector worse).

VERDICT (A): a portable, bit-exact lever exists (intra-row multi-block prefetch at ~56
regs) DISTINCT from the failed Brick-13 multi-ROW ILP; but bounded ~+1.5-2 t/s and does
NOT overtake ds4 raw 16.5 alone (MoE/glue residual is a separate front). Measurement-only,
no production code. Method note: --replay-mode kernel OOM-rebooted the box once (80GB model
in 119GB unified pool); fixed with application replay + avail<12GB watchdog.

FOLLOWING_AGENTS_PROTOCOL
Assisted-by: Claude Code:claude-opus-4-8 [ClaudeCode]

Author: mudler

.agents/specs/ds4-q8-raw-mechanism-2026-07-30.md

@@ -0,0 +1,160 @@
+# ds4 Q8_0 decode-GEMV — the ds4 SIDE ncu, captured and DIFFed (2026-07-31)
+
+Measurement-only (no production code). Closes the one residual of `ds4-q8-ncu-2026-07-30.md`:
+**ds4's ncu side was never captured** ("the agent stalled in prefill"). Captured now under
+`sudo ncu` on the GB10 (sm_121a), and DIFFed against our recorded kernel counters. Answers the
+raw-vs-raw question: *why does ds4 move the byte-identical 6.15 GiB Q8_0 tower at ~90% of the
+240 GB/s roofline while our `QuantDotGemmQ8_0Kernel` moves it at ~67%?*
+
+Base main `0078e85a`. Branch `ds4-q8-raw-mechanism` (records-only, NOT pushed).
+Oracle: `~/w8run/ds4/ds4` + `ds4_cuda.cu`, model `DeepSeek-V4-Flash-…-imatrix.gguf` (80.76 GiB).
+
+---
+
+## 0. Method (and the two prior-agent traps, both fixed)
+
+`sudo ncu` HW counters ARE unblocked on the box (`ncu 2025.3.1.0`, passwordless sudo).
+ds4 decodes under a CUDA graph, so `--graph-profiling node`. Target the decode kernel by
+name; SKIP prefill by launch count. Prompt `"The capital of France is"`, greedy, `-n 30`.
+Worker absent (`docker ps` empty). All runs under `flock $HOME/gpu.lock`, ONE process.
+
+Two traps that stalled the 8 prior attempts, both diagnosed and fixed here:
+
+1. **`--launch-skip 20000` skipped EVERYTHING.** The decode kernels
+   (`matmul_q8_0_preq_warp8_kernel`, …) have DISTINCT symbol names from the prefill
+   variants (`…_batch_`, `…_tok2/4/8_`, `…_kslice_`), so the name filter counts DECODE
+   launches only — of which there are ~259/step × 30 ≈ 8k total. `20000` overshot all of
+   them → header-only CSV. FIX: `--launch-skip 88` (past 2 decode steps of the plain kernel,
+   44/step) / `180` for the pair kernel.
+2. **`--replay-mode kernel` OOM-REBOOTED the box.** First real attempt (kernel replay, default
+   ctx **32768**) hard-rebooted the GB10: ncu's kernel-replay save/restore buffer over an
+   80 GB-resident model in the 119 GB UNIFIED pool tips it OOM (`gb10-unified-memory-oom-reboots-box`).
+   FIX: **`--replay-mode application`** (re-runs the app per pass, NO large save/restore buffer)
+   + `-c 512` + `--gpu-vram 85` + a live watchdog that kills ncu if `avail_mb < 12000`. Under
+   application replay the min avail stayed **26.6 GB** across both captures — box safe.
+
+Captured CSVs: `~/w8run/ncu_out/ds4_plain2.csv`, `ds4_pair2.csv` (8 steady launches each).
+
+---
+
+## 1. THE DIFF (headline deliverable) — ds4 vs ours, byte-identical Q8_0 weights
+
+Our-side numbers are the recorded Brick-13 sudo-ncu (`ds4-q8-ncu-2026-07-30.md`); ds4-side
+are this session's medians over the steady decode launches.
+
+| counter | ds4 `preq` (1:1 kernel) | ds4 `pair` (fused ×2) | **OURS** `QuantDotGemmQ8_0Kernel` |
+|---|---|---|---|
+| **long-scoreboard / issue-active** | **17.2** | 36.3 | **54.4** |
+| registers / thread | **56** | 48 | **39** |
+| achieved occupancy (`sm__warps_active` %) | **60.3 %** | 52.9 % | **71.9 %** |
+| L1 sector hit-rate | 97.4 % | 97.3 % | 96.7 % |
+| **L2 sector hit-rate** | **2.6 %** | 6.7 % | **10.7 %** |
+| bytes/sector (global ld) | 1.70 | 1.67 | 2.0 |
+| grid / block | 4096 / 256 | — / 256 | 256–512 / 256 |
+| achieved DRAM BW (prior nsys byte-acct) | **~217 GB/s = 90 %** | — | **~161 GB/s = 67 %** |
+
+`dram__throughput.avg.pct_of_peak_sustained_elapsed` returns **n/a** on GB10 (unified LPDDR5X
+exposes different metric names) for BOTH sides — the 90 %/67 % are the prior nsys byte-accounting,
+which the counters below now EXPLAIN.
+
+### The monotone signature (three data points, one law)
+`regs`  56 → 48 → 39   (ds4_preq → ds4_pair → ours)
+`longSB` 17 → 36 → 54   **inversely tracks register count**
+`occ`   60 → 53 → 72   **NON-monotone → occupancy is NOT the driver**
+
+Long-scoreboard (the global-memory-latency stall) falls as registers RISE and is INDEPENDENT of
+occupancy. That is the textbook signature of **register-resident memory-level parallelism (MLP)**:
+the kernel with more registers keeps more independent 34-byte weight-block loads in flight, so the
+compute of block *i* hides the DRAM latency of block *i+1…i+k*, keeping LPDDR5X saturated.
+
+---
+
+## 2. What the DIFF REFUTES (evidence-first — two task hypotheses are dead)
+
+- **L2-residency hypothesis — REFUTED.** ds4's L2 hit-rate (**2.6 %**) is LOWER than ours (10.7 %).
+  ds4 is NOT keeping weight blocks / scales resident in L2; it reads MORE from DRAM, not less. The
+  90 % is not an L2 tiling trick. (Consistent with LPDDR5X: weights are read once per step, no reuse.)
+- **Coalescing hypothesis — REFUTED.** ds4's bytes/sector (**1.70**) is slightly WORSE than ours
+  (2.0). Both over-fetch the 34-byte misaligned blocks heavily; both have it L1-ABSORBED (97 % L1
+  hit). ds4 does NOT use a better thread→byte mapping. Our Brick-4 aligned-repack correctly measured flat.
+- **Occupancy — NOT the driver** (ds4 wins at LOWER occupancy, 60 % vs 72 %). Confirms Brick-11's occupancy-negative.
+
+## 3. What the DIFF CONFIRMS — the mechanism, precisely named
+
+**ds4's edge is memory-LATENCY HIDING via register-resident MLP (a deeper per-warp in-flight
+load pipeline), reached at LOWER occupancy by spending registers (56 vs 39), NOT via L2 residency,
+coalescing, or occupancy.** long-scoreboard 17.2 vs 54.4 = 3.2× better latency hiding on the
+byte-identical stream.
+
+### Source grounding (both kernels read, `ds4_cuda.cu:4343` vs `cuda_quant_dot.cu:834`)
+The two plain kernels are **algorithmically identical**: warp-per-SINGLE-row, 8 warps/block (256
+threads), block-strided `for (b = lane; b < blocks; b += 32)`, `#pragma unroll` inner 8×`__dp4a`,
+34-byte blocks {half d; int8 qs[32]}, `warp_sum` reduce, lane-0 write. **NEITHER has
+`__launch_bounds__`** (grep: 0 in ds4_cuda.cu; 0 in our file). Yet nvcc allocates ds4 **56** regs
+and ours **39** — ds4 gets the deeper load pipeline, ours is register-starved into a shallow one.
+Two concrete code deltas that plausibly drive nvcc's choice (candidates to test, not yet proven):
+- **Activation representation.** ds4 pre-quantizes the activation into SEPARATE 4-byte-aligned
+  `int8 xq[]` + `float xscale[]` arrays → a single **aligned** `int32` load per 4 elems
+  (`load_i8x4_i32_aligned`). Ours reads the activation from the INTERLEAVED 34-byte `BlockQ8_0`
+  struct via `GetIntB2` = two 16-bit loads (2-byte-aligned). Interleaved addressing raises our
+  register/address pressure and constrains load hoisting.
+- **Weight load.** ds4 assembles the weight int32 byte-wise (`load_i8x4_i32_unaligned`); ours uses
+  `GetIntB2` (2×`uint16`). Different load-instruction selection changes how many block-loads nvcc
+  keeps outstanding.
+
+### Why the Brick-13 "ILP" attempt failed while ds4 succeeds (the key reconciliation)
+Brick 13 recorded the ILP axis as MEASURED-NEGATIVE and CLOSED the Q8_0 front. But Brick-13's ILP
+was **N-output-ROWS-per-warp** (`QuantDotGemmQ8_0MultiRowKernel`, `cuda_quant_dot.cu:872`) — a
+kernel that CHANGED the row→warp mapping. Measured (ncu, `~/w8run/ncu_out/ilp_ilp2.csv`): it drove
+long-scoreboard to **85–127** (WORSE than the 54 baseline) and dropped occupancy to 42–82 %. **It
+moved the wrong axis.** ds4 proves the target (longSB 17) is reachable on the SINGLE-row structure —
+via intra-row load-pipeline DEPTH, not multi-row. So the front was closed on the wrong mechanism.
+
+---
+
+## 4. VERDICT
+
+**(A) A portable mechanism EXISTS and is newly-evidence-backed — with two dead hypotheses removed
+and a bounded, honestly-modest payoff.**
+
+- **Portable lever (named precisely):** raise our plain Q8_0 GEMV's *intra-row* memory-level
+  parallelism to ds4's operating point — a DEEPER software-pipelined weight-block load pipeline
+  (manually prefetch the next K blocks' scale+qs into registers before consuming the current
+  `__dp4a`), letting/forcing nvcc to ~56 regs at ~60 % occupancy (e.g. `#pragma unroll K` on the
+  block loop, and/or `__launch_bounds__(256, 3)` / `-maxrregcount` to lift the register cap),
+  optionally adopting ds4's pre-quantized 4-byte-aligned activation for a clean aligned int32 load.
+  **Bit-exact** (same integer `__dp4a` order, same f16-scale fold). This is DISTINCT from the failed
+  Brick-13 multi-row ILP and is the 7th, best-targeted attempt — the only axis the 6 prior bricks
+  (4 align / 8 fusion / 11 sub-warp / 12 launch / 13 multi-row-ILP) never touched.
+- **Dead ends (do NOT build):** L2-residency tiling (ds4's L2 hit is LOWER) and coalescing/thread-map
+  rework (ds4's bytes/sector is WORSE) — both REFUTED by the counters.
+
+**Brutally-honest bounds (this is why it is not oversold):**
+1. **Compiler-sensitive.** nvcc chose 39 vs 56 regs on near-identical code; FORCING the deeper
+   pipeline may not linearly reproduce ds4's longSB 17 (it may spill, or plateau at ~30). It is a
+   real candidate, not a guaranteed win, on a front where 6 axes already went flat/negative.
+2. **Bounded payoff — does NOT overtake ds4 raw.** Full capture of the plain-kernel sub-gap
+   (67 %→90 % BW, −10.5 ms/step on the Q8_0 bucket) is worth **~+1.5–2 tok/s** (≈12.1→~13.9 GPU,
+   ~11.7→~13.3 wall) per the prior byte-accounting. That closes the Q8_0 SUB-gap but does **NOT**
+   reach ds4's raw **16.5 tok/s** — the MoE-expert + glue residual is a separate ~15 ms/step front.
+   So **a pure raw-vs-raw OVERALL beat via this one kernel is not on the table**; raw decode stays
+   behind ds4 even with a full Q8_0 capture. The `pair` kernel diff (longSB 36 at 48 regs) also
+   shows ds4's launch-consolidation (2 weight streams/launch) is a SECOND, independent contributor
+   already characterized in `ds4-q8-kernel-profile-2026-07-30.md` §1 — not re-opened here.
+
+**Net:** the raw-vs-raw honest truth is that ds4's Q8_0 GEMV wins by a *real, portable, latency-hiding
+mechanism* (register-resident MLP), worth one more targeted bit-exact attempt for ~+1.5–2 tok/s; but
+even fully captured it narrows rather than erases the raw decode gap. The user's rejection of the
+spec/MTP "trick" stands — this is the legitimate raw lever, correctly identified, with its ceiling
+stated plainly. Decision (build the intra-row-MLP kernel as Brick 14 / accept ~13 t/s raw ceiling)
+surfaced to the user.
+
+## 5. Reproduce
+```
+# on dgx.casa, worker absent, flock held; application replay = memory-safe:
+~/w8run/ncu_out/run_ds4_ncu2.sh "matmul_q8_0_preq_warp8_kernel" 88 8 plain2   # 1:1 kernel
+~/w8run/ncu_out/run_ds4_ncu2.sh "matmul_q8_0_pair_preq_warp8_kernel" 180 8 pair2  # fused ×2
+# metrics: lts/l1tex__t_sector_hit_rate.pct, smsp__…long_scoreboard…ratio,
+#          sm__warps_active.avg.pct_of_peak_sustained_active,
+#          smsp__sass_average_data_bytes_per_sector_mem_global_op_ld.ratio, launch__registers_per_thread
+```

docs/BENCHMARKS.md

@@ -10608,3 +10608,5 @@ speed lever. Gated microbench `VT_FP4_M1_BENCH` (test-only).
 **Q8_0 GEMV sudo-ncu root cause (2026-07-30, measurement-only — NO production code).** `.agents/specs/ds4-q8-ncu-2026-07-30.md`. Our Q8_0 kernel is memory-LATENCY-bound (long-scoreboard 54.4, occ 71.9%, L1 96.7%) — the 4 refuted structural fixes (occupancy/alignment/fusion/launch) never addressed latency; the untried lever is ILP (multiple output-rows/thread). Banked 10.07→~12.7-13 t/s. No benchmark run this record.
 
 **Laguna-S-2.1 (LagunaForCausalLM) W0 scope (2026-07-30, scope-only — NO measurements).** `.agents/specs/laguna-s21-scope-2026-07-30.md`. Novel Poolside `laguna` arch (118B/8B-active MoE) is vLLM-native (`laguna.py`, oracle-gateable) — ~85-90% REUSE, ZERO new kernel: Q4_K decode already landed (`DotSuperblock<kQ4_K>`), MoE ~90% ds2+GLM4-lite (sigmoid noaux, not softplus — card corrected), sliding-window ~95% Gemma-2/3, dual per-layer YaRN RoPE ~90% OLMo-3, GGUF loader ~90% ds4. New = small host ops (per-head softplus attn output gate, ungrouped router, name-map, per-layer variable Q-head runner). Dual-oracle (vLLM NVFP4 + llama.cpp Q4_K same-quant). No production code; bring-up queued after the DeepSeek Q8_0 lever.
+
+**ds4 Q8_0 GEMV — ds4 SIDE ncu captured + DIFFed (2026-07-31, measurement-only — NO production code).** `.agents/specs/ds4-q8-raw-mechanism-2026-07-30.md`. Captured the ds4 side (prior 8 agents stalled): `sudo ncu --replay-mode application --graph-profiling node` on `matmul_q8_0_preq_warp8_kernel`, byte-identical 6.15 GiB Q8_0 tower, GB10 sm_121a. **DIFF ds4/ours — longSB 17.2/54.4 · regs 56/39 · occ 60.3%/71.9% · L2 2.6%/10.7% · B/sect 1.70/2.0.** ds4's ~90% vs our ~67% roofline = **latency-hiding via register-resident MLP** (deeper in-flight load pipeline; longSB inversely tracks regs, occupancy is not the driver). L2-residency + coalescing hypotheses REFUTED (ds4 worse on both). Portable lever = intra-row multi-block prefetch at ~56 regs (bit-exact, DISTINCT from the negative Brick-13 multi-ROW ILP); honest ceiling ~+1.5-2 t/s, does not overtake ds4 raw 16.5 alone. No benchmark run this record.

docs/STATUS.md

@@ -1967,3 +1967,5 @@ The gate models' W4A4 decode is at the achievable roofline. See `.agents/specs/n
 **Tier-A1 fold — 5 bf16 SwiGLU-MLP archs onto the SHARED gate-up seam `layers::UnquantizedMlpGateUpMethod` (2026-07-30, branch `worktree-wf_9054a036-47d-1`, commit `<this>`, NOT pushed).** First EXECUTED tier of the cross-arch merged-GEMM fold plan (`arch-fusion-fold-plan-2026-07-30.md` §A1) AND the empirical REUSE PROOF the keep-quant generalization front (`bb864b0d`) could not run in keep-quant (DeepSeek is the only keep-quant MoE model). Each arch's hand-rolled `{ResidentWeight; MatmulBT[2I,H]; SiluAndMul}` gate-up sequence is REPLACED by the already-shared bf16 gate-up MLP method (exemplar `qwen3.cpp:91`), so multiple real gated archs now route their MLP through ONE tuned method and inherit the nvfp4 `GateUpFusedMarlinD` fused arm FOR FREE the moment a quantized checkpoint ships. **BYTE-FOR-BYTE identical op sequence — no loader concat needed** (every arch already packs a merged `[2I,H] gate_up_proj` `OwnedTensor`). Folded: **OLMo-2** (`olmo2.cpp`), **StableLM** (`stablelm.cpp`), **qwen3_dflash** (`qwen3_dflash.cpp`, was hand-merged but off-seam), **deepseek_v2** (`deepseek_v2.cpp` `DenseMlp` — one fold covers BOTH the dense-layer MLP and the shared-expert epilogue, its two call sites) via the direct `UnquantizedMlpGateUpMethod` arm; **Granite** (`granite.cpp`) via the full `layers::MakeMlpGateUpMethod` factory (its MLP weights ARE `Qwen3DenseMlpWeights`, so it is the exact exemplar and is nvfp4-ready today). The shared method's TU (`linear.h`/`nvfp4.h`) is UNTOUCHED — only new `#include`s in the 5 model TUs — so 27B/35B/qwen3_dense stay strict without re-gate. **EMPIRICAL GATES (DGX GB10 sm_121a, RelWithDebInfo + CUTLASS 4.5.0 + Triton-AOT — the PRODUCTION stack; configure log proves `CUTLASS … enabling sm120a NVFP4 cutlass GEMM` + `Triton AOT … vendored … sm_121a`).** (1) CPU composite-golden RED-first unit case in `test_linear_method` — the fused seam `UnquantizedMlpGateUpMethod::Apply` is BYTE-IDENTICAL to the standalone `{MatmulBT[2I,H]; SiluAndMul}` (RED-first PROVEN on the dev box: a perturbed reference weight fails all 15 byte checks; GREEN after revert) → on the DGX production CUDA stack `test_linear_method` **4/4 cases · 31 assertions** GREEN. (2) **OLMo-2 SACRED paged-engine gate RAN + PASSED: `test_olmo2_paged_engine` 1/1 · 92 assertions, 16/16 prompts (13/16 STRICT token-exact vs vLLM per-prompt greedy + 3/16 near-tie-band, max gap 0.094 nats, 0 forward-divergent) — IDENTICAL to the pre-fold bring-up result (`5c00fc4`), and the test's HARD anchor-drift REQUIRE against the committed pre-fold `our_ids.npy` PASSED, i.e. the folded engine produced the EXACT SAME token sequence → the fold is a proven byte-exact NO-OP on the real GPU forward through the shared seam.** Golden md5 UNCHANGED before/after (no golden overwrite). **HONEST SCOPE:** only OLMo-2 has a committed golden in-tree — Granite and StableLM goldens were NEVER committed (the `granite_greedy_2b`/`stablelm_greedy_1_6b` dirs the test comments name do not exist in the repo or on the box), so their paged-engine tests LINK + LOAD-verify + skip-on-absent-golden (rc=0, build-verified on the production stack) but do NOT run a vs-vLLM token gate THIS pass; capturing their oracle goldens is available follow-up. Granite/StableLM route through the IDENTICAL shared method OLMo-2 just proved token-exact on GPU, and are covered by the CPU composite golden. dflash + deepseek_v2: build-verified on the production stack + covered by the CPU composite golden. **Net: 1 arch empirically token-exact-gated end-to-end on GPU (OLMo-2, the shared-seam anchor) + the reuse proof (CPU composite RED-first) + 5/5 build-verified on the production stack.** No new env flag (the fold is unconditional and bit-exact; the shared method carries the existing nvfp4 arm selection). Next tiers: A2 (MLA A-proj 3→1 merge) / A4 (bf16 grouped-MoE `kMoeGateUpSwiGLU`).
 
 **Q8_0 GEMV sudo-ncu root cause (2026-07-30, measurement-only, `.agents/specs/ds4-q8-ncu-2026-07-30.md`).** sudo ncu (now unblocked) on our `QuantDotGemmQ8_0Kernel`: **LATENCY-bound** (long-scoreboard stall 54.4) at 71.9% occupancy, L1 hit 96.7% (34-byte over-fetch L1-absorbed → not the DRAM bottleneck), L2 10.7%. Explains why Bricks 4/11/12 (alignment/occupancy/launch-count) were all flat — the axis is memory-LATENCY, not those. The one untried lever = ILP (multiple output-rows/thread → independent load streams). ds4 ncu side not captured (graph-decode needs --graph-profiling). No production code changed.
+
+**ds4 Q8_0 GEMV — ds4 SIDE ncu captured + DIFFed (2026-07-31, measurement-only, `.agents/specs/ds4-q8-raw-mechanism-2026-07-30.md`).** Closes the residual of `ds4-q8-ncu-2026-07-30` (ds4 side never captured). sudo ncu `--replay-mode application --graph-profiling node` on ds4's `matmul_q8_0_preq_warp8_kernel` (the 1:1 kernel) on the byte-identical 6.15 GiB Q8_0 tower. **DIFF (ds4 vs ours): long-scoreboard 17.2 vs 54.4, regs 56 vs 39, occupancy 60.3% vs 71.9%, L2-hit 2.6% vs 10.7%, bytes/sector 1.70 vs 2.0.** Mechanism = **register-resident memory-level parallelism** (longSB inversely tracks regs across 3 kernels; occupancy non-monotone). Task's L2-residency + coalescing hypotheses both **REFUTED** (ds4 is LOWER/worse on both). VERDICT (A): a portable, bit-exact lever exists — deepen our single-row load pipeline (intra-row multi-block prefetch at ~56 regs), DISTINCT from the failed Brick-13 multi-ROW ILP; but bounded (~+1.5-2 t/s), does NOT overtake ds4's raw 16.5 alone. No production code. Box OOM-rebooted once on `--replay-mode kernel` (fixed → application replay). Decision surfaced to the user.