chore(scribe): consolidate Marlin int4 GEMM decisions + logs

.squad/agents/chew/history.md

@@ -87,3 +87,21 @@ fusion pursued: the fusion arc concluded that **native int4 decode of Muse-Glimm
 weight-bandwidth/compute-floor bound at ~47.25 tok/s (H200)** (the architectural ceiling), so
 node/launch fusion — cheap or expensive — cannot help. Rule reinforced: bf16 accumulate in fp32,
 oracle-gate against f64.
+
+## 2026-08-14 — PR #960 Marlin int4 M>1 GEMM numerics review 🟡 APPROVE-WITH-NOTES
+Reviewed Deckard's Marlin fp16×int4 tensor-core GEMM (numerics only; Gaff owns quality). Verified by close
+read of the mma B-fragment lane mapping + repack (host `repack_int4_weights` == device repack, byte-identical,
+all group sizes), the **scale-after-accumulate** factoring (exact: scale/zp constant across K within a group),
+asymmetric nibble zp indexing, padding-column guard, split-K disjoint fp32 partials (deterministic fixed-order
+reduce ⇒ capture-stable), and the lm_head cached dense-GEMM plan (cuBLASLt heuristic deterministic for fixed
+shape ⇒ cached algo == dynamic algo). Ran the in-crate GPU parity suite on H200: **11/11 pass** incl.
+`marlin_parity_vs_f64_oracle`, `marlin_splitk_parity_vs_f64_oracle`, `marlin_splitk_is_deterministic`,
+`repack_roundtrip_all_group_sizes`. Confirmed Pris's #961 f64 tolerance is engineering-justified (~8× fp16-ULP
+headroom), not a rubber stamp. **No correctness bug — ship.** Three non-blocking notes: (N1) "byte-identical
+greedy tokens" is a soft argmax-stability guarantee (empirical over 2 models × 24 tokens), NOT a numeric
+invariant — a near-tie on an untested prompt could flip; keep the flag opt-in. (N2) hard Marlin launch errors
+are silently swallowed into the tiled fallback (`Err(_) => fall through`) — numerically safe but log/count it
+so a real kernel fault can't hide behind the slow path. (N3) Marlin assumes nibble-packed int zero-points
+(matches every existing int4 tiled kernel — not a regression; guard dispatch if float-zp int4 is ever added).
+Rule reinforced: a relayout that reorders partial sums cannot be diffed bit-for-bit; the only defensible
+reference is a high-precision f64 oracle with a justified tolerance.

.squad/agents/deckard/history.md

