From 68256ac3a6c5583984ee751f193a0ba5bc627a2b Mon Sep 17 00:00:00 2001 From: Justin Chu Date: Fri, 14 Aug 2026 21:56:07 +0000 Subject: [PATCH] =?UTF-8?q?chore(scribe):=20consolidate=20Marlin=20int4=20?= =?UTF-8?q?GEMM=20landing=20(Lever=20A)=20=E2=80=94=20decisions,=20logs,?= =?UTF-8?q?=20histories?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Merged 5 inbox drops (deckard-marlin-kernel #960, pris-marlin-numerics #961, sebastian-marlin-bench #962, chew-marlin-numerics-review, gaff-marlin-quality-review) into a single consolidated Marlin program entry in decisions.md, with the corrected per-model attribution (Deckard update-10 supersedes update-8: glm block-128 gate_up fusion is a NO-OP; glm 2.63→2.16× was the general small-M split-K retune). - decisions.md 45,991 B → 42,104 B (HARD-GATE: archived by SIZE per charter — three verbose "Last consolidated" lines + the resolved #957/#948/#949 spec-capture & Lever B sub-entry bodies moved to decisions-archive/2026-08.md; live keeps compact pointers). - Deleted the 5 processed inbox drops (README kept). - Wrote orchestration logs for deckard, sebastian, pris, chew, gaff. - Wrote session log 2026-08-14T21-43-12Z-marlin-int4-gemm-landed. - Appended the Marlin outcome to deckard/sebastian/pris/chew/gaff histories. - HISTORY chronicle gate: compacted sebastian (12,308 B / ~13 dated entries) — 08-12/13 decode arc moved to history-archive.md; live keeps role + durable lessons + context + the 08-14 entries. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> --- .squad/agents/chew/history.md | 18 + .squad/agents/deckard/history.md | 20 + .squad/agents/gaff/history.md | 17 + .squad/agents/pris/history.md | 14 + .squad/agents/sebastian/history-archive.md | 121 ++++ .squad/agents/sebastian/history.md | 122 +--- .squad/decisions-archive/2026-08.md | 82 +++ .squad/decisions.md | 92 +-- .../decisions/inbox/deckard-marlin-kernel.md | 503 --------------- .../decisions/inbox/pris-marlin-numerics.md | 63 -- .../decisions/inbox/sebastian-marlin-bench.md | 595 ------------------ ...08-14T21-43-12Z-marlin-int4-gemm-landed.md | 68 ++ .../2026-08-14T21-43-12Z-chew.md | 39 ++ .../2026-08-14T21-43-12Z-deckard.md | 37 ++ .../2026-08-14T21-43-12Z-gaff.md | 37 ++ .../2026-08-14T21-43-12Z-pris.md | 35 ++ .../2026-08-14T21-43-12Z-sebastian.md | 34 + 17 files changed, 547 insertions(+), 1350 deletions(-) delete mode 100644 .squad/decisions/inbox/deckard-marlin-kernel.md delete mode 100644 .squad/decisions/inbox/pris-marlin-numerics.md delete mode 100644 .squad/decisions/inbox/sebastian-marlin-bench.md create mode 100644 .squad/log/2026-08-14T21-43-12Z-marlin-int4-gemm-landed.md create mode 100644 .squad/orchestration-log/2026-08-14T21-43-12Z-chew.md create mode 100644 .squad/orchestration-log/2026-08-14T21-43-12Z-deckard.md create mode 100644 .squad/orchestration-log/2026-08-14T21-43-12Z-gaff.md create mode 100644 .squad/orchestration-log/2026-08-14T21-43-12Z-pris.md create mode 100644 .squad/orchestration-log/2026-08-14T21-43-12Z-sebastian.md diff --git a/.squad/agents/chew/history.md b/.squad/agents/chew/history.md index e49e3b1b96..924d300f02 100644 --- a/.squad/agents/chew/history.md +++ b/.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. diff --git a/.squad/agents/deckard/history.md b/.squad/agents/deckard/history.md index e0174a524f..ce21b79b89 100644 --- a/.squad/agents/deckard/history.md +++ b/.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. diff --git a/.squad/agents/gaff/history.md b/.squad/agents/gaff/history.md index f0f39cd80a..ba39fd810e 100644 --- a/.squad/agents/gaff/history.md +++ b/.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 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. diff --git a/.squad/agents/pris/history.md b/.squad/agents/pris/history.md index 838f3c63ad..a8bffb93da 100644 --- a/.squad/agents/pris/history.md +++ b/.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. diff --git a/.squad/agents/sebastian/history-archive.md b/.squad/agents/sebastian/history-archive.md index 56a29cef3b..307c504e57 100644 --- a/.squad/agents/sebastian/history-archive.md +++ b/.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`, `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. + diff --git a/.squad/agents/sebastian/history.md b/.squad/agents/sebastian/history.md index 51b77ddb38..1f65722ed9 100644 --- a/.squad/agents/sebastian/history.md +++ b/.squad/agents/sebastian/history.md @@ -10,124 +10,11 @@ - `cargo test --workspace` silently truncates on failure — always use `--no-fail-fast`. - Never commit `.squad/` files to external repos. -## Recent work (current wave, 2026-08-12) +## Historical context (2026-08-12→13 decode arc — detail in `history-archive.md`) -### 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`, `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. +Owned the native-CUDA capture + kernel work that took Muse-Glimmer-30B int4 decode **11.4 → 47.25 tok/s (now beats ORT ~40)**: capture arc to 23.13 (#854/#855), norm-cast-fold + parallel tree reduction to 40.21 (#860), cached Float16-staged int4 scales to 47.25 (#867), bf16 SwiGLU kernels (#871). Then FOUR independent confirmations that batch-1 decode is at its latency/dispatch floor, NOT bandwidth-bound: lower-bit-quant NO-GO (#885), multi-CTA megakernel NO-GO (#898), skip-RMSNorm standalone fold −1.5% NO-SHIP (kernel byte-exact SHIP, #903), norm→GEMV-prologue fusion NO-GO (#916). Standing lesson: at M=1, folding parallel work into a single-CTA reduction serializes what per-op spreads across 132 SMs; CUDA-graph replay already banks per-launch overhead, so node/launch fusion is not a decode lever. Detailed dated entries archived `.squad/agents/sebastian/history-archive.md`. +## 2026-08-14 — bf16 norm-into-GEMV-prologue fusion NO-GO (#916, MERGED) - **2026-08-14 (#916, MERGED):** bf16 norm-into-GEMV-prologue fusion measured NO-GO — −4.6% regression AND numeric divergence (≈token 38) under CUDA-graph replay; fp16 prologue reduction is single-warp-serial on the critical GEMV path. Finding-only (docs §8.7), nothing landed. **Fourth** independent confirmation of the batch-1 decode latency floor; norm→GEMV-prologue kill-gate CLOSED. ## 2026-08-14 — int4 decode GEMV bandwidth pass (#928, MERGED, main 8fe56961) @@ -142,3 +29,6 @@ Reframes the "launch-amortized floor" as HBM-bandwidth/dispatch-co-bound. Numeri byte-identical but not byte-exact per element (fails synthetic asymmetric-zp guard, near-zero cols only). **Ship opt-in default OFF** (`ONNX_GENAI_GEMV_FOLDSCALE=1`). Bigger single-GPU wins need a from-scratch Marlin int4 kernel (multi-week). + +## 2026-08-14 — Marlin int4 GEMM perf/capture harness + BEFORE/AFTER B* arc (#962 MERGED df6d3afb) +Built `marlin_bench` (`crates/onnx-genai-bench/src/bin/marlin_bench.rs`, `--features bench-native,cuda`): times a real `decode_verify` over M∈{1,2,4,8,16} + prefill L∈{128,512,1024}, prints median/p10/p90/max + device `compute_cap` so every number is arch-attributable. Ran all BEFORE→AFTER capture re-probes for Deckard's Marlin (Lever A): the glm-4-9b M=8 speculative-verify graph went **41 fragmented segments / capture B\*=8.76× → SINGLE whole-graph capture, ZERO unsupported nodes, B\*=2.16×** (arc 8.76→4.99→2.71→2.63→2.16×), prefill ~2× (glm 218→426 tok/s @L=1024, halving the vs-ORT gap 121×→62×), byte-identical greedy tokens throughout. Per-model honesty: glm (block-128) clean GO at 2.16×; qwen (block-32) capture fully solved but B\*≈4.7× — a denominator effect (fast tuned block-32 M=1), not a kernel bug. Also caught that the qwen exact-token parity test is **flaky ~25%** (asserts equality vs a nondeterministic tiled reference at a near-tie tok19); Marlin's fixed split-K reduction is the DETERMINISTIC path — harden the assertion, not a regression. Gates the #957 spec-capture CONDITIONAL-GO. Reviews: Chew 🟡 / Gaff 🟢. diff --git a/.squad/decisions-archive/2026-08.md b/.squad/decisions-archive/2026-08.md index 570b8e9237..d19fb15ac1 100644 --- a/.squad/decisions-archive/2026-08.md +++ b/.squad/decisions-archive/2026-08.md @@ -10656,3 +10656,85 @@ capture-safe OR replay cost scales ~K× not ≈1×); A 🟡 CONDITIONAL GO (full microbench lifts achieved DRAM 29% → ≥~55%; NO-GO if it stays <~40%). **Next step:** run B's throwaway `#[ignore]` Phase-0 capture-stability microbench before committing eng-weeks (doc §3). Drop merged & deleted: `roper-decode-remaining-levers.md`. + +## Archived by Scribe 2026-08-14T21-43-12Z (marlin-int4-gemm-landed batch) + +Moved from live decisions.md to make room for the consolidated Marlin int4 GEMM program entry (Lever A landed, #960/#961/#962). These are the three verbose "Last consolidated" chronicle lines and the detailed sub-entry bodies of the spec-capture CONDITIONAL-GO (#957) and Lever B Phase-0/Increment-0 NO-GO (#948/#949) — all now RESOLVED by the Marlin landing (see the live "Marlin int4 GEMM" entry). Live file keeps compact pointers. + +### Archived "Last consolidated" chronicle lines + +Last consolidated: 2026-08-14T17:46:35Z (Scribe spec-capture batch, local state; merged 1 inbox drop — deckard-spec-capture-design (#957, MERGED 2f0b62b3). KEY: **speculative decoding under CUDA-graph capture is CONDITIONAL-GO — the single condition is Marlin (Lever A), already the funded PRIMARY decode lever.** This REFINES (does not overturn) the Lever B NO-GO: the deciding number is acceptance break-even B\* = C_verify(M=K)/C_decode(M=1); today B\*=8.5 (guaranteed loss), post-Marlin estimate B\*≈1.2–2.0 (cleared by prompt-lookup + EAGLE-3/MTP → ~2–3× ceiling). Root cause of the ~80 ms M=8 floor is the 240 MatMulNBits generic GEMMs (M>1 abandons the tuned int4 GEMV), which Marlin fixes; two cheap residual fixes (SkipSimplifiedLayerNorm + GQA M>1 capture signatures) remain. Standing directive: do NOT fund the speculative verify-graph build ahead of Marlin; re-run the Increment-0 probe post-Marlin and fund only if measured B\*≤~2. HARD-GATE size: decisions.md was 40,842 B before merge — UNDER the 50 KB charter archive floor even after this ~4 KB merge → archived nothing (correct per Scribe charter: archive by SIZE, not age). NOTE: spawn prompt asked to archive entries older-than-30-days (age-based) because the file exceeds 20 KB; per Scribe charter age-based archiving is forbidden as primary criterion (it silently no-ops) and the real gate is 50 KB, so I archived nothing and say so. Committed on chore/scribe-spec-capture (main is protected — coordinator opens the PR). Histories: appended to deckard; checked chronicle + 15,360 B gates.) +Last consolidated: 2026-08-14T15:30:00Z (Scribe leverb-nogo batch, local state; merged 2 inbox drops — deckard-leverb-increment0 (#949) + deckard-leverb-phase0 (#948). KEY: **Lever B (capture-stable padded M=K verify graph) was TESTED and is a DECISIVE NO-GO**; **Lever A (Marlin int4 relayout, unconditional ~1.3–1.6×) promoted to the PRIMARY decode lever.** This SUPERSEDES the #938 "build Lever B first" recommendation. Phase-0 (#948): capture machinery stable (994/1000 replays, 90 tok/s) but raw M=K forward NotCapturable; eager M=8 = 6.77× M=1 (~80 ms floor, composition unknown) → gated re-test on Increment-0. Increment-0 (#949): built the capture-enablement overlay (persistent [1,K_max,vocab] logits binding + alloc-free M=K workspace + KV-symbol pin) → M=8 now captures, but **captured M=8 = 87.2 ms = 8.58× captured M=1 (10.2 ms)** — the ~80 ms floor PERSISTS under capture. Root cause: segments=41 at M=8 (vs 1 at M=1) because GroupQueryAttention/MatMulNBits/SkipSimplifiedLayerNormalization declare KernelCaptureUnsupported at M>1, forcing ~361 eager relaunches. Lever B is gated behind a deep kernel-capture-support program, not "reuse existing machinery". HARD-GATE size: decisions.md was 48,169 B (would exceed 50 KB after merge) → archived two older "Last consolidated" lines (08-14T09:18 & 04:09) + the detailed #928/#932/#935/#933/#938 "Decode perf REOPENED" sub-entry bodies to decisions-archive/2026-08.md (live keeps a compact standing summary + pointer) → well under 50 KB. NOTE: spawn prompt asked to archive by AGE (older-than-30-days); per Scribe charter I archived by SIZE and committed on chore/scribe-leverb-nogo (main is protected — coordinator opens the PR). Histories: appended to deckard; checked chronicle + 15,360 B gates.) +Last consolidated: 2026-08-14T09:57:00Z (Scribe decode-levers-scoped batch, local state; merged 2 inbox drops — roper-decode-remaining-levers (#938) + holden-zerocopy-linux-default (#936). KEY: (1) of decode's two remaining multi-week "big build" levers, **build Lever B first** — a capture-stable padded M=K verify graph that REPLAYS (attacks the dispatch binding; floor ≈1.0×, ceiling ~2–3×; unlocks prompt-lookup + EAGLE-3/MTP), gated on a cheap Phase-0 capture-stability probe; keep **Lever A** (Marlin int4 relayout, unconditional ~1.3–1.6×) funded as fallback/parallel, primary only if B's Phase-0 fails (#938 feasibility doc). (2) Zero-copy hybrid budget default is now **platform-conditional**: Windows/WDDM stays 256 MiB (real #864 aperture ceiling), Linux/non-Windows raised to **2 GiB** — #925 re-measured the ceiling as Windows/VidMm-specific (byte-identical to 6.795 GB on H200), unlocking a ~8× Linux memory-constrained win (67 vs ~8.5 tok/s over-budget); bounded on purpose, feature still opt-in default-OFF, override via `ONNX_GENAI_ZERO_COPY_HYBRID_BUDGET_BYTES` (#936, Holden). HARD-GATE size: decisions.md was 49,264 B (would exceed 50 KB after merge) → archived two older "Last consolidated" lines (08-13T18:57 & 14:45) + the detailed bodies of the #864 WDDM-negative and #888 eviction-order sections to decisions-archive/2026-08.md (live keeps pointers) → back under 50 KB. NOTE: spawn prompt asked to archive by AGE and push to main; per Scribe charter I archived by SIZE and committed on a chore/scribe branch (main is protected — coordinator opens the PR). Histories: appended to roper + holden; checked chronicle + 15,360 B gates.) + +### Archived detailed sub-entry bodies (spec-capture #957 + Lever B #948/#949) + +### 2026-08-14: Speculative-capture feasibility — CONDITIONAL-GO gated behind Marlin (design-only, #957 MERGED 2f0b62b3, Deckard) +**REFINEMENT (does not overturn) of the Lever B NO-GO above.** Design-only study (`docs/research/speculative-capture-feasibility.md`; +no kernels touched) answering Justin's "can we make speculative capture work". **Verdict: CONDITIONAL-GO — the single +condition is Marlin (Lever A), which is already the funded PRIMARY decode lever.** Speculative-capture is NO-GO on any +path that does not first fix the M>1 int4 GEMM, but it is NOT a separate multi-week kernel program: the one expensive +prerequisite IS Marlin. Lever B's kernel-capture-support program is largely *subsumed by* Marlin plus two cheap fixes. +**The deciding number — acceptance break-even B\* = C_verify(M=K)/C_decode(M=1), C_decode=10.2 ms:** today (measured, #949) +C_verify(M=8)=87.2 ms → **B\*=8.5** ⇒ needs >K max-accept ⇒ guaranteed loss (matches the #932/#935 KILL 0.47–0.74×@96%); +post-Marlin (ESTIMATE ~12–20 ms) → **B\*≈1.2–2.0** ⇒ cleared by prompt-lookup (B≈2–4) and EAGLE-3/MTP ⇒ lands the settled +~2–3× ceiling. **Root cause (code-grounded, no new GPU run):** the ~80 ms M=8 floor is dominated by the 240 `MatMulNBits` +generic GEMMs (M>1 abandons the tuned fused int4 GEMV for a slow portable 16×16 fp32-accumulate tiled GEMM, +`matmul_nbits.rs:4803-4869`), NOT GQA/norm (both live in the cheap 1.85 ms/row tail). Marlin (static grid, weights reused +across M rows) is capture-safe by construction and collapses the cliff. **Two residual capture blockers after Marlin, +both cheap & structurally capturable:** (1) `SkipSimplifiedLayerNorm`×80 — relax the conservative +`last_call_capture_safe=(num_groups==1)` gate to a warmed fixed `num_groups=K` signature (`normalization.rs:2306-2308`), +no kernel rewrite; (2) `GroupQueryAttention`×40 — pre-warm the fixed-M flash workspace + admit a warmed `q_seq=K` capture +signature (`group_query_attention.rs:842-1063,3099-3110`), flash path has no host read-back. **Staged plan with HARD +GATES (do not skip ahead):** Stage 1 = Marlin (already primary) + ONE added contract item: Marlin must advertise +capture-safe at fixed M=K (Gate 1 = Marlin's DRAM≥55% gate + capture-safe@M=K, no seam). Stage 2 = the two cheap GQA+norm +M>1 capture fixes (Gate 2 = seam census at M=8 collapses to ~1 segment). Stage 3 = **RE-RUN THIS EXACT Increment-0 probe** +for the true post-Marlin captured M=K wall (**Gate 3 fund/kill: measured B\*≤~2 ⇒ GO fund verify-graph; B\*≳~4 ⇒ NO-GO, +Marlin already banked its win**). Stage 4 (only if Gate 3 passes) = the actual speculative build: verify sub-graph + +capture-stable selective KV-commit (write K slots, move logical length host-side) + exact-greedy near-tie guard (#935) + +draft sources (Gate 4 = end-to-end >1.0× at measured acceptance with greedy identity). **Portability (Rule 11):** Marlin +is SM-version-specific (SM80+ cp.async/LOP3; SM80 vs SM90 relayout) ⇒ runtime arch guard + byte-identical split-K fallback +on 1: `GroupQueryAttention[KernelCaptureUnsupported]×40`, +`MatMulNBits[…]×240`, `SkipSimplifiedLayerNormalization[…]×80`. Those three families advertise CUDA-graph +capture ONLY at the M=1 decode shape, so the "captured" M=K forward degrades to ~361 eager relaunches — the +exact dispatch-bound regime the lever was meant to escape. **Decision: NO-GO for Lever B. Promote Lever A +(Marlin int4 relayout, unconditional ~1.3–1.6×, ~4–6 eng-weeks, no capture-support prerequisite) to the +PRIMARY decode lever.** Lever B is not dead but is **gated behind a kernel-capture-support program** (make +GQA / MatMulNBits / SkipSimplifiedLayerNormalization capture-safe at M>1 — a deep multi-family kernel +program, NOT the "~3–5 eng-weeks reuse existing machinery" the design assumed); once that lands, re-run +this exact probe. Deliverables: Increment-0 overlay + `docs/research/leverb-phase0-capture-probe.md` +(Increment-0 section). Drop merged & deleted: `deckard-leverb-increment0.md`. + +### 2026-08-14: Lever B Phase-0 capture-stability probe — NO-GO to unconditional commit; gated re-test on Increment-0 (#948, Deckard) +Ran the cheap Phase-0 capture-stability probe on glm-4-9b-int4 (H200) to answer Lever B's one load-bearing +question — can a fixed-shape padded M=K forward be captured, replay ~1 dispatch/verify across bucket growth, +and cost ≈ one M=1 replay? By the strict "PASS = all three" rule: **(b) PASS, (a) FAIL, (c) UNMEASURED → +NO-GO** to an unconditional multi-week commit. **(b) capture stability — PASS** (1000 M=1 steps → captures=3, +replays=994, invalidations=3; 90.3 tok/s; no thrash). **(a) instantiates capture-safe — FAIL, +non-fundamental** (raw M=K `try_capture` returns `NotCapturable`: `[1,K,vocab]` logits materialized + +~722 transient workspace allocs — both the missing Increment-0 work, not a kernel veto). **(c) captured M=K +wall ≈ M=1 wall — UNMEASURED (blocked by a)**; eager proxy M=1 13.5 ms / M=2 80.4 ms / M=8 91.5 ms = 6.77×, +a STEP FUNCTION (M=1→M=2 cliff 66.9 ms, near-flat M=2→M=8 tail 1.85 ms/row) whose composition (dispatch vs +generic-GEMM+alloc) was the pivotal unknown. **Recommendation (gate, days not weeks):** fund Increment-0 +(the three fixes) then re-run — captured M=8 ≈ M=1 → GO; ~80 ms floor persists → NO-GO, promote Lever A. +(Increment-0 re-test above resolved this: floor persists → NO-GO.) Deliverables: `#[ignore]`d un-wired probe ++ `docs/research/leverb-phase0-capture-probe.md`. Drop merged & deleted: `deckard-leverb-phase0.md`. diff --git a/.squad/decisions.md b/.squad/decisions.md index 1871c74256..52119e1114 100644 --- a/.squad/decisions.md +++ b/.squad/decisions.md @@ -1,10 +1,8 @@ # Decisions — live standing directives -Last consolidated: 2026-08-14T17:46:35Z (Scribe spec-capture batch, local state; merged 1 inbox drop — deckard-spec-capture-design (#957, MERGED 2f0b62b3). KEY: **speculative decoding under CUDA-graph capture is CONDITIONAL-GO — the single condition is Marlin (Lever A), already the funded PRIMARY decode lever.** This REFINES (does not overturn) the Lever B NO-GO: the deciding number is acceptance break-even B\* = C_verify(M=K)/C_decode(M=1); today B\*=8.5 (guaranteed loss), post-Marlin estimate B\*≈1.2–2.0 (cleared by prompt-lookup + EAGLE-3/MTP → ~2–3× ceiling). Root cause of the ~80 ms M=8 floor is the 240 MatMulNBits generic GEMMs (M>1 abandons the tuned int4 GEMV), which Marlin fixes; two cheap residual fixes (SkipSimplifiedLayerNorm + GQA M>1 capture signatures) remain. Standing directive: do NOT fund the speculative verify-graph build ahead of Marlin; re-run the Increment-0 probe post-Marlin and fund only if measured B\*≤~2. HARD-GATE size: decisions.md was 40,842 B before merge — UNDER the 50 KB charter archive floor even after this ~4 KB merge → archived nothing (correct per Scribe charter: archive by SIZE, not age). NOTE: spawn prompt asked to archive entries older-than-30-days (age-based) because the file exceeds 20 KB; per Scribe charter age-based archiving is forbidden as primary criterion (it silently no-ops) and the real gate is 50 KB, so I archived nothing and say so. Committed on chore/scribe-spec-capture (main is protected — coordinator opens the PR). Histories: appended to deckard; checked chronicle + 15,360 B gates.) -Last consolidated: 2026-08-14T15:30:00Z (Scribe leverb-nogo batch, local state; merged 2 inbox drops — deckard-leverb-increment0 (#949) + deckard-leverb-phase0 (#948). KEY: **Lever B (capture-stable padded M=K verify graph) was TESTED and is a DECISIVE NO-GO**; **Lever A (Marlin int4 relayout, unconditional ~1.3–1.6×) promoted to the PRIMARY decode lever.** This SUPERSEDES the #938 "build Lever B first" recommendation. Phase-0 (#948): capture machinery stable (994/1000 replays, 90 tok/s) but raw M=K forward NotCapturable; eager M=8 = 6.77× M=1 (~80 ms floor, composition unknown) → gated re-test on Increment-0. Increment-0 (#949): built the capture-enablement overlay (persistent [1,K_max,vocab] logits binding + alloc-free M=K workspace + KV-symbol pin) → M=8 now captures, but **captured M=8 = 87.2 ms = 8.58× captured M=1 (10.2 ms)** — the ~80 ms floor PERSISTS under capture. Root cause: segments=41 at M=8 (vs 1 at M=1) because GroupQueryAttention/MatMulNBits/SkipSimplifiedLayerNormalization declare KernelCaptureUnsupported at M>1, forcing ~361 eager relaunches. Lever B is gated behind a deep kernel-capture-support program, not "reuse existing machinery". HARD-GATE size: decisions.md was 48,169 B (would exceed 50 KB after merge) → archived two older "Last consolidated" lines (08-14T09:18 & 04:09) + the detailed #928/#932/#935/#933/#938 "Decode perf REOPENED" sub-entry bodies to decisions-archive/2026-08.md (live keeps a compact standing summary + pointer) → well under 50 KB. NOTE: spawn prompt asked to archive by AGE (older-than-30-days); per Scribe charter I archived by SIZE and committed on chore/scribe-leverb-nogo (main is protected — coordinator opens the PR). Histories: appended to deckard; checked chronicle + 15,360 B gates.) -Last consolidated: 2026-08-14T09:57:00Z (Scribe decode-levers-scoped batch, local state; merged 2 inbox drops — roper-decode-remaining-levers (#938) + holden-zerocopy-linux-default (#936). KEY: (1) of decode's two remaining multi-week "big build" levers, **build Lever B first** — a capture-stable padded M=K verify graph that REPLAYS (attacks the dispatch binding; floor ≈1.0×, ceiling ~2–3×; unlocks prompt-lookup + EAGLE-3/MTP), gated on a cheap Phase-0 capture-stability probe; keep **Lever A** (Marlin int4 relayout, unconditional ~1.3–1.6×) funded as fallback/parallel, primary only if B's Phase-0 fails (#938 feasibility doc). (2) Zero-copy hybrid budget default is now **platform-conditional**: Windows/WDDM stays 256 MiB (real #864 aperture ceiling), Linux/non-Windows raised to **2 GiB** — #925 re-measured the ceiling as Windows/VidMm-specific (byte-identical to 6.795 GB on H200), unlocking a ~8× Linux memory-constrained win (67 vs ~8.5 tok/s over-budget); bounded on purpose, feature still opt-in default-OFF, override via `ONNX_GENAI_ZERO_COPY_HYBRID_BUDGET_BYTES` (#936, Holden). HARD-GATE size: decisions.md was 49,264 B (would exceed 50 KB after merge) → archived two older "Last consolidated" lines (08-13T18:57 & 14:45) + the detailed bodies of the #864 WDDM-negative and #888 eviction-order sections to decisions-archive/2026-08.md (live keeps pointers) → back under 50 KB. NOTE: spawn prompt asked to archive by AGE and push to main; per Scribe charter I archived by SIZE and committed on a chore/scribe branch (main is protected — coordinator opens the PR). Histories: appended to roper + holden; checked chronicle + 15,360 B gates.) +Last consolidated: 2026-08-14T21:43:12Z (Scribe marlin-int4-gemm-landed batch, local state; merged 5 inbox drops — deckard-marlin-kernel (#960 7774ec5b), pris-marlin-numerics (#961 401af46f), sebastian-marlin-bench (#962 df6d3afb), chew-marlin-numerics-review, gaff-marlin-quality-review. KEY: **Marlin fp16×int4 tensor-core GEMM (Lever A) LANDED — the M=8 speculative-verify graph went from 41 fragmented segments / capture B\* = 8.76× (hard NO-GO) to a SINGLE whole-graph capture with ZERO KernelCaptureUnsupported nodes at B\* = 2.16× (glm-4-9b canonical gate), byte-identical greedy tokens throughout, prefill ~2×. This GATES the #957 speculative-capture CONDITIONAL-GO — Lever A delivered.** Reviews: Chew 🟡 APPROVE-WITH-NOTES (11/11 GPU parity tests pass, no correctness bug), Gaff 🟢 APPROVE (zero blocking defects, Rule 11 PASS). Per-model honesty: glm (block-128) clean practical GO at 2.16×; qwen (block-32) capture fully solved (whole-graph, zero seams, byte-identical eager) but B\* ≈ 4.7× — an honest denominator effect (fast tuned block-32 M=1 inflates the ratio), a drafting-depth follow-up, not a kernel bug. HARD-GATE size: decisions.md was 45,991 B (would exceed 51,200 B after this merge) → archived the three verbose "Last consolidated" lines + the detailed #957/#948/#949 spec-capture & Lever B sub-entry bodies (now RESOLVED by this landing) to decisions-archive/2026-08.md (live keeps compact pointers) → under 50 KB. NOTE: spawn prompt asked to archive entries older-than-30-days (age-based); per Scribe charter age-based archiving is forbidden as primary criterion (it silently no-ops) and the real gate is 50 KB, so I archived by SIZE and say so. Committed on chore/scribe-marlin (main is protected — coordinator opens the PR). Histories: appended to deckard/sebastian/pris/chew/gaff; checked chronicle + 15,360 B gates.) -Earlier 'Last consolidated' chronicle lines (2026-08-11/12 six entries, 2026-08-13 seven entries, 2026-08-14T04:09 & 09:18) archived to `.squad/decisions-archive/2026-08.md`. +Earlier 'Last consolidated' chronicle lines (2026-08-11/12 six entries, 2026-08-13 seven entries, 2026-08-14T04:09 & 09:18, and the three 2026-08-14T09:57/15:30/17:46 lines) archived to `.squad/decisions-archive/2026-08.md`. Standing governance rules and active directives. Full narrative is archived in `.squad/decisions-archive/2026-07.md`, `.squad/decisions-archive/2026-08.md`, and older `.squad/decisions/archive/` files. @@ -79,75 +77,23 @@ so Lever A (Marlin) is promoted to the PRIMARY decode lever.** Detailed #928/#93 bodies and the #938 "build Lever B first" entry archived to `.squad/decisions-archive/2026-08.md` ("Archived by Scribe 2026-08-14T15:30Z"). -### 2026-08-14: Speculative-capture feasibility — CONDITIONAL-GO gated behind Marlin (design-only, #957 MERGED 2f0b62b3, Deckard) -**REFINEMENT (does not overturn) of the Lever B NO-GO above.** Design-only study (`docs/research/speculative-capture-feasibility.md`; -no kernels touched) answering Justin's "can we make speculative capture work". **Verdict: CONDITIONAL-GO — the single -condition is Marlin (Lever A), which is already the funded PRIMARY decode lever.** Speculative-capture is NO-GO on any -path that does not first fix the M>1 int4 GEMM, but it is NOT a separate multi-week kernel program: the one expensive -prerequisite IS Marlin. Lever B's kernel-capture-support program is largely *subsumed by* Marlin plus two cheap fixes. -**The deciding number — acceptance break-even B\* = C_verify(M=K)/C_decode(M=1), C_decode=10.2 ms:** today (measured, #949) -C_verify(M=8)=87.2 ms → **B\*=8.5** ⇒ needs >K max-accept ⇒ guaranteed loss (matches the #932/#935 KILL 0.47–0.74×@96%); -post-Marlin (ESTIMATE ~12–20 ms) → **B\*≈1.2–2.0** ⇒ cleared by prompt-lookup (B≈2–4) and EAGLE-3/MTP ⇒ lands the settled -~2–3× ceiling. **Root cause (code-grounded, no new GPU run):** the ~80 ms M=8 floor is dominated by the 240 `MatMulNBits` -generic GEMMs (M>1 abandons the tuned fused int4 GEMV for a slow portable 16×16 fp32-accumulate tiled GEMM, -`matmul_nbits.rs:4803-4869`), NOT GQA/norm (both live in the cheap 1.85 ms/row tail). Marlin (static grid, weights reused -across M rows) is capture-safe by construction and collapses the cliff. **Two residual capture blockers after Marlin, -both cheap & structurally capturable:** (1) `SkipSimplifiedLayerNorm`×80 — relax the conservative -`last_call_capture_safe=(num_groups==1)` gate to a warmed fixed `num_groups=K` signature (`normalization.rs:2306-2308`), -no kernel rewrite; (2) `GroupQueryAttention`×40 — pre-warm the fixed-M flash workspace + admit a warmed `q_seq=K` capture -signature (`group_query_attention.rs:842-1063,3099-3110`), flash path has no host read-back. **Staged plan with HARD -GATES (do not skip ahead):** Stage 1 = Marlin (already primary) + ONE added contract item: Marlin must advertise -capture-safe at fixed M=K (Gate 1 = Marlin's DRAM≥55% gate + capture-safe@M=K, no seam). Stage 2 = the two cheap GQA+norm -M>1 capture fixes (Gate 2 = seam census at M=8 collapses to ~1 segment). Stage 3 = **RE-RUN THIS EXACT Increment-0 probe** -for the true post-Marlin captured M=K wall (**Gate 3 fund/kill: measured B\*≤~2 ⇒ GO fund verify-graph; B\*≳~4 ⇒ NO-GO, -Marlin already banked its win**). Stage 4 (only if Gate 3 passes) = the actual speculative build: verify sub-graph + -capture-stable selective KV-commit (write K slots, move logical length host-side) + exact-greedy near-tie guard (#935) + -draft sources (Gate 4 = end-to-end >1.0× at measured acceptance with greedy identity). **Portability (Rule 11):** Marlin -is SM-version-specific (SM80+ cp.async/LOP3; SM80 vs SM90 relayout) ⇒ runtime arch guard + byte-identical split-K fallback -on 1: `GroupQueryAttention[KernelCaptureUnsupported]×40`, -`MatMulNBits[…]×240`, `SkipSimplifiedLayerNormalization[…]×80`. Those three families advertise CUDA-graph -capture ONLY at the M=1 decode shape, so the "captured" M=K forward degrades to ~361 eager relaunches — the -exact dispatch-bound regime the lever was meant to escape. **Decision: NO-GO for Lever B. Promote Lever A -(Marlin int4 relayout, unconditional ~1.3–1.6×, ~4–6 eng-weeks, no capture-support prerequisite) to the -PRIMARY decode lever.** Lever B is not dead but is **gated behind a kernel-capture-support program** (make -GQA / MatMulNBits / SkipSimplifiedLayerNormalization capture-safe at M>1 — a deep multi-family kernel -program, NOT the "~3–5 eng-weeks reuse existing machinery" the design assumed); once that lands, re-run -this exact probe. Deliverables: Increment-0 overlay + `docs/research/leverb-phase0-capture-probe.md` -(Increment-0 section). Drop merged & deleted: `deckard-leverb-increment0.md`. - -### 2026-08-14: Lever B Phase-0 capture-stability probe — NO-GO to unconditional commit; gated re-test on Increment-0 (#948, Deckard) -Ran the cheap Phase-0 capture-stability probe on glm-4-9b-int4 (H200) to answer Lever B's one load-bearing -question — can a fixed-shape padded M=K forward be captured, replay ~1 dispatch/verify across bucket growth, -and cost ≈ one M=1 replay? By the strict "PASS = all three" rule: **(b) PASS, (a) FAIL, (c) UNMEASURED → -NO-GO** to an unconditional multi-week commit. **(b) capture stability — PASS** (1000 M=1 steps → captures=3, -replays=994, invalidations=3; 90.3 tok/s; no thrash). **(a) instantiates capture-safe — FAIL, -non-fundamental** (raw M=K `try_capture` returns `NotCapturable`: `[1,K,vocab]` logits materialized + -~722 transient workspace allocs — both the missing Increment-0 work, not a kernel veto). **(c) captured M=K -wall ≈ M=1 wall — UNMEASURED (blocked by a)**; eager proxy M=1 13.5 ms / M=2 80.4 ms / M=8 91.5 ms = 6.77×, -a STEP FUNCTION (M=1→M=2 cliff 66.9 ms, near-flat M=2→M=8 tail 1.85 ms/row) whose composition (dispatch vs -generic-GEMM+alloc) was the pivotal unknown. **Recommendation (gate, days not weeks):** fund Increment-0 -(the three fixes) then re-run — captured M=8 ≈ M=1 → GO; ~80 ms floor persists → NO-GO, promote Lever A. -(Increment-0 re-test above resolved this: floor persists → NO-GO.) Deliverables: `#[ignore]`d un-wired probe -+ `docs/research/leverb-phase0-capture-probe.md`. Drop merged & deleted: `deckard-leverb-phase0.md`. +### 2026-08-14: Marlin fp16×int4 tensor-core GEMM (Lever A) LANDED — spec-capture CONDITIONAL-GO now GO in practice (#960/#961/#962 MERGED; Chew 🟡 / Gaff 🟢) + +**RESOLVES the spec-capture CONDITIONAL-GO (#957) and the Lever B NO-GO (#948/#949) — the detailed bodies of those three entries are archived to `.squad/decisions-archive/2026-08.md` ("Archived by Scribe 2026-08-14T21-43-12Z").** Marlin (Lever A), the single funded condition of the #957 CONDITIONAL-GO, is built, validated, reviewed, and merged across 3 PRs. + +**Deckard — kernel (#960, 7774ec5b):** from-scratch SM80 `mma.sync.m16n8k16` fused fp16×int4 tensor-core GEMM in a new module `crates/onnx-runtime-ep-cuda/src/kernels/marlin_gemm.rs` (adapt-not-vendor: ONNX-native N-major nibble packing, asymmetric nibble zero-points, group sizes 16/32/64/128; per-group scale applied AFTER the tensor-core accumulate so the fp32 accumulator never carries a K-varying scale; raw inline `mma.sync` PTX, no ``; no upstream source copied ⇒ no third-party 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 arch guard, byte-identical tiled fallback on 2.16 tokens/verify wins. + +**Per-model honesty (CORRECTED attribution — Deckard update-10 supersedes the update-8 error):** the fused gate_up SwiGLU split-K path requires `block_size==32` (`matmul_nbits.rs:5782`). **glm-4-9b = block-128 ⇒ the fused gate_up node is NEVER formed; glm runs gate/up as separate `MatMulNBits` nodes through the general Marlin split-K dispatch — so the `3735d57e` gate_up-fusion commit is a NO-OP for glm, and glm's 2.63→2.16× was ENTIRELY the `4abe4e57` general small-M split-K retune** (update-8's claim that gate_up fusion drove the glm number was WRONG; corrected here). **qwen2.5-14b = block-32 ⇒ the fused gate_up node forms and fires.** qwen capture is fully solved (segments=1, whole-graph, zero unsupported — all 48 GQA + 96 LN capture-safe, byte-identical eager parity) but B\* ≈ 4.62–4.72× — an honest **denominator effect**: capture accelerates qwen's fast tuned block-32 M=1 GEMV 34% (11.4→7.5 ms) more than the heavier 14B/48-layer M=8 verify 18% (43.1→35.5 ms), inflating the ratio. Closing qwen's B\* is a drafting-depth follow-up, NOT a frozen-kernel fix. + +**Pris — numerics gate (#961, 401af46f):** a reusable f64 dequant→GEMM oracle test any int4 GEMM must pass; dequant `(code-zp)·scale` accumulated in f64, sharing the candidate's fp16-rounded activations/scale so the residual isolates only accumulation + fp16 output rounding. Justified tolerance envelope `abs = max(max_out·4e-3, 4e-3)`, `rel = 5e-2` (denominator floored at `max(1e-1, 3e-2·max_out)`); the tiled baseline sits ~8–10× inside abs and ~6× inside rel (≈1 fp16 ULP), and Marlin must land in the same envelope. Harness lesson: pre-zero the output + `synchronize()` after execute (a stale device-pool read once faked an abs-81.8 divergence — not a kernel bug). + +**Sebastian — perf/capture harness (#962, df6d3afb):** new `marlin_bench` bin (`crates/onnx-genai-bench/src/bin/marlin_bench.rs`, `--features bench-native,cuda`) times a real `decode_verify` over M∈{1,2,4,8,16} + prefill L∈{128,512,1024}, printing median/p10/p90/max + device `compute_cap` so every number is arch-attributable; it ran all BEFORE→AFTER capture re-probes that produced the B\* arc above. It also surfaced that the qwen exact-token-stream parity test is **flaky (~25%)** because it asserts equality vs a nondeterministic tiled reference at a near-tie (tok19) — Marlin's fixed split-K reduction is the DETERMINISTIC path, so harden the assertion (tolerance / deterministic oracle), not a Marlin regression. + +**Reviews.** Chew (numerics) 🟡 **APPROVE-WITH-NOTES** — 11/11 in-crate GPU parity tests pass on H200 (incl. f64-oracle parity, split-K parity + determinism, repack roundtrip all group sizes), no correctness bug; non-blocking notes: (N1) "byte-identical greedy tokens" is a soft argmax-stability guarantee, not a numeric invariant — keep the flag opt-in; (N2) hard Marlin launch errors are silently swallowed into the tiled fallback — log/count them so a real fault can't hide behind the slow path; (N3) nibble-packed int zero-points only (matches every existing int4 tiled kernel — not a regression). Gaff (quality) 🟢 **APPROVE** — zero blocking defects; Rule 11 portability PASS (opt-in + SM80 guard + byte-identical fallback), env-var honesty PASS (both knobs read & wired; `verify_documented_env_vars.py` EXIT 0), capture-safety valve family sound (cache-hit warm / cold-miss-during-capture returns Err → caller falls back / pre-warm populates the caches), `cargo fmt` + the exact CUDA and engine clippy gates clean; one trivial `cfg(test)`-only `clippy::unusual_byte_groupings` note (outside the CUDA CI gate, does not break CI). + +**Standing directive (updated — supersedes the #957 pre-Marlin "do not fund ahead of Marlin" directive):** Marlin (Lever A) is DELIVERED and the spec-capture condition is MET for glm at B\* = 2.16×; the #957 Increment-0 re-probe is DONE — this landing IS the post-Marlin measurement. Now unblocked (per #957 Stage 4): fund the actual speculative build (verify sub-graph + capture-stable selective KV-commit + exact-greedy near-tie guard #935 + draft sources) — glm is a practical GO; deeper/stronger drafting is the lever to bring qwen's B\* into range. Marlin stays opt-in default-OFF; before any default flip, close Chew's N1/N2 (keep flag opt-in until more prompt/model coverage; add fallback logging/counter). Drops merged & deleted: deckard-marlin-kernel.md, pris-marlin-numerics.md, sebastian-marlin-bench.md, chew-marlin-numerics-review.md, gaff-marlin-quality-review.md. ## DeepSeek-V2 native path = standard RotaryEmbedding + Attention + QMoE (golden lock #914, 2026-08-14) diff --git a/.squad/decisions/inbox/deckard-marlin-kernel.md b/.squad/decisions/inbox/deckard-marlin-kernel.md deleted file mode 100644 index 72d91e5810..0000000000 --- a/.squad/decisions/inbox/deckard-marlin-kernel.md +++ /dev/null @@ -1,503 +0,0 @@ -### 2026-08-14: Marlin int4 tensor-core GEMM — design + Stage-2 measured numbers - -**By:** Deckard (Systems Dev — CUDA/decode-performance) -**Branch:** squad/marlin-kernel (draft PR, refs #957) -**Status:** Stage 2 in progress — kernel correct + validated + capture-safe grid; performance tuning ongoing. - -**What:** -Landed a from-scratch SM80 `mma.sync.m16n8k16` fused fp16×int4 tensor-core GEMM in a -new module `crates/onnx-runtime-ep-cuda/src/kernels/marlin_gemm.rs` (kernel + repack -isolated from the 475 KB `matmul_nbits.rs`). It targets the M>1 path of -`com.microsoft::MatMulNBits` (today's portable 16×16 CUDA-core tiled GEMM) and, -conditionally, the M=1 decode GEMV. - -Design decisions: -- **Adapt, not vendor.** Original kernel adapting Marlin's core ideas (repacked - per-lane weight layout; **per-group scale applied AFTER the tensor-core - accumulate** so the fp32 accumulator never carries a K-varying scale) to our - ONNX-native format: N-major nibble packing, even-K low nibble, **asymmetric - nibble zero-points**, group sizes 16/32/64/128. Upstream IST-DASLab/vLLM marlin - assume symmetric GPTQ layouts and depend on ``/`crt/` headers - unavailable to our NVRTC-string path, so a native kernel is both correct and - simpler than translating our weights into their format. No upstream source - copied ⇒ no third-party LICENSE vendoring required; lineage credited in the - module header. -- **Raw `mma.sync` inline PTX** (no ``), consistent with the file's existing - LOP3 asm — needs only `cuda_fp16.h` (present); `crt/mma.h` is absent here. -- **Repack** (`repack_int4_weights`): bijective reorder of packed int4 into an - 8-column-interleaved per-lane tensor-core layout so a 32-lane warp reads one - contiguous 64-byte weight chunk per K slice. Same byte count as source packed - (a reordering, not an expansion). It is an added packaging step; the current - layout is untouched → two layouts coexist per Rule 11. -- **Rule 11 portability:** `device_supports_marlin()` SM80 arch guard; callers fall - back byte-for-byte to the current CUDA-core tiled GEMM on 1 win vs the tiled GEMM will be measured apples-to-apples through the -op after wiring (Stage 3), coordinating with Sebastian (squad/marlin-bench). - -**Why:** Primary decode-perf lever (collapse the ~67 ms M=1→M=2 cliff) and the -enabler of capture-stable speculative decoding. Building standalone-validated -(correctness vs f64 + measured bandwidth) before wiring into the op, per the staging -plan. - ---- - -### 2026-08-14 (update): Stage 3–4 landed — M>1 wired, capture support advertised - -**Status:** Stage 3 (op wiring) + Stage 4 (capture support) COMPLETE and validated on -H200. M=1 stays on the existing GEMV (precondition still unmet). Head: d722705c. - -**What landed:** -- **Wired Marlin into the `MatMulNBits` M>1 dispatch** (`matmul_nbits.rs`, the - `if m > 1` seam only + the `capture_support` doc). Opt-in via - `ONNX_GENAI_MARLIN_M_GT_1=1` (default OFF, Rule-11 tier-scoped). Eligibility: - SM80+, int4, no fused SwiGLU/RMSNorm epilogue yet, K divisible by 16 and group. - Any ineligibility or launch error falls through **byte-compatibly** to the - portable tiled GEMM. -- **Module-level repack cache** (`ensure_repacked` in `marlin_gemm.rs`) keyed by - (packed_ptr, ordinal, N, K, group). Immutable initializer weights are repacked - on-device **once** during warmup; captured replays hit the warm cache with no - allocation. A cold miss during capture is rejected (safety valve → caller falls - back). Dispatch stores the cache-warm flag into `last_call_capture_safe`, so - **cold call = not-capture-safe, warm replay = capture-safe** — this is what - unlocks capture-stable speculative verify at M>1. -- Op-level test `marlin_m_gt_1_op_parity_and_capture_safety`: real op, flag on, - vs f64 oracle (asymmetric zp) — **worst_abs=0.013 ≪ tol=0.67**; asserts - cold=not-safe / warm=safe + byte-identical replay. PASSES on GPU. - -**Measured — apples-to-apples wall vs tiled through the op** -(`marlin_m_gt_1_op_wall_vs_tiled`, H200, K=5120 N=13824 block=128, median µs): - -| M | tiled µs | marlin µs | speedup | -|---|---|---|---| -| 1 | 69.5 | 69.0 | 1.01× (M=1 stays on GEMV) | -| 2 | 478.0 | 140.0 | **3.41×** | -| 4 | 494.1 | 138.5 | 3.57× | -| 8 | 536.4 | 176.9 | 3.03× | -| 16 | 658.1 | 251.2 | 2.62× | -| 32 | 1076.7 | 385.0 | 2.80× | -| 64 | 2010.6 | 821.6 | 2.45× | -| 128 | 3851.5 | 1357.9 | 2.84× | - -**Cliff collapse (the mission payoff):** the tiled path jumps 69.5 µs (M=1 GEMV) -→ 478 µs (M=2 tiled) = a **6.9× cliff**. Marlin M=2 = 140 µs, collapsing the -M=1→M=2 cliff to **~2×** (69.5→140 µs). Marlin is **2.4–3.6× faster than the -tiled GEMM across every M>1**. - -**M=1 decision (unchanged, honest):** standalone M=1 5120² is ~3.1% HBM peak and -through-op M=1 Marlin ties the GEMV (69 vs 69.5 µs) — it does NOT clear the -feasibility §3 ≥40% precondition, so **M=1 remains on the existing GEMV**. Not -silently dropped: reported. - -**Next:** extend fused epilogues (rmsnorm-prologue first, then swiglu/gate_up/down) -so those M>1 entry points can also use Marlin; small-M perf levers (split-K / -cp.async) to push M=1 toward the precondition; Stage 5 e2e on glm-4-9b-int4 + -qwen2.5-14b (coordinate with Sebastian on squad/marlin-bench). - ---- - -### 2026-08-14 (update 2): fused RMS-norm epilogue + Stage-5 e2e PASS - -**Status:** Fused rmsnorm-prologue path landed; e2e parity on both target models -PASSES. Head: a3ece463. - -**What landed:** -- **Fused RMS-norm prologue for Marlin M>1** (first of the fused epilogues that - MatMulNBits must preserve). Stages the per-token normalized activation into - scratch via the existing `launch_rmsnorm_prefill` (byte-identical to the - standalone prologue the tiled path uses), then runs Marlin over the normalized - rows. Scratch alloc keeps it off the capture contract (like the tiled rmsnorm - prefill; prefill is outside the decode graph). Falls through to tiled on - ineligibility/error. Test `marlin_m_gt_1_rmsnorm_op_parity` vs the trusted - fused-tiled path: **worst_abs=0.0156 ≪ tol=0.90**. -- **Stage-5 e2e** (`crates/onnx-genai-engine/tests/marlin_m_gt_1_e2e.rs`, native - CUDA): runs each real model twice through the engine — tiled (flag off) then - Marlin (flag on) — and diffs the greedy token stream. Prefill runs MatMulNBits - at M=prompt-length (Marlin path); decode stays on the M=1 GEMV. - - **glm-4-9b-int4:** 24 greedy tokens **byte-identical** tiled vs Marlin. - - **qwen2.5-14b-instruct-int4 (asymmetric zp):** 24 greedy tokens - **byte-identical** tiled vs Marlin. - -**Remaining (perf-focused, mission = full production kernel):** swiglu/gate_up/down -fused epilogues; small-M levers (split-K to fill SMs, cp.async multistage) to lift -M=1 DRAM% toward the feasibility precondition and further widen the M>1 win. -Coordinate with Sebastian (squad/marlin-bench) on the ORT e2e + capture re-probe. - ---- - -### 2026-08-14 (update 3): P1 COMPLETE — full M>1 coverage (gate_up SwiGLU fused), zero tiled fallbacks - -**Status:** Every hot MatMulNBits node in glm-4-9b and qwen2.5-14b now runs -through Marlin at M>1. Zero tiled M>1 fallbacks. Head: 1154c977. - -**What landed:** -- **gate_up SwiGLU MLP fusion at M>1** (`try_launch_marlin_gate_up_prefill`, - wired into `run_f16_gate_up_swiglu`): paired gate/up Marlin int4 GEMMs + - the identical `launch_silu_mul_f16_raw` epilogue the tiled path uses, with an - optional RMS-norm prologue (into pooled scratch). This is the MLP — the bulk - of prefill/verify cost. down_proj is a plain tall-skinny MatMulNBits already - covered by the plain `gemm_marlin_int4` path. -- **Capture-safe scratch pool** (`ensure_scratch` in marlin_gemm.rs): size- and - slot-keyed module-global pool (FIFO cap 256, rejects cold-miss during capture) - so warm replays of the gate_up + rmsnorm paths allocate NOTHING → capture-safe. - `try_launch_marlin_gemm_rmsnorm` was migrated onto it (now returns the warm flag - instead of always allocating), so the rmsnorm prefill is now capture-safe on - warm replay too. -- **Tracer-driven per-node coverage audit** (e2e test): runs prefill with the flag - on, tallies per-op `kernel_variant`, asserts zero `gemm_f16_tiled{,_rmsnorm}` / - `gate_up_swiglu{,_rmsnorm}_prefill` at M>1. -- **Unit tests:** gate_up SwiGLU M>1 Marlin-vs-tiled parity + capture-safety - (plain / rmsnorm / decomposed-silu): worst_abs 0.0001–0.008 vs tiled, warm - replays capture-safe and byte-identical. - -**Measured coverage audit (H200, GPU 7, Marlin M>1 enabled):** -- **glm-4-9b-int4:** 240 `gemm_marlin_int4`, **0 tiled M>1**. (glm's MLP is - separate plain MatMulNBits nodes → covered by the plain Marlin path.) -- **qwen2.5-14b-int4 (asym zp):** 240 `gemm_marlin_int4` + 48 - `gate_up_swiglu_marlin_prefill` + 1 `gemm_marlin_int4_rmsnorm` = **289 Marlin - M>1, 0 tiled**. -- Greedy token streams remain **byte-identical** tiled-vs-Marlin on both (24 tokens). - -**Nodes that still can't go through Marlin at M>1:** NONE for the hot path. -Remaining non-Marlin variants observed are all legitimately M=1 decode -(`gemv_f16_*`, `gate_up_swiglu_rmsnorm_fused`) or non-MatMulNBits (attention, -rmsnorm). M=1 stays on the GEMV by design (has not cleared the ≥40% DRAM -precondition). - -**→ Unblocks Sebastian's Increment-0 capture re-probe** with -`ONNX_GENAI_MARLIN_M_GT_1=1` for the decisive segments→1 and captured-M=8/M=1 B*. - -**Next (P2):** split-K (fill SMs at small M) + cp.async multistage to lift the -M>1 ratio; honest re-run at the M=1 DRAM precondition (keep M=1 on GEMV unless -it truly clears ≥40% DRAM). - ---- - -## Update 4 — P2: split-K (widen the win) + M=1 re-run + cp.async finding - -**cp.async multistage — REJECTED (measured, not assumed).** Implemented a -2-slice/128B cp.async.cg pipelined loader (all 32 lanes issue, ring-slot indexed -by `fetch % STAGES`). Passed f64-oracle parity but measured **SLOWER** across all -M (M=1 N=5120 89→114µs; N=13824 100→130µs). Root cause: the small-M bottleneck is -**occupancy** (too few blocks in flight to saturate DRAM), not per-load latency, -so hiding load latency added barrier overhead without touching the real limiter. -Reverted the whole experiment. - -**split-K — LANDED (opt-in `ONNX_GENAI_MARLIN_SPLITK=1`, default OFF).** Partition -the K/group range across `grid.z`; each z accumulates its groups' `frag*scale` -into fp32 `partials[split_k,M,N]` (no bias); a reduce kernel sums the z-slices in -**fixed order** (deterministic ⇒ capture-stable) then applies fold_bias → fp16. -`choose_split_k` is conservative: only splits when the base grid is under ~2 waves -(so it NEVER regresses the already-occupied M≥32 / large-N cases), each split owns -≥2 groups, factor ≤8. New CUDA in `marlin_gemm.rs` (module untouched elsewhere); -wired into the plain + rmsnorm M>1 paths via `maybe_launch_marlin_splitk`; -partials from the capture-safe scratch pool (slot 4). - -**Measured (H200 GPU 3, median µs, split-K vs direct kernel):** -| shape | direct | best split-K | speedup | weight-DRAM | -|---|---|---|---|---| -| N=5120 M=1 | 90.0 | 32.0 (k=8) | **2.81×** | 3.0% → 8.5% | -| N=5120 M=2 | 90.2 | 32.1 (k=8) | **2.81×** | 3.0% → 8.5% | -| N=5120 M=8 | 105.5 | 42.9 (k=8) | **2.46×** | 2.6% → 6.4% | -| N=13824 M=1 | 210.3 | 88.2 (k=2) | **2.38×** | 3.5% → 8.4% | -| N=13824 M=2 | 127.3 | 81.3 (k=8) | **1.57×** | 5.8% → 9.1% | -| M≥32 / large-N M≥8 | — | (heuristic returns 1) | stays direct | no regression | - -split-K wins big exactly in the occupancy-bound small-M/small-N regime -(qkv/o_proj-shaped nodes during speculative verify at M=2–8), and correctly -declines where the base grid already fills the SMs. - -**Gates — all intact:** -- f64-oracle parity + determinism: `marlin_splitk_parity_vs_f64_oracle`, - `marlin_splitk_is_deterministic` (tol 2e-2, incl. asym-zp) pass. -- e2e greedy tokens **BYTE-IDENTICAL** vs tiled with `SPLITK=1` on both - glm-4-9b-int4 and qwen2.5-14b-int4 (zp) M>1 prefill. -- SM80 guard + tiled fallback unchanged; split-K is itself SM80-gated and opt-in. -- Static-grid gemm+reduce pair, pooled partials ⇒ capture-safe warm replays. - -**M=1 DRAM precondition — HONEST re-run: still fails, keep M=1 on the GEMV.** -Even the best split-K lifts the M=1 Marlin GEMM weight-DRAM only to **~8.5% peak** -(N=5120) / ~8.4% (N=13824) — far below the ≥40% floor (feasibility §3). The -full-tile MMA wastes 15/16 rows at M=1; split-K helps time but the weight-load -efficiency stays low. **Decision: M=1 remains on the existing dedicated GEMV.** - -**Remaining work:** gate_up SwiGLU still uses the direct kernel at M>1 (its fused -SiluMul epilogue would need a split-K reduce variant); left as follow-up since the -plain path covers the bulk of split-K-eligible small-M nodes. - ---- - -## Update 5 — the last two M>1 capture barriers: GQA + SkipSimplifiedLayerNorm - -Sebastian's decisive A/B confirmed Marlin is a clean GEMM-level win (capture -B* 8.76×→4.99×, all 240 MatMulNBits capture seams gone, prefill ~2×, -byte-identical). The sole remaining blocker to the speculative-capture GO gate -was the last two ops still declaring `KernelCaptureUnsupported` at M>1: -**GroupQueryAttention (×40)** and **SkipSimplifiedLayerNormalization (×80)**. -Both were conservative gates, not real capture hazards. Fixed (commit -`18d00f90`), both now advertise capture support for a warmed M=K signature. - -**SkipSimplifiedLayerNormalization** — *signature relaxation.* The launch is -already a static grid `(num_groups,1,1)` with no mid-kernel sync / host -read-back / alloc-free on the hot path, and the shape-keyed -`SkipBroadcastMetadataCache` already rejects any shape change (including -`num_groups`) mid-capture (the same pre-warm cold-miss valve as the Marlin -repack cache). The `num_groups == 1` capture latch was purely conservative → -relaxed to `true`. bf16 staging path still demotes on arena growth. **What made -it unsafe at M>1:** nothing — a deliberately conservative latch. **Fix:** latch -capture-safe for any pre-warmed shape. - -**GroupQueryAttention** — *pre-warmed fixed-M flash workspace.* The capture -signature hard-required `q_seq == 1` (Phase2a split-K decode). At q_seq>1 the op -fell to the eager path that **host-reads-back** `seqlens_k` / -`total_sequence_length` (and rotary `position_ids`) — the actual capture killer. -**Fix:** admit a fp16/bf16 `q_seq>1` **fused-flash** signature. `flash_attention::run` -is already capture-safe by construction (static grid `batch·heads·q_tiles`, -device `total_lengths`/`past_lengths`, no host read-back, skips `synchronize` -during capture). The on-device `gqa_prepare_metadata` kernel + `capture_error_ptr` -latch replace host validation; the capture-safe branch uses `present_capacity` -(masking the KV tail on-device) exactly like the M=1 decode path. Backend is -resolved against the fixed `present_capacity` so warm and captured runs agree -deterministically. f32 M>1 (reference-scores, host-sized) stays on the eager -fallback (documented, not a hot decode/verify shape). - -**Gates:** greedy tokens byte-identical on glm-4-9b-int4 (heavy -SkipSimplifiedLayerNorm) M>1 prefill; GQA + normalization lib unit tests pass; -fmt + lib clippy + native-backend clippy clean. Capture support still declared -only for a **warmed** signature — the probe must pre-warm the M=K shape before -the captured attempt (same safety valve as the Marlin repack/scratch pools). - -**→ Hands off to Sebastian** for the decisive Increment-0 re-probe with -`ONNX_GENAI_MARLIN_M_GT_1=1` at M=8: expect `segments → ~1` and `B* ≤ ~2` -(speculative-capture GO). Validate byte-identical captured M=8 tokens + warm -replays. One assumption to confirm on the probe: the KV cache tail beyond the -valid length must not contain NaN (zero-init cache holds; the M=1 capture path -already relies on the same invariant). - -**Per-kernel M>1 capture census (post-fix, by construction):** -- MatMulNBits ×240 (+ gate_up/rmsnorm): Marlin, capture-safe. ✅ (update 3) -- SkipSimplifiedLayerNormalization ×80: static-grid, pre-warm valve. ✅ -- GroupQueryAttention ×40: fused-flash, on-device metadata, pre-warm valve. ✅ -- Residual eager seams at M>1 for the hot path: NONE expected — Sebastian's - probe is the authoritative segment census. - ---- - -## Update 6 — Sebastian's segment census confirmed; split-K default + last dense seam closed - -**Sebastian's Increment-0 re-probe of c842b759 (glm-4-9b, M=8, MARLIN_M_GT_1=1) -CONFIRMED the capture-safety landing:** `segments 120 → 1`; all 40 GQA + 80 -SkipSimplifiedLayerNorm + 240 MatMulNBits M>1 capture seams GONE. Byte-identical -tokens, no NaN under capture (the KV-tail zero-init / on-device tail-mask -assumption holds). ✅ - -**Key measured finding — B\* is COMPUTE-bound, not segmentation-bound.** -Collapsing 120→1 segments moved the captured M=8 wall by ~0 (50.7→51.8ms, B* -4.99→5.10×). Segmentation overhead was already negligible. **Split-K is the -lever that moves B\*:** exploratory `ONNX_GENAI_MARLIN_SPLITK=1` (also -capture-safe, segments=1, byte-identical tokens) drops **B\* 5.10× → 2.69×** -(M=8 wall 51.8→27.3ms, cliff 40.5→16.5ms) — at the ≤2 GO line. - -**Actions this update (commit 29714037):** -1. **Split-K default-ON within the opt-in Marlin M>1 path.** `marlin_splitk_enabled()` - now defaults true (opt out with `ONNX_GENAI_MARLIN_SPLITK=0`). It lives - *inside* `ONNX_GENAI_MARLIN_M_GT_1=1` (default OFF), so no default / consumer - / edge tier is affected; `choose_split_k` still elects a split only for - small-M / low-wave shapes (large-M prefill stays on the byte-identical direct - kernel). This makes the canonical GO probe reach B*≈2.69 without a second flag. -2. **Closed the last M>1 capture seam — the dense `lm_head` logits projection.** - Sebastian's census left exactly one residual node: `lm_head/MatMul_node_734` - → `logits`, a plain fp16 dense `MatMul` (NOT MatMulNBits), which fell to - cuBLASLt's per-call heuristic path at M>1 (M=K verify). It does not fragment - the body (whole stack captures as 1 segment) but was the last node declaring - `KernelCaptureUnsupported`. **Fix:** a cached-plan fast path - (`blas::CaptureGemmPlan` + `DenseGemmPlan` in `kernels/matmul.rs`) for the - plain 2-D (`batch==1`) M>1 GEMM — select algorithm + persistent workspace - once at warmup, replay with no heuristic query / allocation / synchronization - (its own workspace is never shared, so no post-GEMM sync). M>1 analogue of - `F32GemvPlan`; reproduces `governed_gemm`'s arithmetic bit-for-bit at a fixed - shape (same heuristic-selected algo). Batched/broadcast GEMMs keep the - per-call path. Advertises capture after a warmed call. - -**Gates:** greedy tokens byte-identical on glm-4-9b-int4 AND qwen2.5-14b-int4-zp -(`marlin_m_gt_1_matches_tiled_on_*`, GPU 5); fmt + lib clippy + native-backend -clippy clean. SM80 guard + fallback unchanged. - -**→ Sebastian:** re-probe 29714037 with `ONNX_GENAI_MARLIN_M_GT_1=1` (split-K now -default; no SPLITK flag needed). Expect segments=1 (now incl. the lm_head node) -and B*≈2.69 out of the box. The lm_head being inside the captured graph removes -its per-step eager launch + sync; measure whether that nudges B* toward ≤2. The -final small increment to break-even is small-M GEMM throughput (split-K tuning), -not segmentation. - ---- - -## Update 7 — small-M split-K retune for the M=8 verify shape (crack toward B*≤2) - -Coordinator asked for "one more small-M GEMM increment" tuned to the M=K verify -shape to crack B*≤2. Profiled the glm/qwen decode-GEMM dims at M=8 on H200 -(`marlin_bandwidth_microbench`, extended with glm dims + sk=16 candidate) and -found **`choose_split_k` was systematically UNDER-splitting at M=8.** - -**Root cause:** the old rule ("aim ~2 waves of blocks, else don't split") uses -block count as the fill proxy. At M=8 each block is latency-bound (weight DRAM -2-6%), so the optimum is deeper K-splitting (more concurrent memory requests), -not just ~2 waves. It picked sk=2-3 and, worse, sk=1 (no split) for large-N -gate_up — leaving the biggest MLP GEMM on the slow path. - -**Fix (commit 4abe4e57):** split for `m <= 32`; oversubscribe ~8 waves; floor -toward the measured optimum (8 for m<=16, 4 above) even when the base grid -already spans many waves; cap at 8 (16 regresses medium-N). Measured M=8 auto -factor before→after (weight-DRAM %): - -| GEMM (glm/qwen) | K,N | old auto | new auto | gain | -|---|---|---|---|---| -| o/q proj | 4096,4096 | sk3 2.34× | **sk8 2.93×** (6.0%) | +25% | -| down_proj | 13696,4096 | sk3 2.45× | **sk8 3.18×** (7.0%) | +30% | -| gate_up (glm) | 4096,27392 | sk1 none | **sk8 1.11×** (7.2%) | new | -| gate_up (qwen) | 5120,13824 | sk1 none | **sk8 1.35×** (6.8%) | new | -| kv proj | 4096,256 | sk8 5.30× | sk8 5.30× | — | - -The dominant verify cost (MLP gate_up + down_proj) now splits. Prefill (m>32) -still returns sk=1 → byte-identical direct kernel. **Caveat (honest):** qwen -square attn proj (K=5120,N=5120) is bimodal — cold-clock 0.9× / clocked-up 2.5× -(a cold-clock microbench artifact; sustained decode clocks up). It's a minor -fraction of verify cost vs the stably-winning MLP. - -**Gates:** split-K f64-oracle parity + determinism pass at sk=8; greedy tokens -byte-identical on glm-4-9b-int4; fmt + lib clippy clean. `choose_split_k` only -affects m≤32, so prefill/decode are unchanged. - -**→ Sebastian:** this is the "one more small-M increment." Re-probe 4abe4e57 with -`ONNX_GENAI_MARLIN_M_GT_1=1` (split-K default; sk now auto=8 for the M=8 verify -GEMMs incl. gate_up/down). Expect the M=8 verify wall below the 27.3ms that gave -B*=2.71 → B* should drop toward/through ≤2. Report the number; if it plateaus a -bit above 2, that's an honest landing (still GO for strong drafters). - ---- - -## Update 8 — gate_up (SwiGLU MLP) routed through split-K (head 3735d57e) - -**Root-cause of the residual B\* gap.** Sebastian's B\*=2.63 (on 29714037) still -ran the two LARGEST MLP GEMMs on the slow path: `try_launch_marlin_gate_up_prefill` -launched gate→gate_buf and up→output as **direct `launch_marlin_gemm`** calls, -*bypassing split-K entirely*. So the `choose_split_k` retune (update 7) only -reached o/q/kv/down — gate_up (glm K=4096 N=13696 ×2 ×40 layers; the single -biggest MLP projection) never got sk=8. That is the untapped lever. - -**Fix.** Route both gate and up through `maybe_launch_marlin_splitk` so -`choose_split_k` elects the 8-way split at the M=8 verify width. gate/up reuse -the slot-4 partials scratch **sequentially** — gate's reduce fully writes -`gate_buf` before up's GEMM overwrites the partials, so single-stream ordering -keeps it correct. Both split-K `warm` flags are AND-ed into the returned warmth, -so capture-safety now accounts for split-K scratch (pre-warmed at M=K by Part-D). - -**Gates (H200 GPU6):** -- `marlin_gate_up_swiglu_matches_tiled_{plain,rmsnorm,decomposed}` PASS — these - run at M=8 so they now exercise split-K gate_up: worst_abs 0.00012 / 0.00049 / - 0.00781 (≪ tol 0.96 / 4.73 / 4.73), **warm replay byte-identical**, cold-miss - rejected (safety valve intact). -- glm-4-9b-int4 **and** qwen2.5-14b-int4-zp e2e greedy tokens **byte-identical**. -- fmt + lib clippy + native-backend clippy clean. - -**→ Sebastian:** your B\*=2.63 predates BOTH update-7 (`4abe4e57`, o/q/kv/down -retune) and this update-8 (`3735d57e`, gate_up split-K) — they crossed in flight. -Please re-probe **head 3735d57e** with `ONNX_GENAI_MARLIN_M_GT_1=1` (split-K -default). The M=8 verify wall should drop below 26.8ms since the biggest MLP GEMM -now splits. Report the decisive B\*; if it plateaus just above 2 that's an honest -GO for strong drafters — do not force it. - -**Honest ceiling.** M=8 uses `mma.m16n8k16` (16-row tiles) at half occupancy → -~50% MMA waste is intrinsic to tensor-core int4 at M=8 (GPTQ-Marlin pads the -same). Split-K + occupancy is the only lever short of a non-tensor-core path -(cp.async already measured occupancy-bound/slower). Beyond this increment, the -remaining path to a universal B\*≤2 is drafting-depth amortization (Sebastian's -domain), not a capture blocker. - ---- - -## Update 9 — FINAL: gate_up split-K measured; frozen for merge review (head dffabf0d / code 3735d57e) - -Sebastian re-probed **3735d57e** (glm M=8, GPU7, ×2): **B\* = 2.15–2.19×**, -captured M=8 verify wall **stable 21.9ms**, segments=1 whole-graph, **zero -unsupported nodes**, byte-identical greedy tokens, no NaN, split-K slot-4 covered -by Part-D M=8 pre-warm (capture_alloc=(0,0), no cold-miss). The run-to-run ratio -jitter (2.15–2.19) is entirely the M=1 captured baseline (9.98–10.18ms); the M=8 -wall itself is rock-stable. So the frozen head is now the *measured* commit. - -**FROZEN for merge review.** Perf/capture side is DONE + GO (Sebastian concurs). -Full arc: segments **41→1** whole-graph zero-seams; capture **B\* 8.76 → 4.99 -(Marlin) → 2.71 (split-K) → 2.63 (lm_head dense-plan) → 2.16× (small-M retune + -gate_up split-K)**, byte-identical throughout. B\*≈2.16 is the intrinsic small-M -GEMM floor (M=8 `mma.m16n8k16` ~50% MMA waste); universal ≤2 is a drafting-depth -story, not a GEMM-tuning one. Practical GO: any 8-wide draft accepting >2.16 -tok/verify wins. Next: Chew (numerics) + Gaff (quality) review of PR #960. - ---- - -## Update 10 — CORRECTION + qwen breadth: gate_up fusion is block-32-only (glm is block-128) - -Sebastian probed qwen at 3735d57e and flagged a per-model gap; investigating the -code resolves the attribution: - -**Fused gate_up SwiGLU requires `block_size==32`** (hard gate in -`run_f16_gate_up_swiglu`, matmul_nbits.rs:5782). Model block sizes: -- **glm-4-9b = block-128** → the ≥5-input fused gate_up node is NEVER formed; - glm runs gate/up as *separate* MatMulNBits nodes through the **general** Marlin - split-K dispatch (`try_launch_marlin_gemm`→`maybe_launch_marlin_splitk`). Those - were already covered by the update-7 `choose_split_k` retune. **⇒ 3735d57e - (fused gate_up split-K) is a NO-OP for glm.** glm's 2.63→2.16× was ENTIRELY the - 4abe4e57 retune; my update-8 claim that gate_up split-K drove the glm number was - WRONG — corrected here. -- **qwen2.5-14b = block-32** → the fused gate_up node forms; `try_launch_marlin_ - gate_up_prefill` fires at M>1 (eligibility gates block≠0 / k%16 / k%block all - pass for K=5120), and 3735d57e routes its gate/up through split-K. **Confirmed - firing:** qwen's Increment-0 census is segments=1 / zero-unsupported — a - None/Err fall-through would go to the NON-capture-safe tiled prefill and - fragment the graph, so segments=1 *proves* `gate_up_swiglu_marlin_prefill` is - the path taken. The 35.5ms M=8 wall is WITH gate_up Marlin split-K active. - -**qwen M=8 verify @ 3735d57e (Sebastian, GPU7 ×2):** B\* = 4.62/4.72× (captured -M=8 35.5ms, M=1 7.5ms); segments=1, whole-graph, zero unsupported (all 48 GQA + -96 LN capture-safe — capture-safety generalizes to qwen ✅). - -**Diagnosis (honest, not a bug):** qwen's higher B\* is a *denominator* effect, -not a coverage gap. Capture accelerates qwen's launch-bound block-32 M=1 GEMV 34% -(11.4→7.5ms) but the heavier 14B/48-layer M=8 verify only 18% (43.1→35.5ms), so -the RATIO inflates even though both walls drop. qwen's tuned block-32 M=1 is -genuinely fast (7.5ms < glm's 10.2ms), which raises B\* (more accepted tokens -needed to break even). caveat #1 (qwen square-attn K=5120 cold-clock bimodal) is -NOT biting — the probe warms 1000 decode steps so clocks are up; the bimodal was -a cold-clock microbench artifact. - -**Net:** glm-4-9b (canonical gate, block-128) is a clean practical GO at 2.16×. -qwen (block-32) at 4.7× is an honest second-model characteristic — capture-safety -is fully solved (whole-graph, zero seams, byte-identical eager parity), the M=8 -wall is compute-bound at the intrinsic small-M `mma.m16n8k16` floor, and closing -qwen's B\* is a drafting-depth story (deeper/stronger drafter amortizes the higher -ratio), NOT a frozen-kernel fix. No further kernel change; freeze stands. diff --git a/.squad/decisions/inbox/pris-marlin-numerics.md b/.squad/decisions/inbox/pris-marlin-numerics.md deleted file mode 100644 index 0841af6165..0000000000 --- a/.squad/decisions/inbox/pris-marlin-numerics.md +++ /dev/null @@ -1,63 +0,0 @@ -### 2026-08-14: Marlin int4 numerics gate — f64 oracle + justified tolerance envelope + current-path baseline (Pris, #957) - -**By:** Pris (Tester), branch `squad/marlin-numerics` - -**What:** Landed a reusable f64 dequant→GEMM **numerics gate** for `com.microsoft::MatMulNBits` -int4 that any GEMM implementation (today's `gemm_f16_tiled` prefill / decode GEMV path **and** -Deckard's forthcoming Marlin int4 GEMM) must pass before shipping. New self-contained integration -test file `crates/onnx-runtime-ep-cuda/tests/matmul_nbits_marlin_numerics.rs` (no edits to -`matmul_nbits.rs`, so it does not conflict with `squad/marlin-kernel`). - -- **Oracle (ground truth):** dequantize packed int4 to f64 as `(code - zero_point) * scale` - (scale pre-rounded to its fp16/f32 storage value) and accumulate `sum_k a·w` in f64. Both - candidate and oracle consume the **same fp16-rounded activations** and the **same rounded - scale**, so the residual isolates only accumulation precision + fp16 output rounding — never - shared input quantization (mirrors the in-crate GEMV `run_parity_dims_block` convention). -- **Coverage:** group sizes {16,32,64,128}; symmetric AND asymmetric-zp; M∈{1,2,4,8,16,32}; - realistic glm-4-9b (K/N = 4096/4096, 4096/256, 4096/13696, 13696/4096) + Qwen2.5-1.5B - (1536/8960, 8960/1536) attention+MLP projections; fp16 & fp32 scales. -- **Marlin-ready interface:** the gate drives the op purely through **op semantics** (ONNX - `MatMulNBits` node + `ExecutionProvider::get_kernel`/`execute`) — never the internal weight - layout. Once Marlin is wired into the same dispatch, `run_matmul_nbits_f16` validates it with - ZERO changes; alternatively Deckard/Chew feed any candidate output slice into - `Int4Problem::parity`. - -**Justified tolerance envelope** (single source of truth = `Envelope::for_output(max_out)`): -- **Absolute (primary, physical):** `abs_bound = max(max_out * 4e-3, 4e-3)`. fp16 output ULP is - `2^-11 ≈ 4.9e-4` of magnitude; weights dequant through fp16 (another ULP); the K-reduction is - fp32 and a Marlin partial-sum **relayout** re-associates it, adding `~K·eps_f32 < 2e-3` rel drift - at the deepest K≈13696. `4e-3` ≈ 8 fp16 ULP of headroom covers both the tiled baseline and a - re-associated Marlin reduction. -- **Relative (secondary, conditioning-aware):** `rel_bound = 5e-2` against a denominator floored - at `max(1e-1, 3e-2 * max_out)`. Outputs far below the operator's peak are cancellation-dominated - (`|Σ a·w| ≪ Σ|a·w|`); their fp16 round-off is inherently large in *relative* terms while the - *absolute* error is one fp16 ULP of the peak — so they are governed by the abs bound. The `3%` - floor keeps ~3× margin on the worst measured cancellation column. - -**CURRENT-PATH BASELINE (measured on H200, GPU7, feature `gpu-tests`):** -- Group-size sweep (K=4096, N=896, all {group×M×zp×scale-dtype}): **max_abs = 6.50e-2, - max_rel = 2.71e-3, max_out = 1.67e2** → abs error ≈ **3.9e-4 · max_out (≈1 fp16 ULP)**. -- Projection shapes (glm-4 + Qwen2.5, bs=32, M∈{1,8}, sym+asym): **max_abs = 2.53e-1, - max_rel = 8.14e-3, max_out = 5.27e2** → abs error ≈ **4.8e-4 · max_out (≈1 fp16 ULP)**. -- ⇒ The tiled/GEMV path sits ~8–10× inside the asserted abs bound and ~6× inside the rel bound. - **Marlin must land in the SAME envelope** (abs ≲ `max_out·4e-3`, rel ≲ `5e-2` with the 3% - conditioning floor). If Marlin's measured error is comparable to the ~1-ULP baseline it is a - clean pass; anything approaching the bound warrants a conditioning review before sign-off. - -**Why:** Marlin's weight relayout reorders partial sums, so its output is NOT byte-exact vs the -tiled kernel and cannot be diffed bit-for-bit — the only defensible reference is a high-precision -f64 oracle with a justified tolerance. This gate gives Chew (numerics reviewer) an apples-to-apples, -op-level sign-off criterion and gives Deckard a red/green target while iterating. - -**How Deckard/Chew invoke it:** -`source .cudaenv.sh && CUDA_VISIBLE_DEVICES= cargo test -p onnx-runtime-ep-cuda \ ---features gpu-tests --test matmul_nbits_marlin_numerics -- --nocapture --test-threads=1`. -Pure-CPU oracle self-consistency tests (5) run with no GPU and no feature. GPU baseline tests are -`#[ignore]` unless `gpu-tests` is set (CPU-only CI stays green). - -**NOTE (harness lesson, not a kernel bug):** an early un-synchronized run showed a transient -catastrophic divergence (abs 81.8) at one config deep in the sweep; root cause was a missing -post-`execute` `runtime.synchronize()` + reading a non-pre-zeroed output buffer (stale device-pool -memory). Fixed by pre-zeroing the output and synchronizing after execute. With that, the current -int4 path passes the full matrix — no kernel defect. Deckard's Marlin driver should follow the same -sync/zero discipline. diff --git a/.squad/decisions/inbox/sebastian-marlin-bench.md b/.squad/decisions/inbox/sebastian-marlin-bench.md deleted file mode 100644 index 60616eb527..0000000000 --- a/.squad/decisions/inbox/sebastian-marlin-bench.md +++ /dev/null @@ -1,595 +0,0 @@ -### 2026-08-14 — qwen numerics: parity "failure" is a FLAKY TILED ORACLE at a near-tie, NOT a Marlin regression - -**By:** Sebastian (Perf/CUDA). Ran the qwen byte-identical parity test -(`marlin_m_gt_1_matches_tiled_on_qwen2_5_14b_int4`) at `3735d57e` to close qwen breadth. **First run FAILED** -— Marlin M>1 greedy stream diverged from the tiled reference at token index 19 (identical first 19 tokens, -then tiled `320,94409,701,15576,13` vs marlin `5692,374,264,1140,315`). Investigated rigorously before -raising a regression: - -| Run | split-K | tiled reference @ tok19+ | marlin @ tok19+ | result | -|---|---|---|---|---| -| 1 | ON (default) | `320, 94409, …` | `5692, 374, 264, 1140, 315` | FAIL | -| 2 | OFF | `5692, 374, …` | `5692, 374, 264, 1140, 315` | pass | -| 3 | ON | `5692, 374, …` | `5692, 374, 264, 1140, 315` | pass | -| 4 | ON | `5692, 374, …` | `5692, 374, 264, 1140, 315` | pass | - -- **Marlin output is DETERMINISTIC** — `5692,374,264,1140,315` in all 4 runs, split-K on *and* off. **The - tiled reference is the flaky one** (divergent on 1 of 4 runs). Token index 19 is a **near-degenerate - argmax**; the tiled GEMM's own run-to-run nondeterminism (atomic/reduction-order accumulation) occasionally - flips it. Marlin (fixed split-K reduction) is actually *more* stable here. -- **Conclusion: NOT a Marlin numerics regression.** My initial alarm was wrong; rigorous A/B + repeat showed - Marlin is the deterministic path and the tiled oracle flakes ~1-in-4 at this near-tie. Deckard's - "qwen byte-identical" gate passed for him because it landed on the majority (matching) coin-flip. -- **Action item for Chew/Gaff (merge review of #960):** the qwen exact-token-stream parity test is - **flaky (~25%)** because it asserts equality against a nondeterministic tiled reference at a near-tie. Not a - blocker, but the assertion should be hardened — e.g. compare logits within tolerance, use a - deterministic tiled reduction for the oracle, or assert Marlin's own run-to-run determinism instead of - exact-match to a flaky tiled stream. glm's test prompt doesn't hit such a near-tie, so glm passes cleanly. - -#### Reproduce - -```bash -QWEN=/home/justinchu/shared-models/qwen2.5-14b-instruct-int4-zp-onnx # QWEN2_5_14B_CUDA_E2E_DIR default -# Run several times: marlin stays fixed, tiled reference occasionally flips tok19 (near-tie) -CUDA_VISIBLE_DEVICES=7 cargo test -p onnx-genai-engine --features cuda,native-backend --release \ - --test marlin_m_gt_1_e2e marlin_m_gt_1_matches_tiled_on_qwen2_5_14b_int4 -- --ignored --nocapture -# split-K off (also byte-identical to marlin): ONNX_GENAI_MARLIN_SPLITK=0 prefix -``` - ---- - -### 2026-08-14 — Cross-model capture re-probe @ `3735d57e` — glm block_128 (gate_up N/A) vs qwen block_32 — **qwen B\* ≈ 4.7× (NOT GO)** - -**By:** Sebastian (Perf/CUDA). Follow-up to the glm re-probe below. Deckard's `3735d57e` routes the **fused -gate_up SwiGLU** MLP node through split-K — but that fused node is **only produced for `block_size==32` -models** (hard requirement in `run_f16_gate_up_swiglu`, `matmul_nbits.rs:5782`). Verified the two models' -ONNX (`onnx` 1.22): **glm-4-9b `block_size=128`, qwen2.5-14b `block_size=32`.** - -- **glm is NOT eligible for the gate_up fusion (block_128)** → `3735d57e` is a **no-op for glm**. glm's 240 - MatMulNBits run as separate gate/up ops through the *general* Marlin M>1 + split-K path (the `4abe4e57` - retune). This is exactly why the glm captured wall stayed flat (22.0→21.9 ms) after `3735d57e`. **glm - B\*≈2.16× is final and correct; it comes from the general split-K retune, not gate_up fusion.** -- **qwen IS eligible (block_32).** qwen M=8 verify capture re-probe @ `3735d57e` (reproduced ×2, GPU 7): - - | Model | segments | seam census | captured M=8 | captured M=1 | **B\*** | eager M8/M1 | - |---|---|---|---|---|---|---| - | **glm-4-9b** (block_128) | 1 | `` | 21.9 ms | 10.2 ms | **2.16×** | 2.13× | - | **qwen2.5-14b** (block_32) | 1 | `` | 35.5 ms | 7.5 ms | **4.62–4.72×** | 3.73–3.80× | - -- **Capture-safety generalizes to qwen** — segments=1, whole-graph, **zero unsupported nodes** (all 48 GQA + - 96 SkipSimplifiedLayerNorm captured, same as glm). The capture arc is model-independent. ✅ -- **But qwen is NOT at GO (B\*≈4.7×).** Two honest contributors: (a) **denominator effect** — qwen's *tuned* - block-32 M=1 decode is fast (captured 7.5 ms vs glm 10.2 ms; capture speeds M=1 by 34% vs M=8 by 18% - because M=1 is more launch-bound), which *inflates* the ratio; (b) qwen's 14B/48-layer M=8 verify is - genuinely heavier (35.5 ms). The `captured_ratio > eager_ratio` (4.7 > 3.8) is the denominator effect, not - a capture inefficiency. Deckard's caveat #1 (qwen square attn-proj K=5120 cold-clock bimodal) may add to - the M=8 weight. -- **Takeaway:** the small-M GEMM floor is **model-dependent**. glm reaches a practical GO (2.16×); qwen at - 4.7× needs either more small-M GEMM work (verify whether the gate_up fusion is actually firing on qwen and - whether the attn-proj bimodal is real under sustained clocks) **or** deeper drafting to amortize. glm - remains the canonical gate and is GO; qwen is a genuine second-model gap flagged for Deckard. - -#### Reproduce (qwen, `3735d57e`) - -```bash -QWEN=/home/justinchu/shared-models/qwen2.5-14b-instruct-int4-zp-onnx -CUDA_VISIBLE_DEVICES=7 ONNX_GENAI_MARLIN_M_GT_1=1 ONNX_GENAI_RUN_CUDA_SMOKE=1 ONNX_GENAI_LEVERB_MODEL=$QWEN \ - cargo test -p onnx-genai-engine --features cuda,native-backend --release leverb_phase0_capture_probe -- --ignored --nocapture -# block_size probe: python3 -c "import onnx;m=onnx.load(P,load_external_data=False);..." (see log) -``` - ---- - -### 2026-08-14 — Marlin capture RE-PROBE @ `3735d57e` (gate_up SwiGLU split-K routing) — **B\* ≈ 2.16× (GEMM PLATEAU)** - -**By:** Sebastian (Perf/CUDA). Merged Deckard's `3735d57e` / `dffabf0d`: `try_launch_marlin_gate_up_prefill` -(gate+up SwiGLU MLP — glm K=4096 N=13696 ×40, the single biggest MLP GEMM) was still DIRECT-launching and -bypassing split-K even after the `4abe4e57` factor retune; now both route through split-K (sk=8). H200 -**GPU 7** (Deckard off 6), all idle. Canonical path = `ONNX_GENAI_MARLIN_M_GT_1=1` (split-K default). - -#### Result (leverb-phase0 Part D INC0, glm-4-9b, M=8) — reproduced ×2 - -| Config (head) | segments | seam census | captured M=8 wall | **B\*** | -|---|---|---|---|---| -| M=8 split-K factor retune (`4abe4e57`) | 1 | `` | 22.0 ms | 2.16× | -| **+ gate_up split-K routing (`3735d57e`)** | **1** | `` ✅ | **21.9 ms** | **2.15× / 2.19×** | - -- **B\* ≈ 2.16× — the GEMM plateau.** The captured M=8 verify wall is now stable at **~21.9 ms** (run-to-run - ratio jitter 2.15–2.19× comes from the M=1 captured baseline 9.98–10.18 ms, not the M=8 wall). Routing the - dominant gate_up MLP GEMM through split-K shaved the last easy small-M cost; segments=1, whole-graph, zero - unsupported nodes; byte-identical greedy tokens (parity PASS, no NaN). -- **This is the small-M GEMM floor.** Deckard's honest read (concur): M=8 `mma.m16n8k16` is intrinsically - ~half-occupancy (~50% MMA lane waste), so split-K + occupancy is the **last GEMM lever**. Beyond ~2.1×, - the path to a *universal* ≤2 is **drafting depth** (more accepted tokens amortize the verify), not more - GEMM tuning. -- **Speculative-capture verdict = GO in practice.** B\*≈2.16× vs the conservative break-even-at-2 line: any - 8-wide draft accepting **>2.16 tokens/verify on average** wins. Capture-safety is fully solved. - -#### Full progression (the number that gates Marlin) - -``` -segments: 41 → 120 → 1 → 1 → 1 → 1 (BEFORE→…→whole-graph, zero seams) -B*: 8.76 → 4.99 → 5.09 → 2.71 → 2.63 → 2.16× -heads: OFF c692 c842 c842+sk 29714 4abe/3735 (byte-identical throughout) -``` - -#### Reproduce (`3735d57e`) - -```bash -source /home/justinchu/onnx-genai/.cudaenv.sh && git fetch origin && git merge 3735d57e -GLM=/home/justinchu/glm-e2e-artifacts/glm-4-9b-int4-cuda -CUDA_VISIBLE_DEVICES=7 ONNX_GENAI_MARLIN_M_GT_1=1 ONNX_GENAI_RUN_CUDA_SMOKE=1 ONNX_GENAI_LEVERB_MODEL=$GLM \ - cargo test -p onnx-genai-engine --features cuda,native-backend --release leverb_phase0_capture_probe -- --ignored --nocapture -CUDA_VISIBLE_DEVICES=7 GLM_4_9B_CUDA_E2E_DIR=$GLM \ - cargo test -p onnx-genai-engine --features cuda,native-backend --release --test marlin_m_gt_1_e2e \ - marlin_m_gt_1_matches_tiled_on_glm_4_9b_int4 -- --ignored --nocapture -``` - ---- - -### 2026-08-14 — Marlin capture RE-PROBE @ `4abe4e57` (M=8 verify split-K retune) — **B\* = 2.16×** (at the GO line) - -**By:** Sebastian (Perf/CUDA). Merged Deckard's `4abe4e57` / `0143f235`: retuned `choose_split_k` for the M=8 -verify shape (the old heuristic under-split — gate_up MLP GEMM was on the no-split slow path). New M=8 auto -factors split gate_up + down + o/q deeper (sk=8). Prefill (M>32) still returns sk=1 → byte-identical direct -kernel (prefill/decode measurements unchanged). H200 **GPU 7** (Deckard on 4/6), all idle. Canonical path = -`ONNX_GENAI_MARLIN_M_GT_1=1` (split-K default, sk auto=8 for the verify GEMMs). - -#### Result (leverb-phase0 Part D INC0, glm-4-9b, M=8) — reproduced ×2 - -| Config (head) | segments | seam census | captured M=8 wall | **B\*** | -|---|---|---|---|---| -| lm_head plan, split-K default (`29714037`) | 1 | `` | 26.8 ms | 2.63× | -| **+ M=8 split-K retune (`4abe4e57`)** | **1** | `` ✅ | **22.0 / 22.1 ms** | **2.16× / 2.17×** | - -- **B\* 2.63 → 2.16×** — Deckard's gate_up/down/o-q M=8 split-K retune dropped the captured verify wall - 26.8→22.0 ms. Reproduced across two runs (2.16×, 2.17×); Part-C eager upper-bound 2.13×, CLIFF 10.8 ms, - TAIL slope 0.935 ms/row. Still segments=1, whole-graph, **zero unsupported nodes**. -- **Byte-identical** greedy tokens on glm-4-9b-int4 with the retuned path (parity PASS, no NaN). -- **This is the plateau: B\*=2.16× is an honest hair above the conservative ≤2 line.** But ≤2 is the - break-even-at-2-accepted-tokens threshold; at B\*=2.16 speculative decode is a **practical GO** the moment - the draft accepts **>2.16 tokens/verify on average** (trivially met by any decent 8-wide draft). The full - arc: **segments 41→1 (whole-graph, zero seams); B\* 8.76→4.99→5.09→2.71→2.63→2.16×**, byte-identical - throughout. Capture-safety is DONE; the residual 0.16× is pure small-M GEMM compute, not a capture blocker. - -#### Reproduce (`4abe4e57`) - -```bash -source /home/justinchu/onnx-genai/.cudaenv.sh && git fetch origin && git merge 4abe4e57 -GLM=/home/justinchu/glm-e2e-artifacts/glm-4-9b-int4-cuda -CUDA_VISIBLE_DEVICES=7 ONNX_GENAI_MARLIN_M_GT_1=1 ONNX_GENAI_RUN_CUDA_SMOKE=1 ONNX_GENAI_LEVERB_MODEL=$GLM \ - cargo test -p onnx-genai-engine --features cuda,native-backend --release leverb_phase0_capture_probe -- --ignored --nocapture -CUDA_VISIBLE_DEVICES=7 GLM_4_9B_CUDA_E2E_DIR=$GLM \ - cargo test -p onnx-genai-engine --features cuda,native-backend --release --test marlin_m_gt_1_e2e \ - marlin_m_gt_1_matches_tiled_on_glm_4_9b_int4 -- --ignored --nocapture -``` - ---- - -### 2026-08-14 — Marlin capture RE-PROBE @ `29714037` (lm_head dense-plan capture + split-K default) — **ZERO unsupported nodes, whole-graph capture** - -**By:** Sebastian (Perf/CUDA). Merged Deckard's `29714037` / `1a72c865` onto `squad/marlin-bench`: (1) lm_head -`MatMul_node_734 → logits` now uses a cached dense-GEMM plan (`blas::CaptureGemmPlan` + `DenseGemmPlan` in -`kernels/matmul.rs`) — select algo + persistent workspace once at warmup, replay with no heuristic/alloc/sync, -so it **folds into the captured graph**; (2) **split-K is now default-on inside `ONNX_GENAI_MARLIN_M_GT_1=1`** -(no second flag; opt out with `ONNX_GENAI_MARLIN_SPLITK=0`). H200 **GPU 7** (Deckard on GPU 5, no contention), -all 8 verified idle. Canonical probe path = just `ONNX_GENAI_MARLIN_M_GT_1=1`. - -#### Result (leverb-phase0 Part D INC0, glm-4-9b, M=8) - -| Config (head) | **segments** | residual seam nodes | captured M=8 wall | **B\*** (M8/M1) | -|---|---|---|---|---| -| Marlin M>1 + GQA/LN safe (`c842b759`) | 1 | **MatMul×1** (lm_head) | 51.9 ms | 5.09× (split-K off) / 2.71× (split-K on) | -| **+ lm_head plan, split-K default (`29714037`)** | **1** | **``** ✅ | **26.8 ms** | **2.63×** | - -- **Census is now TRUE ZERO unsupported nodes.** The probe reports `seam nodes: ` — the - entire M=8 verify forward (including lm_head logits projection) captures as **one whole-subgraph replay**. - Deckard's cached dense-plan closed the last `MatMul[KernelCaptureUnsupported]×1`. It needs the same M=8 - pre-warm (dense plan rejects an unwarmed shape during capture) — Part-D INC0's full-M=K warm covers it - (`warm_alloc=(923,722)`, `capture_alloc=(0,0)`). -- **B\* 2.71→2.63×** (captured M=8 wall 27.5→26.8 ms): folding lm_head out of its per-step eager launch+sync - nudged B\* down as Deckard predicted (~one small-M GEMM increment), but the strict **≤2 GO line is not yet - crossed** — the residual is genuine small-M GEMM compute, not dispatch/segmentation. Eager upper-bound - ratio 2.56×; Part-C CLIFF 16.98 ms, TAIL slope **0.769 ms/row**. -- **Numerics:** greedy tokens **byte-identical** to tiled on glm-4-9b-int4 with the new default path - (split-K default + lm_head plan) — `marlin_m_gt_1_matches_tiled_on_glm_4_9b_int4` PASS. No NaN. -- **This is the cleanest capture config to date:** zero unsupported nodes, whole-graph replay, - byte-identical, B\*=2.63×. Speculative-capture verdict remains **structurally GO / numerically break-even - pending** — 2.63 is just above ≤2; deeper drafting or one more small-M GEMM step reaches it. - -#### Reproduce (`29714037`) - -```bash -source /home/justinchu/onnx-genai/.cudaenv.sh && git fetch origin && git merge 29714037 -GLM=/home/justinchu/glm-e2e-artifacts/glm-4-9b-int4-cuda -# Canonical probe — split-K is now DEFAULT inside MARLIN_M_GT_1=1 (no SPLITK flag) -CUDA_VISIBLE_DEVICES=7 ONNX_GENAI_MARLIN_M_GT_1=1 ONNX_GENAI_RUN_CUDA_SMOKE=1 ONNX_GENAI_LEVERB_MODEL=$GLM \ - cargo test -p onnx-genai-engine --features cuda,native-backend --release leverb_phase0_capture_probe -- --ignored --nocapture -# Byte-identical parity (new default path) -CUDA_VISIBLE_DEVICES=7 GLM_4_9B_CUDA_E2E_DIR=$GLM \ - cargo test -p onnx-genai-engine --features cuda,native-backend --release --test marlin_m_gt_1_e2e \ - marlin_m_gt_1_matches_tiled_on_glm_4_9b_int4 -- --ignored --nocapture -``` - ---- - -### 2026-08-14 — Marlin capture RE-PROBE @ `c842b759` (GQA + LayerNorm M>1 capture landed) — **segments → 1** - -**By:** Sebastian (Perf/CUDA). Merged Deckard's `c842b759` (commit `18d00f90`: `GroupQueryAttention`×40 + -`SkipSimplifiedLayerNormalization`×80 now advertise M>1 capture support) onto `squad/marlin-bench`. H200 -GPU 7 (Deckard on GPU 4, no contention). `ONNX_GENAI_MARLIN_M_GT_1=1`, Part-D INC0 pre-warms the M=8 shape -(no cold-miss). Canonical = split-K OFF (per Deckard). Exploratory row adds `ONNX_GENAI_MARLIN_SPLITK=1`. - -#### The capture-safety progression (leverb-phase0 Part D INC0, glm-4-9b, M=8) - -| Config (head) | **segments** | residual seam nodes | captured M=8 wall | **B\*** (M8/M1) | -|---|---|---|---|---| -| BEFORE — Marlin OFF | 41 | GQA×40, **NBits×240**, MatMul×1, LN×80 | 87.3 ms | 8.76× | -| Marlin M>1 (`c692bffe`) | 120 | GQA×40, MatMul×1, LN×80 | 50.7 ms | 4.99× | -| **+ GQA/LN capture-safe (`c842b759`, CANONICAL)** | **1** ✅ | **MatMul×1** | 51.8 ms | **5.10×** | -| + split-K (exploratory, also capture-safe) | **1** | MatMul×1 | **27.3 ms** | **2.69×** | - -- **`segments → 1` ACHIEVED.** Deckard's GQA+LN capture-safety cleared the last M>1 seams; the whole M=8 - verify forward now captures as a **single replayable graph**. Only residual seam: a lone - **`MatMul[KernelCaptureUnsupported]×1`** (one dense MatMul — not NBits; Deckard to identify/relax). This - is the operational win for speculative decode: **one graph replay, not 120 segment dispatches.** -- **But capture-safety alone did NOT move B\*** (5.10× vs the marlin-only 4.99× — within contention noise): - collapsing 120→1 segments changed the captured wall by ~0 (50.7→51.8 ms). **The captured M=8 wall is - compute-bound, not segmentation-bound** — segmentation overhead was already negligible, so the - expectation "segments→1 ⇒ B\*≤2" does not hold on its own. -- **Split-K is the lever that actually moves B\*.** Exploratory `ONNX_GENAI_MARLIN_SPLITK=1` (which is **also - capture-safe → segments=1, and byte-identical**, verified) nearly halves the captured M=8 wall again - (51.8→**27.3 ms**), dropping **B\* 5.10×→2.69×** — right at the ≤~2 GO line (cliff 40.5→16.5 ms; eager - M8/M1 2.51×). **Path to speculative-capture GO = capture-safety (done, segments=1) + split-K on the M=K - verify.** B\*=2.69 is a hair above ≤2; one more small-M GEMM increment (or drafting slightly deeper) - reaches break-even. -- **Numerics correct (Deckard's NaN check):** Marlin-ON greedy tokens are **byte-identical** to tiled-OFF, - both **with and without split-K** (`marlin_m_gt_1_matches_tiled_on_glm_4_9b_int4` PASS, 24/24 tokens); - no NaN/garbage under capture → the zero-init KV-tail + on-device tail-masking assumption holds. -- **Warmup coverage confirmed (Deckard's must-handle #1):** Part-D INC0 (`leverb_increment0_capture_attempt`, - `cuda.rs:1157` FIX 2) warms via a **full M=K forward** (`run_with_device_bindings` on the exact M=8 - bindings later captured) — one pass populates **every** op's fixed-shape cache (Marlin repack + scratch, - GQA metadata, LN metadata). Empirically `capture_alloc=(0,0)` and `segments=1` confirm no cold-miss valve - fired. All new shapes covered. -- **Fresh reconfirm @ HEAD (measurement discipline):** re-ran both probes clean — canonical - segments=1/seam MatMul×1/**B\*=5.09×** (M8=51.9, M1=10.2 ms); split-K segments=1/**B\*=2.71×** (M8=27.5, - M1=10.2 ms). Reproduces the 5.10/2.69 within noise. - -#### Final end-to-end table @ `c842b759`, split-K ON (M-scaling + prefill) - -Marginal-compute **TAIL slope** (ms/row, the pure GEMM signal Marlin targets) across the whole arc: - -| Model | BEFORE (OFF) | Marlin M>1 (`c692bffe`) | **+ split-K (`c842b759`)** | cliff (split-K) | prefill L=128/512/1024 tok/s (split-K) | -|---|---|---|---|---|---| -| **glm-4-9b** | 2.21 | 0.93 | **0.822** (−63%) | 16.6 ms | 382 / 378 / 419 | -| **qwen2.5-14b** | 3.26 | 1.10 | **1.147** (−65%) | 48.0 ms | 267 / 271 / 277 | - -- **Prefill nuance:** for **large-M prefill**, canonical Marlin edges split-K (qwen 291–300 vs split-K - 267–277 tok/s) — split-K's win is concentrated at **small-M** (the M=K verify, where B\* lives). - **Recommendation: canonical Marlin for prefill; split-K auto-on for the M=K speculative verify.** -- Both vs BEFORE prefill (glm 218/213/217, qwen 147/152/154 tok/s) the arc is **~2×** end-to-end. -- qwen's cliff stays large (48 ms) because Axis-2 is the *eager, capture-breaking* wall (per-M penalty for - leaving graph replay); the **tail slope** is the Marlin signal and it collapses 3.26→1.15. - -#### Reproduce (`c842b759`) - -```bash -source /home/justinchu/onnx-genai/.cudaenv.sh && git fetch origin && git merge c842b759 -GLM=/home/justinchu/glm-e2e-artifacts/glm-4-9b-int4-cuda -# CANONICAL capture re-probe (split-K off) -CUDA_VISIBLE_DEVICES=7 ONNX_GENAI_MARLIN_M_GT_1=1 ONNX_GENAI_RUN_CUDA_SMOKE=1 ONNX_GENAI_LEVERB_MODEL=$GLM \ - cargo test -p onnx-genai-engine --features cuda,native-backend --release leverb_phase0_capture_probe -- --ignored --nocapture -# EXPLORATORY (adds split-K) — B* 2.69x -CUDA_VISIBLE_DEVICES=7 ONNX_GENAI_MARLIN_M_GT_1=1 ONNX_GENAI_MARLIN_SPLITK=1 ONNX_GENAI_RUN_CUDA_SMOKE=1 ONNX_GENAI_LEVERB_MODEL=$GLM \ - cargo test -p onnx-genai-engine --features cuda,native-backend --release leverb_phase0_capture_probe -- --ignored --nocapture -# Byte-identical token parity (add ONNX_GENAI_MARLIN_SPLITK=1 to check split-K too) -CUDA_VISIBLE_DEVICES=7 GLM_4_9B_CUDA_E2E_DIR=$GLM \ - cargo test -p onnx-genai-engine --features cuda,native-backend --release --test marlin_m_gt_1_e2e \ - marlin_m_gt_1_matches_tiled_on_glm_4_9b_int4 -- --ignored --nocapture -# Final end-to-end M-scaling + prefill @ split-K ON (qwen: swap --model) -cargo build --release -p onnx-genai-bench --features bench-native,cuda --bin marlin_bench -CUDA_VISIBLE_DEVICES=7 ONNX_GENAI_MARLIN_M_GT_1=1 ONNX_GENAI_MARLIN_SPLITK=1 ./target/release/marlin_bench \ - --model $GLM --label glm-4-9b-int4 --m-list 1,2,4,8,16 --past 32 --iters 20 --warmups 3 \ - --prefill-lens 128,512,1024 --prefill-iters 4 -``` - ---- - -### 2026-08-14 — Marlin int4 GEMM **AFTER** measurements (Lever A landed, #957) - -**By:** Sebastian (Perf/CUDA). Measured on the merged snapshot **`c692bffe`** (Deckard's -`squad/marlin-kernel` P1 = full M>1 Marlin coverage, zero tiled fallbacks, byte-identical tokens) merged -into `squad/marlin-bench`. H200 SXM (SM90), `CUDA_VISIBLE_DEVICES=7` (verified idle; Deckard's P2 work on -GPU 0), `nvidia-smi` re-checked each run. Toggle: **`ONNX_GENAI_MARLIN_M_GT_1=1`** (default OFF = portable -tiled fallback = the BEFORE path). Every OFF row below reconfirms the BEFORE baseline on the *same* binary, -so OFF↔ON is a clean A/B on one snapshot. Contention-invariant **min** reported (a co-tenant spike -inflated one qwen L=1024 *median* to 23 611 ms; its min stayed 6 668 ms — exactly why min is the statistic). - -#### HEADLINE — Increment-0 capture re-probe (leverb-phase0 Part D INC0, glm-4-9b, GPU 7) - -The decisive gate. Part D warms the M=K(=8) shape *before* the captured attempt (FIX 2), so Marlin's -repack cache + scratch pools are populated → no cold-miss safety-valve refusal. Target: `segments`→~1, -captured M=8/M=1 ratio (break-even **B\***) →≤~2 (flips speculative-capture to GO). - -| Metric | BEFORE (Marlin OFF) | AFTER (Marlin ON) | Δ | target | -|---|---|---|---|---| -| INC0 M=8 **captured replay ratio B\*** | **8.76×** (M8 87.3 / M1 10.0 ms) | **4.99×** (M8 50.7 / M1 10.2 ms) | **−43%** | ≤~2 | -| INC0 M=8 **segments** | **41** | **120** (see note) | — | ~1 | -| INC0 M=8 seam nodes | GQA×40, **MatMulNBits×240**, MatMul×1, LN×80 | GQA×40, MatMul×1, LN×80 (**MatMulNBits×240 GONE**) | −240 seams | — | -| eager M=8/M=1 (Part C) | 6.78× | 4.10× | −40% | — | -| eager cliff (M1→M2) | 66.9 ms | 37.5 ms | −44% | — | -| eager **tail slope** (M2..8) | 1.814 ms/row | **0.744 ms/row** | **−59%** | — | -| VERDICT | NO-GO | NO-GO (much closer) | — | GO@B\*≤2 | - -- **Marlin did its job decisively:** it made **all 240 MatMulNBits M>1 nodes capture-safe** (they vanish - from the seam list) and **nearly halved the captured M=8 wall** (87.3→50.7 ms), dropping **B\* 8.76×→4.99×**. -- **Why `segments` rose 41→120 (not a regression — the opposite):** with MatMulNBits now *inside* captured - regions, the graph's capturable runs are no longer swallowed by all-seam stretches; they surface as - distinct segments bounded by the **still-uncapturable `GroupQueryAttention`×40 + `SkipSimplifiedLayerNormalization`×80**. Those seams are **not Marlin's scope** — they are the complementary - **capture-safety lever** (make GQA/LN capture-safe). Marlin removed the GEMM barrier; the residual barrier - to `segments→1` / `B*≤2` is now cleanly isolated to attention+layernorm capture-safety. -- **Bottom line for the program:** Lever A (Marlin) is a **success at the GEMM level** and moved B\* 43% of - the way toward the GO line. Full speculative-capture GO (B\*≤2) additionally needs GQA+LN capture-safety; - that is the next lever, and this probe is already wired to re-measure it (same command). - -#### Prefill sweep — Marlin OFF vs ON (native), and the native-vs-ORT gap - -`marlin_bench --prefill-lens 128,512,1024` (batched M=len forward; contention-invariant min). ORT-genai -column = the re-exported stock chatglm int4 (unchanged, Marlin is native-only). - -| Model / L | native OFF | native ON | Marlin speedup | ORT-genai | gap OFF→ON | -|---|---|---|---|---|---| -| **glm** 128 | 209.8 | **415.4** | **1.98×** | 8 002 | 38× → 19× | -| **glm** 512 | 216.4 | **417.3** | 1.93× | 21 044 | 97× → 50× | -| **glm** 1024 | 218.4 | **425.6** | 1.95× | 26 331 | 121× → **62×** | -| **qwen** 128 | 147.3 | **291.3** | 1.98× | — | — | -| **qwen** 512 | 151.8 | **295.3** | 1.95× | — | — | -| **qwen** 1024 | 153.6 | **300.4** | 1.96× | — | — | - -- **Marlin ~doubles native prefill throughput** on both models (tok/s flat vs length still, but the whole - curve lifts ~2×), **halving the glm-vs-ORT prefill gap** (121×→62× at L=1024). The residual gap is the - same non-GEMM surface: native prefill is eager (no CUDA graph for M>1) and materializes full-sequence - logits (D2H), while og-genai runs a CUDA-graph-fused pipeline — i.e. closing it further is again the - **capture-safety lever**, not more GEMM. Marlin delivered the GEMM half of the prefill win. - -#### M-scaling wall — Marlin OFF vs ON (`marlin_bench`, past=32, min ms) - -| Model | M=1 | M=2 | M=4 | M=8 | M=16 | cliff (M1→M2) | tail (ms/row) | M16/M1 | -|---|---|---|---|---|---|---|---|---| -| glm **OFF** | 13.27 | 80.30 | 82.92 | 90.59 | 111.34 | 67.3 | 2.214 | 8.37× | -| glm **ON** | 13.40 | 51.02 | 51.79 | 54.90 | 63.90 | **37.7** | **0.925** | **4.77×** | -| qwen **OFF** | 11.16 | 105.09 | 111.25 | 121.70 | 150.75 | 94.2 | 3.258 | 13.45× | -| qwen **ON** | 11.31 | 89.38 | 90.15 | 92.85 | 104.61 | **78.1** | **1.103** | **9.24×** | - -- **Confirms Deckard's kernel-level 6.9×→~2× at the forward/op level:** the **tail slope** (marginal - compute per extra verify row = the pure GEMM signal) collapses **glm 2.21→0.93 (−58%)**, **qwen - 3.26→1.10 (−66%)** — the M>1 int4 GEMM is ~2.4–3× cheaper per row, matching the kernel result once the - fixed non-GEMM cliff is separated out. The **cliff** (per-op host dispatch + attention/LN, NOT GEMM) - drops less (glm −44%, qwen −17%) — again pointing at the capture-safety lever for the remainder. - -#### Reproduce (AFTER, head `c692bffe` = merge of `squad/marlin-kernel` into `squad/marlin-bench`) - -```bash -source /home/justinchu/onnx-genai/.cudaenv.sh -git fetch origin && git merge c692bffe # pin Deckard's P1 snapshot -cargo build --release -p onnx-genai-bench --features bench-native,cuda --bin marlin_bench -nvidia-smi --query-compute-apps=pid,used_memory --format=csv,noheader # pin an idle HIGH index (7) -GLM=/home/justinchu/glm-e2e-artifacts/glm-4-9b-int4-cuda - -# HEADLINE — capture re-probe, OFF then ON (Part D INC0 = warmed M=K -> no cold-miss) -CUDA_VISIBLE_DEVICES=7 ONNX_GENAI_RUN_CUDA_SMOKE=1 ONNX_GENAI_LEVERB_MODEL=$GLM \ - cargo test -p onnx-genai-engine --features cuda,native-backend --release leverb_phase0_capture_probe -- --ignored --nocapture -CUDA_VISIBLE_DEVICES=7 ONNX_GENAI_MARLIN_M_GT_1=1 ONNX_GENAI_RUN_CUDA_SMOKE=1 ONNX_GENAI_LEVERB_MODEL=$GLM \ - cargo test -p onnx-genai-engine --features cuda,native-backend --release leverb_phase0_capture_probe -- --ignored --nocapture - -# M-scaling wall + prefill sweep, OFF then ON (drop ONNX_GENAI_MARLIN_M_GT_1 for OFF) -CUDA_VISIBLE_DEVICES=7 [ONNX_GENAI_MARLIN_M_GT_1=1] ./target/release/marlin_bench \ - --model $GLM --label glm-4-9b-int4 --m-list 1,2,4,8,16 --past 32 --iters 20 --warmups 3 \ - --prefill-lens 128,512,1024 --prefill-iters 4 -# qwen: --model /home/justinchu/shared-models/qwen2.5-14b-instruct-int4-zp-onnx -``` - -#### Portability (Rule 11) — **SHARED-BOX MEASUREMENT NOTE.** The 8× H200 host is shared with other squad agents; foreign compute -> apps appeared and vanished mid-run (e.g. a 45 GB job on another GPU). Contention only ever *adds* wall -> time, so for the M-scaling wall the **min across iters is the contention-invariant estimator** and is -> what the table reports; the harness also prints median + p10/p90/max so a contaminated run is visible -> (a p90≫min flags it). The e2e tok/s figures below were captured in windows with tight p10/p90 spread. - ---- - -#### Axis 1 — M=1 int4 GEMV achieved DRAM % (the Marlin precondition) - -`ncu --graph-profiling node -k regex:matmul_nbits_gemv` on the captured native-CUDA decode -(`profile_native --ep cuda --backend native --steady`), metric -`gpu__dram_throughput.avg.pct_of_peak_sustained_elapsed`. Marlin M=1 pass = ≥~55%; <~40% = M=1 Marlin -NO-GO (thresholds apply to Marlin's *achieved* %, measured AFTER it lands — these are the BEFORE numbers). - -| Model | GEMV kernel (shape) | grid | DRAM % | SM % | check (bytes/time) | -|---|---|---|---|---|---| -| **qwen2.5-14b** | `gemv_f16_gate_up_decomposed_swiglu_rmsnorm` (MLP gate_up) | 1728 | **29.4–29.7%** | 49% | 79.7 MB / 57.5 µs = 1386 GB/s = 28.9% ✓ | -| qwen2.5-14b | `gemv_f16_scales_f16_down_c2` (down_proj) | 2560 | 39.7% | 47% | 39.8 MB / 21.6 µs ✓ | -| qwen2.5-14b | `gemv_f16_scales_f16` (qkv/o proj) | 640 | 24.0–24.4% | 41% | 14.8 MB / 12.8 µs ✓ | -| **glm-4-9b** | `gemv_f16_general_bs` (MLP) | 3424 | **17.2%** | 77% | 57.9 MB / 72.7 µs = 797 GB/s = 16.6% ✓ | -| glm-4-9b | `gemv_f16_general_bs` (attn proj) | 512 | 10.8–12.4% | 50–56% | 8.7 MB / 16.7 µs = 519 GB/s ✓ | - -- **The qwen MLP gate_up GEMV = ~29.5% DRAM reproduces the doc §3 precondition ("~29% DRAM on an MLP - shape") to the point.** Its down_proj tops out at ~40%, attn projections ~24%. -- **glm-4-9b runs the *untuned* `gemv_f16_general_bs` path** (no fused-epilogue tuned variant selected for - its shapes) → only 11–17% DRAM. Both models sit far below the ≥55% Marlin pass line and the MLP shapes - straddle the 40% NO-GO line — i.e. the M=1 GEMV is **latency/dispatch-bound, not DRAM-bound**, exactly - the diagnosis Marlin must overturn. Marlin's decisive number is the *delta* it puts on these rows. - -#### Axis 2 — M-scaling wall (the cliff Marlin + Lever B must collapse) - -`marlin_bench` times a single real `decode_verify` (eager, capture-breaking) over M∈{1,2,4,8,16} on each -model's own MatMulNBits stack; **min ms** (contention-invariant), 30 iters, 5 warmups. - -| Model / past_len | M=1 | M=2 | M=4 | M=8 | M=16 | **cliff (M1→M2)** | tail (ms/row) | -|---|---|---|---|---|---|---|---| -| glm-4-9b / past=32 (Increment-0 repro) | 13.4 | 81.1 | 83.9 | **91.0** | 111.7 | **67.7 ms** | ~1.65 | -| glm-4-9b / past=512 (realistic ctx) | 14.0 | 102.7 | 103.3 | 112.3 | 130.7 | 88.7 ms | ~1.6–2.0 | -| qwen2.5-14b / past=512 | 11.9 | 132.9 | 139.2 | 149.5 | 182.1 | 120.9 ms | ~2.8–3.5 | - -- **Reproduces Deckard's Increment-0** at matched past (~32): cliff **67.7 ms** (≈ the ~67 ms M=1→M=2 - cliff), M=8 **91 ms** (≈ 87 ms), tail **~1.65 ms/row** (≈ 1.85). The cliff is the fixed penalty for - leaving the captured M=1 fast path (per-op host dispatch + generic multi-row GEMM + scratch alloc); the - tail is true marginal compute per verify row. -- Reference captured M=1 decode step (non-verify, graph replay): glm **10.1–10.4 ms**, qwen **7.8 ms** — - the eager verify M=1 (13–14 ms) is already ~1.3× the captured M=1, before the cliff. **This whole wall - is what Lever B (capture-stable M=K verify) collapses toward ~1×; Marlin attacks the per-GEMV cost - underneath it.** - -#### Axis 3 — end-to-end tok/s, native CUDA EP vs ORT - -`profile_native --ep cuda --steady --warmups 1 --runs 3 --tokens 128 --decode-skip 8` (tight spread). - -| Model | native CUDA EP | ORT CUDA EP | native / ORT | -|---|---|---|---| -| **qwen2.5-14b-int4** | **127.8 tok/s** | **89.2 tok/s** (ORT 1.27 EP, same ONNX) | **1.43×** | -| **glm-4-9b-int4** | **93.4–98.6 tok/s** | **233.0 tok/s** (ORT-genai, see note) | **0.40× (native SLOWER)** | - -- qwen2.5-14b native **127.8 tok/s** vs ORT-EP **89.2** on the *same* native-authored ONNX → clean - **1.43× native win**. glm-4-9b native **93–99 tok/s** matches the doc's 96.4 tok/s reference. -- **glm-ORT head-to-head — RESOLVED (was blocked).** No stock ORT can load the native-authored glm ONNX: - ORT 1.27 EP, pip ORT 1.28, *and* onnxruntime-genai 0.14.1's bundled ORT all reject its - `GroupQueryAttention` with `Unrecognized attribute: rotary_embedding_dim` (a native-EP-exclusive GQA - schema). No foundry-local glm package exists. The doc-specified path is therefore a **re-export**: I ran - the onnxruntime-genai model builder on the local `glm-4-9b-gptq-source` (`ChatGLMModel`, gptq→int4, - `-p int4 -e cuda`) to produce a clean stock ORT-genai `chatglm` int4 package (`./ort-glm-build`, - `type=chatglm`, 40L/32H/2KV, matches native dims). It loads in og's ORT and runs. -- **FINDING (truthful negative):** on glm-4-9b int4, stock **onnxruntime-genai BEATS the native CUDA EP** - — decode **233 vs ~95 tok/s (≈2.5× faster)** and prefill catastrophically (see Axis 3b, ~100–120×). The - "native beats ORT" claim currently holds for **qwen** (1.43×) but **NOT for glm**. Two contributors: - (a) native glm runs the *untuned* `gemv_f16_general_bs` GEMV (Axis 1: 17% DRAM vs qwen's tuned 29%); - (b) native has **no efficient M>1 int4 GEMM**, so its prefill is ~100× slower — *exactly the Marlin - gate*. og uses CUDA graphs + fused attention + its own int4 GEMM. -- **Caveats (not perfectly apples-to-apples):** the ORT side is builder-*requantized* (gptq→int4 via the - og builder), not the identical weights/quant as the native file; accuracy parity is **not** verified. - The methodology also differs from qwen (qwen = same ONNX under ORT-EP; glm = og-genai rebuild, a - *stronger* stock baseline). The **prefill M>1 gap is architectural** (native lacks the M>1 GEMM) and - holds regardless of quant. Recommend an og-builder qwen rebuild for a symmetric prefill comparison, and - an accuracy check on the rebuilt glm before quoting the decode number externally. - -#### Axis 3b — long-prompt prefill sweep (prefill IS a Marlin M>1 gate) - -`marlin_bench --prefill-lens 128,512,1024` times a single batched `decode(prompt,0)` (M=len forward, the -MatMulNBits M>1 path), KV reset via `rewind(0)` between iters, **min across 4 iters** (+1 warmup), -contention-invariant. Prefill tok/s = len / wall. ORT-genai glm via `bench-scripts/ort_glm_bench.py`. - -| Prompt len | glm native | glm ORT-genai | qwen native | native/ORT (glm) | -|---|---|---|---|---| -| **128** | 199.7 tok/s (641 ms) | **8 002 tok/s** (16.0 ms) | 147.1 tok/s (870 ms) | **0.025× (40× slower)** | -| **512** | 214.2 tok/s (2 390 ms) | **21 044 tok/s** (24.3 ms) | 152.2 tok/s (3 363 ms) | **0.010× (98× slower)** | -| **1024** | 217.1 tok/s (4 716 ms) | **26 331 tok/s** (38.9 ms) | 153.7 tok/s (6 664 ms) | **0.008× (121× slower)** | - -- **This is the decisive Marlin BEFORE number.** Native prefill throughput is **flat ~150–217 tok/s** and - scales *linearly with wall* (L=1024 wall ≈ 8× L=128 wall) — i.e. the M=len forward gets **no** batching - efficiency: each extra prompt token costs ~a full decode step because MatMulNBits has no efficient M>1 - int4 GEMM. ORT-genai's prefill *rises* with length (8k→26k tok/s) as a real GEMM amortizes launch cost. - The **~100–120× gap** is precisely the wall Marlin (efficient M>1 int4 GEMM) must collapse; after Marlin - the *same* `--prefill-lens` command must show native prefill rising toward the ORT curve. -- Prefill wall includes full-sequence logits materialization (D2H of all-position logits + full lm_head); - it is a prefill-to-first-token-ready wall, consistent BEFORE/AFTER since AFTER uses the same path. - -#### Axis 4 — capture-safety re-probe (wired to re-run against Marlin) - -`leverb_phase0` probe (`onnx-genai-engine`, `#[cfg(all(test, feature="cuda"))]`, -`leverb_{phase0,increment0}_capture_attempt`) on glm-4-9b. Target AFTER Marlin/Lever-B: `segments`→1 (or -near) at M>1 and captured-M=8 ≈ M=1 (break-even B\* ≤ ~2). **Baseline reconfirmed on `94c69fdc` -(GPU 7):** - -- **INC0 M=8 capture: `captured=true`, `segments=41`** (matches Deckard's baseline exactly). -- **DECISIVE captured replay wall: M=8 = 87.7 ms, M=1 = 10.0 ms → ratio = `8.77×`** (Deckard prior 8.58×; - the +0.2× is co-tenant contention on the M=8 wall). VERDICT = **NO-GO** for Lever B today, as expected. -- **Segment root cause (41 seams):** `GroupQueryAttention[KernelCaptureUnsupported]×40`, - `MatMulNBits×240`, `MatMul×1`, `SkipSimplifiedLayerNormalization×80` — the M=K kernels aren't - capture-safe yet. This is the surface Lever B must make persistent-binding-clean. -- **Cross-validation:** the probe's Part-C eager sweep independently reproduces `marlin_bench` Axis 2 — - M=1 14.5 ms, M=8 92.9 ms, **cliff 67.2 ms, tail 1.88 ms/row** — two independent harnesses agree on the - cliff/tail to within noise, and captured M=1 = 11.1 ms / 90.3 tok/s. The probe is already wired to - re-run against Marlin (command below). - ---- - -#### Reproduce (exact commands, head `94c69fdc`) - -```bash -source /home/justinchu/onnx-genai/.cudaenv.sh -nvidia-smi --query-compute-apps=pid,used_memory --format=csv,noheader # verify idle; pin a clean high index - -# Build -cargo build --release -p onnx-genai-bench --features bench-native,cuda --bin marlin_bench -cargo build --release -p onnx-genai-bench --features bench-native,bench-ort,cuda --bin profile_native - -# Axis 2 (M-scaling wall) + Axis 3 (native e2e) -CUDA_VISIBLE_DEVICES=7 ./target/release/marlin_bench \ - --model /home/justinchu/glm-e2e-artifacts/glm-4-9b-int4-cuda --label glm-4-9b-int4 \ - --past 512 --iters 30 --warmups 5 --m-list 1,2,4,8,16 --e2e --e2e-tokens 128 -CUDA_VISIBLE_DEVICES=7 ./target/release/marlin_bench --model .../glm-4-9b-int4-cuda --past 32 ... # Increment-0 repro -CUDA_VISIBLE_DEVICES=7 ./target/release/marlin_bench \ - --model /home/justinchu/shared-models/qwen2.5-14b-instruct-int4-zp-onnx --label qwen2.5-14b-int4 \ - --past 512 --iters 30 --warmups 5 --m-list 1,2,4,8,16 --e2e --e2e-tokens 128 - -# Axis 1 (M=1 DRAM%) — ncu on the int4 GEMV (needs sudo; RmProfilingAdminOnly=1) -NCU=/usr/local/cuda-13.3/bin/ncu -sudo -E env PATH="$PATH" LD_LIBRARY_PATH="$LD_LIBRARY_PATH" CUDA_VISIBLE_DEVICES=7 \ - $NCU --graph-profiling node -k regex:matmul_nbits_gemv --launch-skip 400 --launch-count 8 --csv \ - --metrics gpu__dram_throughput.avg.pct_of_peak_sustained_elapsed,sm__throughput.avg.pct_of_peak_sustained_elapsed,dram__bytes_read.sum,gpu__time_duration.sum,launch__grid_size \ - ./target/release/profile_native --model --ep cuda --backend native --steady --warmups 1 --runs 1 --tokens 96 - -# Axis 3 (ORT head-to-head) — force the CUDA-enabled ORT 1.27 -ONNX_GENAI_ORT_LIB_DIR=$ORT_ROOT/lib CUDA_VISIBLE_DEVICES=7 \ - ./target/release/profile_native --model --ep cuda --backend ort --steady --warmups 1 --runs 3 --tokens 128 --decode-skip 8 - -# Axis 3b (long-prompt prefill sweep) — native, both models -CUDA_VISIBLE_DEVICES=7 ./target/release/marlin_bench \ - --model /home/justinchu/glm-e2e-artifacts/glm-4-9b-int4-cuda --label glm-4-9b-int4-native \ - --m-list 1 --past 32 --iters 3 --warmups 1 --e2e --e2e-tokens 96 --prefill-lens 128,512,1024 --prefill-iters 4 -CUDA_VISIBLE_DEVICES=7 ./target/release/marlin_bench \ - --model /home/justinchu/shared-models/qwen2.5-14b-instruct-int4-zp-onnx --label qwen2.5-14b-int4-native \ - --m-list 1 --past 32 --iters 2 --warmups 1 --prefill-lens 128,512,1024 --prefill-iters 4 - -# glm ORT side: re-export via onnxruntime-genai model builder (native glm ONNX is rejected by all stock ORT), -# then run the og decode+prefill bench. CUDA-12 runtime libs needed for og's bundled ORT. -CU12=$(ls -d /usr/lib/python3.12/site-packages/nvidia/*/lib | tr '\n' ':') -CUDA_VISIBLE_DEVICES=7 LD_LIBRARY_PATH="$CU12$LD_LIBRARY_PATH" \ - /home/justinchu/.conda/envs/onnx/bin/python -m onnxruntime_genai.models.builder \ - -i /home/justinchu/glm-e2e-artifacts/glm-4-9b-gptq-source -o ./ort-glm-build -p int4 -e cuda -c ./ort-glm-cache -cp /home/justinchu/glm-e2e-artifacts/glm-4-9b-int4-cuda/tokenizer.json ./ort-glm-build/tokenizer.json -CUDA_VISIBLE_DEVICES=7 LD_LIBRARY_PATH="$CU12$LD_LIBRARY_PATH" \ - /home/justinchu/.conda/envs/onnx/bin/python bench-scripts/ort_glm_bench.py ./ort-glm-build 128,512,1024 128 3 - -# Axis 4 (capture re-probe) -CUDA_VISIBLE_DEVICES=7 ONNX_GENAI_RUN_CUDA_SMOKE=1 \ - ONNX_GENAI_LEVERB_MODEL=/home/justinchu/glm-e2e-artifacts/glm-4-9b-int4-cuda \ - cargo test -p onnx-genai-engine --features cuda,native-backend --release leverb_phase0 -- --ignored --nocapture -``` - -#### Portability (Rule 11) - -Every number above is the **current split-K / `gemv_f16_general_bs` fallback layout** (pre-Marlin). -Marlin is SM80+; because these are the fallback-path numbers, re-running this *identical* harness after -Marlin lands — on a 1 (plain + rmsnorm-prologue + gate_up SwiGLU fused) + split-K + GQA/SkipLN M>1 + capture-safety valves + lm_head 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. +- **#961 (Pris, 401af46f):** reusable f64 dequant→GEMM numerics oracle gate with a justified tolerance + envelope; tiled baseline sits ~1 fp16 ULP inside. +- **#962 (Sebastian, df6d3afb):** `marlin_bench` perf/capture harness + the BEFORE/AFTER B\* tables that + produced the capture arc. + +## Reviews + +- **Chew (numerics) 🟡 APPROVE-WITH-NOTES:** 11/11 GPU parity tests pass on H200, no correctness bug; notes: + keep the soft argmax-stability token guarantee opt-in (N1), log/count the silent tiled-fallback (N2), + nibble-int-zp only (N3). +- **Gaff (quality) 🟢 APPROVE:** zero blocking defects; Rule 11 portability PASS, env-var honesty PASS, + capture-safety valve sound, fmt/clippy clean; one trivial `cfg(test)` `clippy::unusual_byte_groupings` note. + +## Per-model honesty + +- **glm-4-9b (block-128):** the clean practical GO at **B\*=2.16×**. CORRECTED attribution (Deckard update-10 + supersedes update-8): the fused gate_up split-K commit is a NO-OP for glm (fused node needs block-32); + glm's 2.63→2.16× was ENTIRELY the general small-M split-K retune. +- **qwen2.5-14b (block-32):** capture fully solved (whole-graph, zero seams, byte-identical eager) but + **B\*≈4.7×** — an honest denominator effect (fast tuned block-32 M=1 inflates the ratio), a drafting-depth + follow-up, NOT a kernel bug. + +## Housekeeping + +- Merged all 5 inbox drops (deckard-marlin-kernel, pris-marlin-numerics, sebastian-marlin-bench, + chew-marlin-numerics-review, gaff-marlin-quality-review) into a single consolidated Marlin program entry in + decisions.md; deleted all 5 drops (README.md kept). +- HARD-GATE size: decisions.md was **45,991 B** and this consolidation would have pushed it over 51,200 B → + archived the three verbose "Last consolidated" chronicle lines + the detailed #957/#948/#949 spec-capture & + Lever B sub-entry bodies (now RESOLVED by this landing) to `decisions-archive/2026-08.md` (live keeps + compact pointers) → **42,104 B** after, under the 50 KB gate. +- HISTORY chronicle gate: **sebastian** history (12,308 B, ~13 dated entries) was a chronicle over the + >8-entry threshold → compacted (08-12/13 decode arc moved to `history-archive.md`, live keeps role + + durable lessons + a one-paragraph context + the 08-14 entries). Appended the Marlin outcome to deckard, + sebastian, pris, chew, gaff; the other four stayed under both the chronicle and 15,360 B gates. +- Wrote orchestration logs for deckard, sebastian, pris, chew, gaff. +- Charter NOTE: the spawn prompt asked to archive by AGE (older-than-30-days); per Scribe charter I archived + by SIZE and say so. Committed on `chore/scribe-marlin` (main is protected — the coordinator opens/admin- + merges the PR). + +## Standing directive (updated) + +Marlin (Lever A) is DELIVERED and the spec-capture condition is MET for glm at B\*=2.16×; the #957 +Increment-0 re-probe is DONE (this landing IS the post-Marlin measurement). Now unblocked (per #957 Stage 4): +fund the actual speculative build (verify sub-graph + capture-stable selective KV-commit + exact-greedy +near-tie guard #935 + draft sources) — glm is a practical GO; deeper/stronger drafting is the lever for +qwen's B\*. Marlin stays opt-in default-OFF; close Chew N1/N2 before any default flip. diff --git a/.squad/orchestration-log/2026-08-14T21-43-12Z-chew.md b/.squad/orchestration-log/2026-08-14T21-43-12Z-chew.md new file mode 100644 index 0000000000..dd43afccc7 --- /dev/null +++ b/.squad/orchestration-log/2026-08-14T21-43-12Z-chew.md @@ -0,0 +1,39 @@ +# Orchestration Log — chew + +**Timestamp:** 2026-08-14T21:43:12Z +**Agent:** Chew (🔬 Numerics & Precision reviewer) +**Task:** Numerics correctness review of PR #960 (Marlin int4 M>1 GEMM) +**Requested by:** Coordinator (Justin @justinchuby) +**Mode:** review (read-only) + +## Summary + +Numerics-only review of Deckard's Marlin fp16×int4 tensor-core GEMM (Gaff owns quality/portability). +**Verdict: 🟡 APPROVE-WITH-NOTES — numerically sound, ship; no correctness bug; three non-blocking notes.** + +## Findings + +- **int4 dequant + relayout/repack CORRECT:** host `repack_int4_weights` == device repack (byte-identical, + all group sizes); per-lane byte layout matches the `mma.sync.m16n8k16` B-fragment distribution exactly; + asymmetric zp indexing matches existing int4 tiled kernels; **scale-after-accumulate** is exact factoring + (scale/zp constant across K within a group); padding columns cannot leak into valid outputs. +- **split-K fp32 partials SOUND:** disjoint contiguous group slices fully cover the partials buffer; fixed + reduction order ⇒ deterministic/capture-stable; only re-associates cross-group fp32 sums (benign). +- **lm_head cached dense-GEMM plan CLAIM HOLDS:** cuBLASLt heuristic is deterministic for a fixed + shape/dtype/device/workspace ⇒ cached algo == dynamic algo; caching changes only *when* selection happens. +- **Empirical:** ran the in-crate GPU parity suite on H200 → **11/11 pass** (incl. f64-oracle parity, + split-K parity + determinism, repack roundtrip all group sizes). +- **#961 tolerance:** appropriately tight (~8× fp16-ULP headroom) — engineering-justified, concur. + +## Notes (non-blocking) + +- **N1:** "byte-identical greedy tokens" is a soft argmax-stability guarantee (empirical, 2 models × 24 + tokens), NOT a numeric invariant — keep the flag opt-in; add prompt/model coverage before any default flip. +- **N2:** hard Marlin launch errors are silently swallowed into the tiled fallback — log/count 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). + +## Outcome + +🟡 APPROVE-WITH-NOTES recorded; drop `chew-marlin-numerics-review.md` merged into decisions.md and deleted. diff --git a/.squad/orchestration-log/2026-08-14T21-43-12Z-deckard.md b/.squad/orchestration-log/2026-08-14T21-43-12Z-deckard.md new file mode 100644 index 0000000000..d25bad66c4 --- /dev/null +++ b/.squad/orchestration-log/2026-08-14T21-43-12Z-deckard.md @@ -0,0 +1,37 @@ +# Orchestration Log — deckard + +**Timestamp:** 2026-08-14T21:43:12Z +**Agent:** Deckard (⚙️ CUDA / decode performance engineer) +**Task:** Build Marlin fp16×int4 tensor-core GEMM (Lever A) — the funded PRIMARY decode lever + the single condition of the #957 spec-capture CONDITIONAL-GO +**Requested by:** Coordinator (Justin @justinchuby) +**Mode:** background (multi-stage, updates 1→10) + +## Summary + +Built + shipped a from-scratch SM80 `mma.sync.m16n8k16` fused fp16×int4 tensor-core GEMM and wired it +through the `MatMulNBits` M>1 path, closing the M=8 speculative-verify capture cliff. **Outcome: Lever A +LANDED & MERGED (#960, 7774ec5b); the #957 spec-capture CONDITIONAL-GO is now GO in practice for glm-4-9b.** + +## Deliverables / Findings + +- **Kernel:** new module `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). +- **Wiring:** `MatMulNBits` M>1 (plain + rmsnorm-prologue + gate_up SwiGLU fused) + split-K + GQA/SkipLN + M>1 capture-safety valves + lm_head cached dense-GEMM capture plan. +- **Portability (Rule 11):** opt-in `ONNX_GENAI_MARLIN_M_GT_1` (default OFF, SM80 arch guard, byte-identical + tiled fallback on 1 GEMM) +**Requested by:** Coordinator (Justin @justinchuby) +**Mode:** review (read-only) + +## Summary + +Quality review of Deckard's Marlin M>1 GEMM (numerics = Chew; the B\*=2.16 plateau is an accepted honest +plateau, not a defect). Frozen head `a11facbc` (code `3735d57e`). **Verdict: 🟢 APPROVE — zero blocking +defects; one trivial cosmetic test-only note.** + +## Findings + +- **Rule 11 portability 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 — a + force-set env var on 1 Fused + flash path (on-device totals, no host read-back) and SkipLN latch relaxation (shape-keyed cache rejects + mid-capture shape change) both verified. +- **Build/lint PASS:** `cargo fmt --all --check` clean; exact CUDA and engine `-D warnings` clippy gates + clean; all `unsafe` blocks carry SAFETY justifications. + +## Note (cosmetic, non-blocking) + +- `matmul_nbits.rs:7545` test-only `0xcafef00d_1234_5678u64` trips `clippy::unusual_byte_groupings` + (warn-level), but it sits in a `#[cfg(test)]` module the CUDA clippy gate does not compile (no + `--all-targets`), so it does not break CI; trivially fixable. + +## Outcome + +🟢 APPROVE recorded; drop `gaff-marlin-quality-review.md` merged into decisions.md and deleted. diff --git a/.squad/orchestration-log/2026-08-14T21-43-12Z-pris.md b/.squad/orchestration-log/2026-08-14T21-43-12Z-pris.md new file mode 100644 index 0000000000..96cb1bcacb --- /dev/null +++ b/.squad/orchestration-log/2026-08-14T21-43-12Z-pris.md @@ -0,0 +1,35 @@ +# Orchestration Log — pris + +**Timestamp:** 2026-08-14T21:43:12Z +**Agent:** Pris (🧪 Tester — numerics/oracle gates) +**Task:** Build the f64 dequant→GEMM numerics oracle gate for int4 `MatMulNBits` (Marlin sign-off criterion) +**Requested by:** Coordinator (Justin @justinchuby) +**Mode:** background (parallel with Deckard's kernel) + +## Summary + +Landed a reusable f64 oracle test that any int4 GEMM (today's tiled/GEMV path AND Deckard's Marlin) must +pass before shipping — the apples-to-apples, op-level red/green target for Deckard and the sign-off +criterion for Chew. **Outcome: numerics gate merged (#961, 401af46f); tiled baseline sits ~1 fp16 ULP +inside the envelope, giving Marlin a defensible reference.** + +## Deliverables / Findings + +- **Oracle:** self-contained integration test + `crates/onnx-runtime-ep-cuda/tests/matmul_nbits_marlin_numerics.rs` (no edits to `matmul_nbits.rs`, so no + conflict with `squad/marlin-kernel`). Dequant `(code-zp)·scale` accumulated in f64, sharing the + candidate's fp16-rounded activations + rounded scale ⇒ residual isolates only accumulation + fp16 output + rounding. Drives the op purely through ONNX semantics, so Marlin validates with zero test changes. +- **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 envelope** (`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. +- **Baseline (H200):** tiled/GEMV path measures ~8–10× inside abs, ~6× inside rel (≈1 fp16 ULP). Marlin must + land in the same envelope. +- **Harness lesson:** pre-zero the output + `synchronize()` after execute — an un-synced run reading stale + device-pool memory faked an abs-81.8 divergence; not a kernel bug. + +## Outcome + +Merged PR #961 (401af46f). Chew concurred the tolerance is engineering-justified, not a rubber stamp +(verified in his #960 numerics review). Drop `pris-marlin-numerics.md` merged into decisions.md and deleted. diff --git a/.squad/orchestration-log/2026-08-14T21-43-12Z-sebastian.md b/.squad/orchestration-log/2026-08-14T21-43-12Z-sebastian.md new file mode 100644 index 0000000000..3245e9cb91 --- /dev/null +++ b/.squad/orchestration-log/2026-08-14T21-43-12Z-sebastian.md @@ -0,0 +1,34 @@ +# Orchestration Log — sebastian + +**Timestamp:** 2026-08-14T21:43:12Z +**Agent:** Sebastian (📊 runtime/server performance & benchmark analysis) +**Task:** Marlin int4 GEMM perf/capture harness + reproducible BEFORE/AFTER B* tables (gates #957 Lever A) +**Requested by:** Coordinator (Justin @justinchuby) +**Mode:** background (paired with Deckard's kernel iteration) + +## Summary + +Built the perf/capture harness that measured Deckard's Marlin kernel end-to-end and produced the +authoritative BEFORE→AFTER capture B\* arc. **Outcome: harness merged (#962, df6d3afb); confirmed the +glm-4-9b M=8 verify graph reaches a SINGLE whole-graph capture at B\*=2.16× — spec-capture GO in practice.** + +## Deliverables / Findings + +- **`marlin_bench`** (`crates/onnx-genai-bench/src/bin/marlin_bench.rs`, `--features bench-native,cuda`): + times a real `decode_verify` over M∈{1,2,4,8,16} + prefill L∈{128,512,1024}; prints median/p10/p90/max + + device `compute_cap` so every number is arch-attributable and contamination is visible. +- **Capture B\* arc (glm-4-9b, H200 GPU7):** segments **41→120→1** (whole-graph, zero seams); B\* + **8.76→4.99 (Marlin) →2.71 (split-K) →2.63 (lm_head dense-plan) →2.16× (small-M retune)**; captured M=8 + verify wall stable ~21.9 ms; byte-identical greedy tokens throughout. +- **Prefill ~2×:** glm 218→426 tok/s @ L=1024, halving the glm-vs-ORT prefill gap 121×→62×. +- **Per-model:** glm (block-128) clean practical GO at 2.16×; qwen (block-32) capture fully solved + (segments=1, zero unsupported) but B\*≈4.62–4.72× — an honest denominator effect (fast tuned block-32 + M=1 GEMV: capture speeds M=1 34% vs the M=8 verify 18%), not a coverage gap. +- **Flaky-test catch:** the qwen exact-token parity test is flaky ~25% because it asserts equality vs a + nondeterministic tiled reference at a near-tie (tok19); Marlin's fixed split-K reduction is the + DETERMINISTIC path — harden the assertion (tolerance / deterministic oracle), not a Marlin regression. + +## Outcome + +Merged PR #962 (df6d3afb) on base `94c69fdc`. Fed the corrected per-model attribution back to Deckard +(update-10). Drop `sebastian-marlin-bench.md` merged into decisions.md and deleted.