@@ -54,3 +54,23 @@ deep kernel-capture-support program, not "reuse existing machinery"); **Lever A
 unconditional ~1.3–1.6×) promoted to the primary decode lever.** Supersedes #938's Lever-B-first call.
 
 - 2026-08-14 (#957, MERGED 2f0b62b3): Spec-capture feasibility (design-only) → **CONDITIONAL-GO gated behind Marlin** — refines (not overturns) the Lever B NO-GO. Break-even B\*=8.5 today → ~1.2–2.0 post-Marlin; ~80 ms M=8 floor is the 240 MatMulNBits generic GEMMs (Marlin fixes it) + two cheap GQA/norm M>1 capture fixes. Do NOT fund the speculative build ahead of Marlin.
+
+## 2026-08-14 — Marlin fp16×int4 tensor-core GEMM (Lever A) LANDED & MERGED (#960, 7774ec5b)
+Built + shipped the from-scratch SM80 `mma.sync.m16n8k16` fused fp16×int4 tensor-core GEMM
+(`crates/onnx-runtime-ep-cuda/src/kernels/marlin_gemm.rs`) — adapt-not-vendor (ONNX N-major nibble
+packing, asymmetric nibble zero-points, groups 16/32/64/128; per-group scale applied AFTER the
+tensor-core accumulate; raw inline `mma.sync` PTX; no upstream source ⇒ no LICENSE vendoring). Wired into
+`MatMulNBits` M>1 (plain + rmsnorm-prologue + gate_up SwiGLU fused) + split-K + the GQA/SkipLN M>1
+capture-safety valves + lm_head cached dense-GEMM capture plan. Opt-in `ONNX_GENAI_MARLIN_M_GT_1` (default
+OFF, SM80 guard, byte-identical tiled fallback — Rule 11); split-K default-ON (`ONNX_GENAI_MARLIN_SPLITK=0`
+opts out). **Result (glm-4-9b canonical gate):** M=8 verify graph **41 fragmented segments / B\*=8.76×
+(hard NO-GO) → SINGLE whole-graph capture, ZERO unsupported nodes, B\*=2.16×** (arc 8.76→4.99→2.71→2.63→
+2.16×), byte-identical greedy tokens, prefill ~2×. B\*≈2.16 is the intrinsic small-M `mma.m16n8k16` floor,
+not a tuning gap. **Attribution correction (my update-10 fixes update-8):** fused gate_up split-K needs
+`block_size==32`; glm is block-128 so its fused gate_up node never forms — glm's 2.63→2.16× was ENTIRELY
+the `4abe4e57` general small-M split-K retune, NOT gate_up fusion. qwen (block-32) forms/fires the fused
+node; its capture is fully solved (segments=1, zero unsupported) but B\*≈4.7× is a denominator effect (fast
+tuned block-32 M=1), a drafting-depth follow-up not a kernel bug. Gates the #957 spec-capture
+CONDITIONAL-GO — Lever A delivered. Reviews: Chew 🟡 APPROVE-WITH-NOTES, Gaff 🟢 APPROVE. Lesson: keep the
+flag opt-in until broader prompt/model parity coverage; add a fallback logging/counter (Chew N2) before any
+default flip.

.squad/agents/gaff/history.md

@@ -70,3 +70,20 @@ Reviewed HEAD `fbf322f76b` — the evidence-accuracy fix commit. Built from clea
 ## 2026-08-12 — PR #31973 evidence-accuracy review (focused delta)
 
 Focused evidence review of Mariette's B1+B2 fixes (HEAD `fbf322f76b`). Reproduced all four accuracy figures to 4 significant figures. Confirmed `nullptr` MeanOut matches production. Dispatch assertion fires. No stale claims found. Test counts 41/2/43 confirmed on fresh build. One nit: RMSNorm ~3.3x at NormSize 256 is ~14% above measured ~2.84x; body says ~0.83x at NormSize 15, measured 1.00x (body is conservative). Coordinator widened variance disclosure to ~15%. Verdict: ready to leave draft, no blockers.
+
+## 2026-08-14 — PR #960 Marlin int4 M>1 GEMM quality/portability review 🟢 APPROVE
+Reviewed Deckard's Marlin M>1 GEMM for quality / Rule 11 portability / capture-safety / build hygiene (numerics
+= Chew; the B\*=2.16 plateau is an accepted honest plateau, not a defect). Frozen head `a11facbc` (code
+`3735d57e`). **Zero blocking defects.** Rule 11 PASS: genuinely opt-in (`ONNX_GENAI_MARLIN_M_GT_1` default OFF,
+checked first at all three M>1 dispatch seams), SM80 guard correct (`device_supports_marlin` ANDed at every call
+site — even a force-set env var on <SM80 falls through to the portable tiled GEMM), and Marlin-OFF is provably
+byte-identical to the prior tiled path; any runtime ineligibility/launch error also falls through. Env-var honesty
+PASS: both knobs read & wired (`verify_documented_env_vars.py` EXIT 0). Capture-safety valve family sound: four
+caches share one contract — cache-hit → warm (no alloc/sync ⇒ capture-safe); cold miss while `is_capturing()` →
+return Err so caller falls back, never alloc inside capture; pre-warm forward populates them. GQA M>1 Fused flash
+path + SkipLN latch relaxation both verified capture-safe (on-device totals, no host read-back, shape-keyed cache
+rejects mid-capture shape change). `cargo fmt` + the exact CUDA and engine `-D warnings` clippy gates clean; all
+`unsafe` blocks carry SAFETY justifications. One trivial `cfg(test)`-only `clippy::unusual_byte_groupings` note
+(`0xcafef00d_1234_5678u64`) — outside the CUDA CI gate (no `--all-targets`), does not break CI, trivially fixable.
+Rule reinforced: env-var honesty + a byte-identical default-OFF fallback are the portability contract for any
+tier-scoped kernel.

.squad/agents/pris/history.md

@@ -52,3 +52,17 @@ Earlier detailed history archived in `history-archive.md`. Pre-compaction summar
 **Action:** Fixed stale FEAT_FP16 comment in `mlas.h`. Pushed commit `a0a9d98`. Proposed PR body rewrite.
 
 **Head SHA:** `a0a9d98`
+
+## 2026-08-14 — Marlin int4 numerics gate: reusable f64 dequant→GEMM oracle (#961, MERGED 401af46f)
+Landed a reusable f64 oracle test any int4 `MatMulNBits` GEMM must pass before shipping — the red/green
+target for Deckard's Marlin kernel and Chew's sign-off. Oracle dequantizes `(code-zp)·scale` and accumulates
+`Σ a·w` in f64, sharing the candidate's fp16-rounded activations + rounded scale so the residual isolates
+ONLY accumulation precision + fp16 output rounding (never shared input quant). Coverage: groups {16,32,64,128},
+sym + asym-zp, M∈{1,2,4,8,16,32}, real glm-4-9b + Qwen2.5 projection shapes, fp16/fp32 scales. Justified
+tolerance envelope (single source `Envelope::for_output(max_out)`): `abs=max(max_out·4e-3,4e-3)`, `rel=5e-2`
+(denominator floored `max(1e-1,3e-2·max_out)`) — ~8 fp16-ULP headroom, tight enough that a structural
+relayout/dequant bug blows it, loose enough to pass fp32 reassociation drift. Current tiled path measured
+~8–10× inside abs / ~6× inside rel (≈1 fp16 ULP). **Harness lesson (durable): pre-zero the output buffer +
+`runtime.synchronize()` after `execute`** — an un-synced run reading stale device-pool memory once faked an
+abs-81.8 catastrophic divergence; not a kernel bug. Marlin driver must follow the same sync/zero discipline.
+Chew concurred the tolerance is engineering-justified, not a rubber stamp.

.squad/agents/sebastian/history-archive.md

@@ -64,3 +64,124 @@ Declared `n % 8 != 0` SAFE — incorrect. Chew revised and confirmed routing con
 
 ### 2026-08-12 — PR #31973: Fix architecture-specific dispatch threshold in LayerNorm tests
 Renamed `kAvx2DispatchThreshold` → `kKernelDispatchThreshold`. Added conditions with `#if MLAS_TARGET_AMD64 || MLAS_TARGET_IX86`, else 1. Fixed `CatastrophicCancellationPasses` (condition 1e4/1e5 reachable). Renamed `AdversarialPrecisionReport` → `DISABLED_`. Removed N=7 benchmark. Head `72e02cd92c`.
+
+## Archive batch 2026-08-14T21-43-12Z (Scribe marlin — 2026-08-12/13 decode arc)
+
+## Recent work (current wave, 2026-08-12)
+
+### 2026-08-12 — PR #31973 v2: architecture-specific dispatch threshold fix
+Renamed `kAvx2DispatchThreshold` → `kKernelDispatchThreshold`. Fixed `CatastrophicCancellationPasses` to exercise accuracy branch. Renamed `AdversarialPrecisionReport` → `DISABLED_`. Removed N=7 benchmark. Head `72e02cd92c`.
+
+### 2026-08-12 — PR #762 S1/S2/S3 resolution (commit a5448fa36)
+S1: `production_scratch_alloc(numel, dtype)` helper + 2 new canary tests (`scratch_buffer_wider_write_absorbed_by_padding`, `scratch_buffer_detects_oversized_write`).
+S2: `TensorMut::validate_write_dtype()` — exact match for present, byte-size gate for absent. `mark_absent()` invariant documented.
+S3: `NodeOutputSink::Absent` variant — `build_subgraph_routing` no longer allocates phantom slots.
+Nits: removed 4 no-op identity transmutes.
+280 passed / 0 failed. Clippy clean. fmt clean. Miri: 4/4 canary tests clean.
+
+### 2026-08-12 — PR #832 H200 CUDA validation build fix (MERGED `2b62c620`)
+Added the missing `bf16_scratch` field (`Mutex<Bf16Scratch>`, `Mutex::new(Bf16Scratch::new(runtime.clone()))`) to 11 `MatMulNBitsKernel` test initializers in `crates/onnx-runtime-ep-cuda/src/kernels/matmul_nbits.rs`. Verified `cargo test --no-run -p onnx-runtime-ep-cuda --features cuda` green. Merged as part of the H200 (Muse-Glimmer-30B) CUDA EP validation wave.
+
+Full pre-compaction history in `history-archive.md`.
+
+### 2026-08-12 — CUDA-graph capture escalation (background, agent sebastian-3)
+Redirected post-#840 to investigate why CUDA-graph capture does not engage for
+Muse-Glimmer native decode. Delivered a cross-domain escalation: **3 stacked
+blockers** — (1) LOAD (engine native pipeline can't load the model), (2) CLASSIFY
+(vestigial SWA mis-classification on decode path), (3) CAPTURE (infra proven, gated
+behind 1+2). No perf PR (model can't load on engine native path yet). Coordinator
+dispatched Batty (LOAD) + Deckard (CLASSIFY); I pair on CAPTURE + re-measure once
+unblocked. Shared team goal: **beat ORT 40 tok/s via CUDA-graph capture**. Prior
+#840 (629fbf90) merged: real cudaMemGetInfo device-capacity + CudaFoldConstantCast,
+native decode 10.2→11.4 tok/s (+11.8%).
+
+## 2026-08-12/13 — CUDA capture arc COMPLETE (shared: 11.4 → 23.13 tok/s)
+Owned diagnosis + escalation + the two CUDA-EP kernel blockers. **#855** (`1022b912`)
+`gqa_decode_bf16` capture-safe kernel (fp32 accumulation; Chew-gated, max_abs 1.953e-3):
+segments 54 → 2, 22.52 tok/s. **#854** (`f85a82f0`) skip-norm capture-safety (persistent
+`NormBf16Scratch`, demote on `grew` only when `is_capturing()`): segments 2 → 1, 0 seams,
+23.13 tok/s (+33% vs capture OFF). Built on #848 (Deckard) / #850 (Batty) / #852 (Leon).
+**Corrected my own diagnosis:** with the step captured, decode is now **kernel-bound**
+(Cast 40%, MatMulNBits 21%, GQA 14%), not pure dispatch-bound. Next lever = Cast
+round-trip elimination to reach ORT's ~40 tok/s.
+
+## 2026-08-12/13 — PR #860 MERGED: 23→40 tok/s, CUDA goal MET (11.4 → 40.21 tok/s)
+Closed the final lever. Generalized ep-cuda `CudaDropNormalizationCasts` to fold **bf16**
+casts around `RMSNormalization` (op-swap `RMSNormalization`→`SimplifiedLayerNormalization`
+for re-inference stability — both map to the same `RmsNormFactory→RmsNormKernel`), and
+rewrote `rmsnorm_bf16` with a **parallel f32 tree reduction** (fp32 accumulation; only
+summation order changes). Native CUDA decode **23.16 → 40.21 tok/s** (H200, 3-run median),
+1 segment / 0 seams, first-16 greedy ids match reference. Cast removal is ~free under
+capture; the parallel reduction is the real lever (serial `fadd` chain ≈40% of captured
+decode; serial floor ≈33 tok/s). Chew-gated 🟢. Escape hatch
+`ONNX_GENAI_CUDA_DISABLE_NORM_CAST_FOLD=1`. **Multi-session CUDA goal MET: matches ORT ~40 tok/s.**
+
+## 2026-08-13 — PR #867 MERGED: 40→47 tok/s, native CUDA now BEATS ORT
+Cached the Float16-staged constant int4 scales. `MatMulNBitsKernel::run_bf16` was re-casting
+the immutable int4 scale slots bf16→f16 into an ephemeral arena **every** decode step
+(~3.3 GB/token, ~25% of int4 weight traffic, 417 redundant cast launches). Added a persistent
+per-kernel `Bf16ConstCache` that stages the constant scale slots once (general path input 2;
+SwiGLU-fusion inputs 2 and 4) and reuses them; dynamic slots (activation input 0, bias
+residual) stay on the ephemeral arena. Capture-safe (alloc+cast on pre-capture warmup; replays
+hit lookups only). Native decode **40.21 → 47.25 tok/s** (H200, 3-run median), MatMulNBits eager
+share ~44%→~31%, 1 segment / 0 seams, full 128-token sequence byte-identical. **Byte-exact by
+construction — no Chew gate**; test `bf16_scale_cache_is_bit_exact_to_inline_staging`.
+**MILESTONE: native CUDA EP now clearly beats ORT — 47.25 vs ~40 tok/s, +18%.** Next lever
+(deferred): GQA is now the largest eager share (~41%); full bf16-native GEMV / accuracy_level=4
+DP4A int8 remain open, both numerics-gated.
+
+## 2026-08-13 — bf16 SwiGLU kernels (#871) + GQA null finding (#870); 47.25 is the CEILING
+Shipped **#871** (bf16 decomposed SiLU/SiLU-Mul kernels `decomposed_silu_mul_bf16`/
+`decomposed_silu_bf16` in `elementwise.rs`) — fixes a real portability **hard crash**
+(bf16 decomposed SiLU previously errored `"requires float16"`); byte-exact **0 ulp** vs f64
+oracle, **Chew 🟢**, 5/5 silu tests on H200. Its graph SwiGLU-Mul fold is FLAT (−104 cheapest
+nodes). **#870** (doc-only): decode is 2568 nodes/token × ~8.17 µs/node; cheapening any single
+kernel inner loop (GQA seq-loop, GEMV depth-loops) is flat → GQA not a viable lever. Full gate_up
+bf16 fold went non-deterministic (f16-staging); safe version needs a bf16-native fused
+`gate_up_swiglu` kernel + Chew — deferred. **CONCLUSION (with Batty #872/#873): native int4 decode
+of Muse-Glimmer-30B is weight-bandwidth/compute-floor bound at ~47.25 tok/s (H200), NOT
+dispatch-bound. 47.25 is the architectural ceiling; beat it via fewer weight bytes/token or a
+megakernel, NOT node fusion.**
+
+## 2026-08-13 — Lower-bit quant NO-GO + mechanism correction (#885, docs-only)
+Researched whether lower-bit quant could beat ~47 tok/s. Byte budget: decoder reads **15.325
+GB weights/token** = 724 GB/s at 47 tok/s = only **~15% of H200 HBM roofline**. Ran a controlled
+weight-DRAM byte-fold probe (`ONNX_GENAI_WEIGHT_FOLD=D`, throwaway/reverted, node-count
+byte-identical): full **47.29** → half **47.98** (+1.5%) → quarter **48.62** (+2.8%). Weight-DRAM-
+bound fraction ≈ **3–4%**; int2-everywhere (−45% bytes) projects ≈+1.6%. **VERDICT: lower-bit quant
+(int3/int2/mixed/2:4/NF4) is a MEASURED 🟥 NO-GO.** **KEY CORRECTION:** this REFUTES the earlier
+"weight-bandwidth/compute-floor bound" attribution (#870/#872/#873) — decode is **LATENCY-bound
+on the ~2568-node serial chain (~8.2 µs/node)**, not bandwidth-bound. Ceiling VALUE (~47) and
+"marginal fusion isn't a lever" still stand; only the WHY changes. **Megakernel / drastic per-layer
+node-collapse REOPENED as the true next lever.** Brief: `docs/research/lowbit-quant-feasibility.md`.
+
+## 2026-08-13 — Dense-decode megakernel: multi-CTA GEMV megakernel MEASURED 🟥 NO-GO (#898)
+Built the persistent multi-CTA cooperative one-layer megakernel (MLP triple-GEMV, 1056 co-resident
+CTAs = 8/SM × 132 SMs, grid.sync seams, L2-resident scratch, production int4 GEMV math) and measured
+vs identical-math per-op baseline (H200, median 200 iters × 3). **Per-op 0.656 → megakernel
+0.676–0.680 ms/layer-MLP = −3.2% recovered (~3% SLOWER), byte-exact 0-ulp; grid.sync = 2.23
+µs/barrier.** Projected whole-model gain ≈ 0% (stays ~47 tok/s). **This CLOSES the megakernel lever
+the #885 correction reopened** — mechanism: (1) CUDA-graph replay already banks the per-launch
+overhead, (2) multi-CTA must PAY a grid.sync tax that cancels savings, (3) GEMVs are genuine
+full-device weight reads already fanned across 132 SMs. **Redirect: the only remaining
+recoverable-overhead lever is graph-side glue node-collapse (Batty, `optimizer.rs`);** my role is
+the already-landed fused epilogues (#867, #854) that enable node deletion. Merged PR #898
+(main @ 0790849c). Staged arc (superseded): Phase A/B feasibility GO → P1.5 (single-CTA 926× slower,
+grid.sync capturable) → P2 multi-CTA NO-GO. Doc: `docs/research/dense-decode-megakernel-feasibility.md` §7.
+
+## 2026-08-13 — bf16 skip-RMSNorm KERNEL byte-exact SHIP; standalone FOLD −1.5% NO-SHIP (#903)
+Closed #900's blocker: built the missing byte-exact **bf16 skip-RMSNorm kernel** for Gemma3
+sandwich-norm (`skip_rmsnorm_bf16` NVRTC, `crates/onnx-runtime-ep-cuda/src/kernels/normalization.rs`).
+`sum = __float2bfloat16_rn(f32(residual)+f32(x))` (bit-for-bit a standalone bf16 `Add`) then the
+identical `rmsnorm_bf16` block-tree reduction. **BYTE-EXACT 0-ulp** (GPU unit tests bit-identical;
+real-model fold OFF vs ON bit-identical, 128/128 tokens). **KERNEL: SHIP.** But the standalone
+`CudaSkipRmsNormFusion` fold **REGRESSES −1.5%** under graph replay (fold OFF 47.77 → ON 47.06
+tok/s, 104 seams): at M=1 the single-CTA RMS reduction serializes the residual `Add` that the
+standalone spread across 132 SMs; replay already banks the launch saving. **FOLD: NO-SHIP as
+default** — retained opt-in behind `ONNX_GENAI_CUDA_ENABLE_SKIP_RMSNORM_FUSION` (default OFF).
+This is the THIRD independent confirmation of the batch-1 decode LATENCY FLOOR (with #898 megakernel
+NO-GO and #899/#900 glue-collapse +0.9% ceiling). **Next lever (NOT funded):** fold the bf16 norm
+into the neighbouring multi-CTA int4 GEMV prologue/epilogue (bf16 analogue of
+`CudaSkipRmsNormMatMulFusion`) — keeps the reduction distributed. Do NOT self-merge (Chew gates
+numerics). Doc §8.6.
+