docs(research): speculative-capture feasibility — CONDITIONAL-GO gated behind Marlin by justinchuby · Pull Request #957 · justinchuby/onnx-genai

justinchuby · 2026-08-14T17:59:34Z

Speculative decoding under CUDA-graph capture — feasibility (DESIGN-ONLY)

Answers Justin's question — "想办法把 speculative decoding 的 capture 做好？有办法吗？" — with a concrete, honest staged path. No kernels touched; adds one research doc + one decision drop.

Verdict: CONDITIONAL-GO, and the single condition is Marlin (Lever A) — which we are already building as the primary decode lever.

Today (measured, research(decode): Lever B Phase-0 capture-stability probe — gated-GO on Increment-0 #948/research(decode): Lever B Increment-0 — decisive captured M=K number → NO-GO, promote Marlin #949): captured M=8 verify = 87.2 ms = 8.58× captured M=1 (10.2 ms) ⇒ acceptance break-even B* = 8.5 ⇒ physically impossible ⇒ guaranteed loss (matches the settled speculative KILL).
Root cause (code-grounded, no new GPU run): the ~80 ms floor is dominated by the 240 MatMulNBits generic 16×16 CUDA-core tiled GEMMs — M>1 leaves the tuned fused int4 GEMV (matmul_nbits.rs:4803-4869). GQA (efficient flash) and norm (same kernel) live in the cheap 1.85 ms/row tail.
Marlin is the enabler, not just a speedup: static launch grid ⇒ capture-safe by construction (kills the 240 seams) + weights reused across the M rows ⇒ collapses the 67 ms cliff. Estimated post-Marlin verify ~12–20 ms ⇒ break-even B* ≈ 1.2–2.0, cleared by prompt-lookup (B≈2–4) and EAGLE-3/MTP.
Residual blockers after Marlin are cheap: SkipSimplifiedLayerNorm = a num_groups==1 signature relaxation (normalization.rs:2306); GQA = pre-warmed fixed-M flash workspace (group_query_attention.rs:3099). Not kernel rewrites.

Staged plan with hard gates

Marlin (already primary; ref docs/research/decode-remaining-levers-feasibility.md) + a new "capture-safe at fixed M=K" contract item.
GQA + norm M>1 capture-support (the two cheap fixes) — gate: seam census collapses to ~1 segment.
Re-run THIS Increment-0 probe for the true post-Marlin captured M=K wall — fund/kill gate: measured B* ≤ ~2 ⇒ GO, ≳4 ⇒ NO-GO.
Only then build the verify-graph + capture-stable KV-commit + near-tie guard + draft sources.

Context

Refines (does not overturn) the settled Lever B NO-GO — research(decode): Lever B Phase-0 capture-stability probe — gated-GO on Increment-0 #948 (Phase-0) / research(decode): Lever B Increment-0 — decisive captured M=K number → NO-GO, promote Marlin #949 (Increment-0).
Marlin scoping: docs/research/decode-remaining-levers-feasibility.md.
Portability (Rule 11): Marlin is SM-version-specific ⇒ arch guard + split-K fallback + degrade-to-disabled on non-Marlin tiers. H200/glm-4-9b-int4-scoped.

⚠️ Design/research doc — no production code changed. A squad member should sanity-check the post-Marlin verify-wall estimate (clearly labeled as an estimate; Stage-3 measures it).

…d behind Marlin Design-only feasibility doc answering whether CUDA-graph-captured speculative decoding can be a net win. Grounds the answer in the settled Lever B NO-GO (#948/#949: captured M=8 verify = 8.58x captured M=1) and the M>1 kernel dispatch paths. Root cause (code-grounded): the ~80 ms M=8 floor is dominated by the 240 MatMulNBits generic 16x16 CUDA-core tiled GEMMs (M>1 leaves the tuned fused int4 GEMV), not GQA/norm. Marlin (Lever A, already primary) is the enabler: capture-safe by construction + weights reused across M rows collapses the cliff. Estimated post-Marlin verify ~12-20 ms moves acceptance break-even from B*=8.5 (impossible today) to B*~1.2-2.0 (cleared by prompt-lookup/EAGLE-3). Residual blockers after Marlin are cheap: SkipSimplifiedLayerNorm = signature relaxation; GQA = pre-warmed fixed-M flash workspace. Staged plan with hard Stage-3 re-measure gate before funding any verify-graph build. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

Merge Deckard's speculative-capture feasibility decision drop into decisions.md (CONDITIONAL-GO gated behind Marlin; break-even B* = 8.5 today → ~1.2–2.0 post-Marlin). Adds orchestration + session logs. State-only; follows #957. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

codecov · 2026-08-14T18:05:50Z

Codecov Report

✅ All modified and coverable lines are covered by tests.
✅ Project coverage is 79.01%. Comparing base (4d4c227) to head (f82d9eb).
⚠️ Report is 2 commits behind head on main.

Additional details and impacted files

@@            Coverage Diff             @@
##             main     #957      +/-   ##
==========================================
- Coverage   79.60%   79.01%   -0.60%     
==========================================
  Files         356      356              
  Lines      150096   150096              
  Branches   150096   150096              
==========================================
- Hits       119486   118596     -890     
- Misses      25955    26849     +894     
+ Partials     4655     4651       -4

Flag	Coverage Δ
cli-ort-linux	`83.52% <ø> (ø)`
cli-ort-windows	`83.11% <ø> (ø)`
mlas	`80.96% <ø> (ø)`
offline	`78.81% <ø> (-0.63%)`	⬇️

Flags with carried forward coverage won't be shown. Click here to find out more.
see 6 files with indirect coverage changes

🚀 New features to boost your workflow:

❄️ Test Analytics: Detect flaky tests, report on failures, and find test suite problems.

github-actions · 2026-08-14T18:25:37Z

🔴 Benchmark Regression Detected

Comparison of criterion micro-benchmarks: PR head vs merge-base, measured on the same runner in the same job (base first → PR second).

ℹ️ Absolute times are informational only — they vary with runner load. The % change column is the reliable signal because both sides ran under identical conditions.

Status	Scenario	Base	PR	Change
🔴	`block_quantized_matmul_cached_dense/mxfp4_cached_dense_repeated_call/1x1024x1024`	45.17 µs	129.03 µs	+185.7%
🔴	`block_quantized_moe_cached_dense/mxfp4_uncached_expert_dequant_each_call/rows=1,H=256,I=256,E=4,top_k=1`	358.95 µs	690.44 µs	+92.3%
🔴	`gather/medium_f32_threads=1-internal/32768`	3.67 µs	5.18 µs	+41.2%
🔴	`block_quantized_matmul_cached_dense/mxfp4_preexpanded_dense_oncelock_like_proxy/1x1024x1024`	73.62 µs	101.67 µs	+38.1%
⚠️	`gather/large_bf16_threads=1-internal/131072`	11.46 µs	13.89 µs	+21.2%
⚠️	`add/large_f32_threads=1-internal/4194304`	37.44 ms	45.20 ms	+20.7%
⚠️	`qwen3_sampling_processors/top_p_fast_after_top_k`	479.45 µs	576.77 µs	+20.3%
⚠️	`matmul/small_generic_bf16_threads=8/1x256x256`	34.75 µs	41.07 µs	+18.2%
⚠️	`matmul/medium_generic_f16_threads=1/32x512x512`	30.27 µs	35.70 µs	+18.0%
⚠️	`add/small_bf16_threads=1-internal/1024`	12.43 µs	14.59 µs	+17.4%
⚠️	`gather/small_bf16_threads=1-internal/4096`	471.8 ns	552.0 ns	+17.0%
✅	`add/large_f16_threads=1-internal/4194304`	39.37 ms	45.24 ms	+14.9%
✅	`gather/small_f32_threads=1-internal/4096`	621.1 ns	712.0 ns	+14.6%
✅	`qwen3_sampling_processors/top_k_top_p_full_sort_baseline`	5.24 ms	5.93 ms	+13.3%
✅	`add/medium_bf16_threads=1-internal/262144`	2.43 ms	2.75 ms	+13.2%
✅	`matmul/medium_generic_f16_threads=8/32x512x512`	31.23 µs	35.17 µs	+12.6%
✅	`reduce_mean/small_f32_threads=1-internal/4096`	14.66 µs	16.40 µs	+11.9%
✅	`matmul/medium_generic_f32_threads=8/32x512x512`	1.44 ms	1.61 ms	+11.7%
✅	`reduce_mean/medium_f32_threads=1-internal/65536`	249.71 µs	278.23 µs	+11.4%
✅	`add/medium_f16_threads=1-internal/262144`	2.55 ms	2.83 ms	+10.9%
✅	`add/large_bf16_threads=1-internal/4194304`	41.75 ms	45.76 ms	+9.6%
✅	`block_quantized_moe_cached_dense/mxfp4_cached_dense_expert_repeated_call/rows=1,H=256,I=256,E=4,top_k=1`	81.41 µs	89.02 µs	+9.3%
✅	`matmul/large_generic_bf16_threads=8/32x1024x1024`	1.70 ms	1.86 ms	+9.3%
✅	`sampling_latency/top_k_per_token`	51.25 µs	55.90 µs	+9.1%
✅	`qwen3_sampling_processors/top_k_top_p_fast`	606.59 µs	659.97 µs	+8.8%
✅	`reduce_mean/large_f32_threads=1-internal/262144`	1.04 ms	1.13 ms	+8.7%
✅	`qwen3_sampling_processors/top_k_full_sort_baseline`	2.08 ms	2.25 ms	+8.3%
✅	`qwen3_sampling_processors/top_k_partial_selection`	137.56 µs	147.94 µs	+7.5%
✅	`tokenization/decode_tokens_per_second`	6.36 ms	6.76 ms	+6.3%
✅	`qwen3_sampling_processors/top_p_full_sort_after_top_k_baseline`	3.36 ms	3.57 ms	+6.2%
✅	`gather/medium_bf16_threads=1-internal/32768`	2.46 µs	2.59 µs	+5.5%
✅	`tokenization/encode_tokens_per_second`	380.16 µs	398.61 µs	+4.9%
✅	`add/medium_f32_threads=1-internal/262144`	2.54 ms	2.65 ms	+4.3%
✅	`add/small_f16_threads=1-internal/1024`	14.07 µs	14.61 µs	+3.8%
✅	`kv_cache/alloc_dealloc_pages`	36.18 µs	37.54 µs	+3.8%
✅	`gather/medium_f16_threads=1-internal/32768`	2.65 µs	2.74 µs	+3.8%
✅	`matmul/medium_generic_f32_threads=1/32x512x512`	2.63 ms	2.69 ms	+2.1%
✅	`gather/small_f16_threads=1-internal/4096`	501.0 ns	505.9 ns	+1.0%
✅	`grammar_masking/llguidance_compute_mask/32`	74.13 µs	74.29 µs	+0.2%
✅	`logit_processing/seven_processor_chain_per_step`	311.92 µs	310.61 µs	-0.4%
✅	`sampling_latency/greedy_per_token`	3.33 µs	3.29 µs	-0.9%
✅	`matmul/small_generic_f32_threads=1/1x256x256`	41.42 µs	40.96 µs	-1.1%
✅	`sampling_latency/top_p_per_token`	380.44 µs	364.97 µs	-4.1%
✅	`matmul/small_generic_bf16_threads=1/1x256x256`	43.26 µs	40.92 µs	-5.4%
✅	`matmul/medium_generic_bf16_threads=1/32x512x512`	594.31 µs	559.96 µs	-5.8%
✅	`sampling_latency/min_p_per_token`	211.07 µs	198.28 µs	-6.1%
✅	`matmul/small_generic_f32_threads=8/1x256x256`	54.25 µs	50.05 µs	-7.7%
✅	`add/small_f32_threads=1-internal/1024`	249.3 ns	225.3 ns	-9.6%
✅	`block_quantized_matmul_cached_dense/mxfp4_uncached_dequant_each_call/1x1024x1024`	657.87 µs	588.21 µs	-10.6%
✅	`matmul/large_generic_f32_threads=1/32x1024x1024`	10.88 ms	9.72 ms	-10.7%
✅	`matmul/small_generic_f16_threads=1/1x256x256`	41.17 µs	35.89 µs	-12.8%
✅	`matmul/large_generic_bf16_threads=1/32x1024x1024`	2.13 ms	1.84 ms	-13.5%
✅	`gather/large_f16_threads=1-internal/131072`	16.98 µs	14.62 µs	-13.9%
🟢	`gather/large_f32_threads=1-internal/131072`	40.07 µs	32.74 µs	-18.3%
🟢	`matmul/small_generic_f16_threads=8/1x256x256`	49.21 µs	39.27 µs	-20.2%
🟢	`matmul/medium_generic_bf16_threads=8/32x512x512`	452.69 µs	357.56 µs	-21.0%
🟢	`matmul/large_generic_f16_threads=1/32x1024x1024`	99.36 µs	73.96 µs	-25.6%
🟢	`matmul/large_generic_f16_threads=8/32x1024x1024`	118.76 µs	80.18 µs	-32.5%
🟢	`matmul/large_generic_f32_threads=8/32x1024x1024`	6.07 ms	3.87 ms	-36.2%

Visual flags: ⚠️ ≥ 15% slower, 🔴 ≥ 30% slower — calibrated against measured runner noise (~27% worst-case on multi-threaded matmul)

Host info

CPU: Apple M1 (Virtual)
Cores: 3
OS: Darwin 25.5.0 arm64
Rust: rustc 1.97.1 (8bab26f4f 2026-07-14)
Load avg: { 2.88 3.30 5.70 }

What this cannot catch

Regressions in code paths not covered by these benchmarks (e.g., end-to-end decode with a real model)
Sub-threshold regressions that compound over multiple PRs
Performance changes that only manifest under GPU execution
Latency changes in the ORT integration path (these benchmarks exercise the native Rust kernels)

Extend the Marlin int4 M>1 path to cover the gate/up SwiGLU MLP fusion so the MLP no longer falls back to the portable tiled two-GEMM at M>1. This is the bulk of prefill/verify cost and the remaining gap toward zero M>1 tiled-fallback nodes (the capture segments->1 unlock). - marlin_gemm.rs: add `ensure_scratch` size-keyed scratch pool (slot-keyed, FIFO cap 256, rejects cold-miss during capture) so warm replays allocate nothing and stay capture-safe. - matmul_nbits.rs: `try_launch_marlin_gemm_rmsnorm` now uses the pooled scratch (returns Option<bool> warm flag); add `try_launch_marlin_gate_up_prefill` (paired gate/up Marlin GEMMs + shared SiluMul, optional rmsnorm prologue) and wire it into `run_f16_gate_up_swiglu` M>1 branch. - 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 & byte-identical. Refs #957, #960. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

…xport marlin_bench: new --prefill-lens/--prefill-iters mode (Axis 3b) times a single batched decode(prompt,0) M=len forward (the MatMulNBits M>1 path), KV reset via rewind(0), contention-invariant MIN. Prefill is a direct Marlin gate (M>1 is where Marlin wins). Resolve the glm ORT head-to-head: all stock ORT (1.27 EP, pip 1.28, og 0.14.1 bundled) reject the native-authored glm ONNX GQA (rotary_embedding_dim). Re-export a clean stock ORT-genai chatglm int4 via the onnxruntime-genai model builder from glm-4-9b-gptq-source and benchmark it (bench-scripts/ort_glm_bench.py). BEFORE finding: native prefill is flat ~150-217 tok/s (no M>1 GEMM efficiency) vs ORT-genai glm 8k-26k tok/s -> ~100-120x gap = the decisive Marlin gate. glm decode native ~95 vs ORT-genai 233 tok/s. Decision drop updated with the prefill sweep table + glm-ORT result + reproduce commands. Refs #957, #962; docs/research/decode-remaining-levers-feasibility.md §3/§4 Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

Add a tracer-driven coverage audit to the M>1 e2e test that runs a prefill-heavy generation with ONNX_GENAI_MARLIN_M_GT_1=1, tallies the per-op kernel_variant annotations, and asserts NO MatMulNBits node fell back to a tiled M>1 kernel (gemm_f16_tiled{,_rmsnorm}, gate_up_swiglu{,_rmsnorm}_prefill). Zero tiled M>1 fallbacks is the prerequisite for the capture segments->1 collapse. Measured (H200, GPU 7): - glm-4-9b-int4 : 240 gemm_marlin_int4, 0 tiled M>1 - qwen2.5-14b : 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. Refs #957, #960. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

P2 "widen the win": adds a split-K variant of the Marlin int4 GEMM to attack the occupancy-bound small-M regime, where the direct kernel launches only ~1-1.5 waves of blocks (N=5120 => ~160 blocks on 132 SMs) and stalls at ~3% weight-DRAM. - matmul_nbits_marlin_gemm_f16_splitk: partitions the K/group range over grid.z, each z accumulates its groups' frag*scale into fp32 and writes partials[split_k, M, N] (no bias). - matmul_nbits_marlin_splitk_reduce: sums the z-slices in FIXED order (deterministic => run-to-run identical => capture-stable) then applies fold_bias -> fp16. - choose_split_k: conservative heuristic; only splits when the base grid is under ~2 waves (so it NEVER regresses the already-occupied M>=32 / large-N cases), caps each split at >=2 groups and factor <= 8. - launch_marlin_gemm_splitk: static-grid gemm + reduce pair, capture-safe once partials is pre-allocated (scratch pool). Numerics: split-K reorders fp32 partials vs the single-block kernel, so NOT byte-identical to direct, but passes the f64-oracle parity guard (tol 2e-2) incl. asymmetric-zp, and is deterministic (new tests marlin_splitk_parity_vs_f64_oracle, marlin_splitk_is_deterministic). Measured (H200 GPU 3, median us, split-K vs direct): N=5120 M=1 : 90.0 -> 32.0 us 2.81x (weight 3.0% -> 8.5% peak) N=5120 M=2 : 90.2 -> 32.1 us 2.81x N=5120 M=8 : 105.5 -> 42.9 us 2.46x N=13824 M=1 : 210.3 -> 88.2 us 2.38x (split_k=2) N=13824 M=2 : 127.3 -> 81.3 us 1.57x (split_k=8) Correctly regresses (heuristic returns 1, stays on direct) at M>=32 and large-N M>=8, where base occupancy already suffices. Standalone/opt-in only; not yet wired into the op dispatch. Direct kernel remains the default (byte-identical greedy-token path unchanged). Refs #957, #960. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

…EFORE Merged Deckard's P1 snapshot c692bff (full M>1 Marlin, zero tiled fallbacks) and ran every gate MARLIN-OFF vs ON (ONNX_GENAI_MARLIN_M_GT_1=1) on H200 GPU7. HEADLINE (leverb-phase0 Part D INC0, glm): captured M=8/M=1 ratio B* 8.76x->4.99x (-43%); all 240 MatMulNBits capture seams eliminated; captured M=8 wall 87.3->50.7ms. segments 41->120 because Marlin exposed the residual GQA x40 + LayerNorm x80 seams (capture-safety lever, not Marlin scope). Still NO-GO at strict B*<=2, but much closer. Prefill sweep: native ~2x on both models (glm 218->426, qwen 154->300 tok/s), halving the glm-vs-ORT gap (121x->62x at L=1024). M-scaling tail slope collapses glm 2.21->0.93, qwen 3.26->1.10 ms/row — confirms Deckard's kernel 6.9x->~2x at the forward level. Portability: Marlin-OFF rows == portable tiled fallback == unchanged BEFORE baseline, proving <SM80 does not regress. Decision drop updated with full AFTER tables + commands. Refs #957, #962; docs/research/decode-remaining-levers-feasibility.md §3/§4 Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

Routes the plain and rmsnorm-prologue Marlin M>1 GEMMs through split-K when ONNX_GENAI_MARLIN_SPLITK=1 AND choose_split_k elects to split (base grid under ~2 waves). Partials come from the capture-safe scratch pool (slot 4, distinct from rmsnorm slot 0 / gate_up slot 1), so warm replays stay allocation-free; the reduce kernel applies the same fold_bias epilogue. Default OFF => the byte-identical direct kernel remains the greedy-token path. New helper Kernel::maybe_launch_marlin_splitk centralizes the direct-vs-split-K choice and propagates the pooled-scratch warm flag into the existing capture-safety return contract. gate_up SwiGLU stays on the direct kernel for now (its fused SiluMul epilogue would need a split-K reduce variant); documented as remaining work. Validated (H200 GPU 3): - f64-oracle parity + determinism (marlin_splitk_* lib tests) pass, incl. asymmetric-zp, tol 2e-2. - 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 itself is SM80-gated. Refs #957, #960. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

The last two MatMulNBits-adjacent capture barriers at query-width > 1: GroupQueryAttention (x40) and SkipSimplifiedLayerNormalization (x80) declared KernelCaptureUnsupported for any M>1 shape, keeping the speculative-verify (M=K) forward segmented even after Marlin removed all 240 MatMulNBits seams. Both were conservative gates, not real hazards. SkipSimplifiedLayerNormalization: - The launch is already a static grid (groups_u,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 (incl. num_groups) mid-capture — the same pre-warm cold-miss safety valve the Marlin repack cache uses. The `num_groups == 1` capture latch was purely conservative. Relax it to advertise capture support for a fixed, pre-warmed M=K shape (bf16 staging path still demotes on arena growth). GroupQueryAttention: - Broaden the capture signature to admit a fp16/bf16 q_seq>1 fused-flash verify/prefill shape, not just q_seq==1 Phase2a decode. The fused flash kernel is already capture-safe by construction (static grid, 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 the host-side seqlens/position validation, and the capture-safe branch uses present_capacity (masking the 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. 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 is still declared only for a warmed signature (pre-warm the M=K shape before capture, same safety valve as the Marlin repack/scratch pools). This is Stage 2 of the merged spec-capture design (#957): "SkipSimplified LayerNorm = signature relaxation, GQA = pre-warmed fixed-M flash workspace." With Marlin M>1 + these two, no hot decode/verify node declares KernelCaptureUnsupported at M>1. Refs #957, #960. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

…afe) Merged Deckard's c842b75 (GroupQueryAttention + SkipSimplifiedLayerNorm M>1 capture support). leverb-phase0 Part D INC0 (glm, M=8, pre-warmed): - segments 120 -> 1: whole M=8 verify now captures as ONE replayable graph. Only residual seam = MatMul[KernelCaptureUnsupported]x1 (lone dense MatMul). - B* unchanged by capture-safety alone (5.10x vs marlin-only 4.99x): the captured M=8 wall is compute-bound, not segmentation-bound (120->1 segments moved the wall ~0: 50.7->51.8ms). - Split-K (exploratory, ALSO capture-safe + byte-identical) is the lever that moves B*: 5.10x -> 2.69x (M=8 wall 51.8->27.3ms) -- right at the <=2 GO line. - Numerics: Marlin greedy tokens byte-identical to tiled with AND without split-K (24/24), no NaN under capture. Zero-init KV tail assumption holds. Path to speculative-capture GO = capture-safety (done, segments=1) + split-K on the M=K verify. Decision drop updated with progression table + commands. Refs #957, #960, #962 Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

…confirm Refreshed M-scaling + prefill sweep at c842b75 with split-K ON, and re-ran the decisive capture probe clean for reproducibility: - Capture re-probe reconfirmed: canonical segments=1 / MatMul x1 seam / B*=5.09x; split-K segments=1 / B*=2.71x. Reproduces prior 5.10/2.69 within noise. - Warmup coverage confirmed: Part-D INC0 warms via a FULL M=K forward, so every fixed-shape cache (Marlin repack+scratch, GQA meta, LN meta) is populated in one pass; capture_alloc=(0,0) + segments=1 prove no cold-miss valve fired. - Final tail slope (ms/row) BEFORE -> Marlin -> +split-K: glm 2.21 -> 0.93 -> 0.822 (-63%); qwen 3.26 -> 1.10 -> 1.147 (-65%). - Prefill nuance: canonical Marlin edges split-K at large-M prefill (qwen 291-300 vs 267-277); split-K's win is small-M M=K verify. Recommend canonical Marlin for prefill, split-K auto-on for the speculative verify. Refs #957, #960, #962 Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

…B* 2.63x Merged Deckard's 2971403/1a72c865 (lm_head cached dense-GEMM capture plan + split-K default within MARLIN_M_GT_1=1). leverb-phase0 Part D INC0, glm, M=8, GPU7 (Deckard on GPU5): - Census now TRUE ZERO unsupported nodes: seam nodes '<none: whole-graph>'. The whole M=8 verify (incl. lm_head logits projection) captures as ONE whole-subgraph replay. Last MatMul[KernelCaptureUnsupported]x1 is gone. - B* 2.71 -> 2.63x (captured M=8 wall 27.5 -> 26.8ms): folding lm_head out of its per-step eager launch+sync nudged B* down as predicted, but strict <=2 GO line not yet crossed. Residual is genuine small-M GEMM compute. - Byte-identical greedy tokens on glm-4-9b-int4 with the new default path (split-K default + lm_head plan); no NaN. Cleanest capture config to date. Refs #957, #960, #962 Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

Merged Deckard's 4abe4e5/0143f235 (choose_split_k retune for M=8 verify: gate_up/down/o-q now split sk=8; prefill still sk=1 direct kernel). Re-probed glm M=8, GPU7 (Deckard on 4/6), reproduced x2: - B* 2.63 -> 2.16x / 2.17x (captured M=8 verify wall 26.8 -> 22.0ms). Gate_up MLP GEMM off the no-split slow path is the dominant win. - Still segments=1, whole-graph, ZERO unsupported nodes. Byte-identical greedy tokens on glm-4-9b-int4 (parity PASS), no NaN. - 2.16x is an honest plateau a hair above the conservative <=2 line; but <=2 is break-even-at-2-accepted-tokens, so this is a PRACTICAL GO once the draft accepts >2.16 tok/verify (trivial for an 8-wide draft). Full arc: segments 41->1; B* 8.76->4.99->5.09->2.71->2.63->2.16x, byte-identical. Refs #957, #960, #962 Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

…plateau) Merged Deckard's 3735d57/dffabf0d (gate_up SwiGLU MLP GEMM was DIRECT-launching, bypassing split-K even after the 4abe4e5 factor retune; now routes sk=8). Re-probed glm M=8, GPU7, reproduced x2: - B* 2.15x / 2.19x; captured M=8 verify wall stable ~21.9ms (ratio jitter is from the M=1 captured baseline 9.98-10.18ms, not the M=8 wall). segments=1, whole-graph, zero unsupported nodes, byte-identical (parity PASS), no NaN. - This is the small-M GEMM floor: M=8 mma.m16n8k16 is ~half-occupancy (~50% MMA waste), so split-K+occupancy is the last GEMM lever; path to universal <=2 is drafting depth, not more GEMM tuning. - Practical GO: any 8-wide draft accepting >2.16 tok/verify wins. Capture-safety fully solved. Full arc: segments 41->1 whole-graph; B* 8.76->4.99->5.09->2.71->2.63->2.16x. Refs #957, #960, #962 Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

…wen block_32 B*~4.7x Key finding: the fused gate_up SwiGLU Marlin path requires block_size==32 (matmul_nbits.rs run_f16_gate_up_swiglu). Verified via onnx: glm=block_128, qwen=block_32. - glm (block_128) is NOT eligible for the gate_up fusion -> 3735d57 is a NO-OP for glm; glm's B*=2.16x comes from the 4abe4e5 general split-K retune, not gate_up fusion. Explains the flat glm wall (22.0->21.9ms) after 3735d57. - qwen (block_32) M=8 verify re-probe: segments=1, whole-graph, zero unsupported (48 GQA + 96 LN captured -> capture-safety generalizes), but B*=4.62-4.72x (captured M=8 35.5ms / M=1 7.5ms). NOT at GO. - Honest read: captured_ratio > eager_ratio (4.7>3.8) is a DENOMINATOR effect (capture speeds the launch-bound M=1 by 34% vs M=8 by 18%), not a capture inefficiency. qwen's tuned block-32 M=1 is fast (inflates ratio) and its 14B/48-layer M=8 verify is genuinely heavier. Deckard caveat #1 (qwen attn proj cold-clock bimodal) may add. Small-M GEMM floor is model-dependent: glm practical GO (2.16x), qwen 4.7x needs more small-M work or deeper drafting. glm remains canonical gate. Refs #957, #960, #962 Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

…OT a Marlin regression Ran qwen byte-identical parity at 3735d57; first run failed (Marlin greedy stream diverged from tiled at token 19). Investigated before raising a regression: - 4 runs (split-K on x3 + off x1): Marlin output DETERMINISTIC every time (5692,374,264,1140,315). The TILED reference diverged (320,94409,...) on only 1 of 4 runs. Token 19 is a near-degenerate argmax; the tiled GEMM's own nondeterministic (atomic/reduction-order) accumulation occasionally flips it. - Conclusion: NOT a Marlin numerics regression. Marlin (fixed split-K reduction) is the MORE deterministic path; the tiled oracle flakes ~1-in-4 at this tie. Split-K on and off both produce identical deterministic Marlin tokens. - Action for Chew/Gaff: the qwen exact-token parity test is flaky (~25%) because it asserts equality vs a nondeterministic tiled reference at a near-tie. Not a blocker; harden the assertion (logit tolerance / deterministic oracle / assert Marlin self-determinism). glm's prompt doesn't hit such a tie -> glm clean. Refs #957, #960, #962 Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

…960) ## Marlin-style fused fp16xint4 tensor-core GEMM (WIP -- draft) Builds the primary decode-perf lever from `docs/research/decode-remaining-levers-feasibility.md` and the capture enabler from the spec-capture feasibility doc (#957): a Marlin-style fused fp16xint4 tensor-core GEMM to replace the slow M>1 path of `com.microsoft::MatMulNBits` (today a portable 16x16 CUDA-core tiled GEMM) and, conditionally, the M=1 decode GEMV. Refs #957. ### What's landed (Stage 1-4) - New module `crates/onnx-runtime-ep-cuda/src/kernels/marlin_gemm.rs` (kernel + host/device repack isolated from the 475 KB `matmul_nbits.rs`). - From-scratch SM80 `mma.sync.m16n8k16` kernel using raw inline PTX (no `<mma.h>`), adapting Marlin ideas (repacked per-lane weights, per-group scale AFTER the tensor-core accumulate) to our ONNX-native asymmetric-zp / nibble format, group sizes 16/32/64/128. - `repack_int4_weights` + on-device `matmul_nbits_marlin_repack`: bijective 8-column-interleaved reorder for coalesced 64-byte/warp loads. Device repack is byte-identical to the host reference. - `device_supports_marlin()` SM80 guard for the Rule 11 byte-identical fallback. Static launch grid => capture-safe by construction. - Wired into the `MatMulNBits` M>1 dispatch (opt-in `ONNX_GENAI_MARLIN_M_GT_1`, default OFF, tier-scoped). Module-level repack cache repacks immutable-initializer weights once during warmup; warm replays do no allocation and are capture-safe. Cold-miss-during-capture is rejected. Any ineligibility / launch error falls through byte-compatibly to the portable tiled GEMM. - `capture_support` advertises M>1 support on a warm cache -- the property that unlocks capture-stable speculative verification at M>1. ### Validation - Numerics: `marlin_parity_vs_f64_oracle` (standalone) and `marlin_m_gt_1_op_parity_and_capture_safety` (through the real op) pass vs an f64 dequant->GEMM oracle across M in {1,2,7,16,33,64}, group in {16,32,64,128}, fp16/fp32 scales, sym+asym zero-points (abs <= 2e-2*max|out|; op worst_abs=0.013 << tol=0.67). The op test asserts cold=NOT-capture-safe / warm=capture-safe + byte-identical replay. - Wall vs tiled (through the op), `marlin_m_gt_1_op_wall_vs_tiled`, H200, K=5120 N=13824 block=128 (median us): | M | tiled us | marlin us | speedup | |---|---|---|---| | 1 | 69.5 | 69.0 | 1.01x (M=1 stays on GEMV) | | 2 | 478.0 | 140.0 | 3.41x | | 4 | 494.1 | 138.5 | 3.57x | | 8 | 536.4 | 176.9 | 3.03x | | 16 | 658.1 | 251.2 | 2.62x | | 32 | 1076.7 | 385.0 | 2.80x | | 64 | 2010.6 | 821.6 | 2.45x | | 128 | 3851.5 | 1357.9 | 2.84x | The M=1->M=2 cliff collapses from 6.9x (69.5->478 us tiled) to ~2x (69.5->140 us). Marlin is 2.4-3.6x faster than the tiled GEMM across every M>1. ### Honest status - M>1: correct, capture-safe (warm), 2.4-3.6x faster than tiled -- landed opt-in. - M=1: through-op M=1 Marlin ties the GEMV (69 vs 69.5 us), ~3.1% HBM peak -- does NOT clear the feasibility section-3 >=40% precondition, so M=1 stays on the existing GEMV (reported, not silently dropped). - Next: fused epilogues (rmsnorm-prologue -> swiglu/gate_up/down); small-M perf levers (split-K / cp.async); e2e on glm-4-9b-int4 / qwen2.5-14b (with Sebastian). Draft -- do not merge. Numerics gated by Chew; e2e/capture re-probe by Sebastian. ### Update — fused RMS-norm epilogue + e2e PASS - **Fused RMS-norm prologue for Marlin M>1** landed (first fused epilogue). Stages the normalized activation into scratch (byte-identical to the tiled prologue), then runs Marlin. Test `marlin_m_gt_1_rmsnorm_op_parity`: worst_abs=0.0156 << tol=0.90. - **Stage-5 e2e** (`tests/marlin_m_gt_1_e2e.rs`, native CUDA): tiled (flag off) vs Marlin (flag on) greedy token diff on the real models. Both **byte-identical**: - glm-4-9b-int4: 24 greedy tokens match. - qwen2.5-14b-instruct-int4 (asymmetric zp): 24 greedy tokens match. - Remaining: swiglu/gate_up/down fused epilogues; small-M perf levers (split-K / cp.async). Numerics gated by Chew; ORT e2e + capture re-probe by Sebastian. ### Update — P1 COMPLETE: full M>1 coverage (gate_up SwiGLU fused), zero tiled fallbacks Every hot MatMulNBits node in glm-4-9b and qwen2.5-14b now runs through Marlin at M>1. - **gate_up SwiGLU MLP fusion at M>1** (`try_launch_marlin_gate_up_prefill`): paired gate/up Marlin int4 GEMMs + the identical SiluMul epilogue the tiled path uses, with an optional RMS-norm prologue. This is the MLP — the bulk of prefill/verify cost. `down_proj` is a plain tall-skinny MatMulNBits already covered by `gemm_marlin_int4`. - **Capture-safe scratch pool** (`ensure_scratch`, size+slot-keyed, rejects cold-miss during capture) so warm replays of gate_up + rmsnorm allocate nothing. The rmsnorm prefill was migrated onto it and is now capture-safe on warm replay too. - **Per-node coverage audit** (tracer-driven e2e test): asserts zero `gemm_f16_tiled{,_rmsnorm}` / `gate_up_swiglu{,_rmsnorm}_prefill` at M>1. Measured coverage audit (H200, Marlin M>1 enabled): - glm-4-9b-int4: 240 `gemm_marlin_int4`, **0 tiled M>1**. - 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 tokens **byte-identical** tiled-vs-Marlin on both (24 tokens). gate_up unit parity (Marlin vs tiled): worst_abs 0.0001–0.008; warm replays capture-safe and byte-identical. No hot M>1 node remains on the tiled path — this unblocks the Increment-0 capture re-probe (segments→1). Next: P2 split-K / cp.async + honest M=1 precondition re-run. --- ## P2 — widen the win: split-K (landed, opt-in) + honest M=1 re-run **split-K** (`ONNX_GENAI_MARLIN_SPLITK=1`, default OFF): partitions the K/group range over `grid.z` to fill idle SMs when M is small (the occupancy-bound regime), fixed-order fp32 partial reduce ⇒ deterministic/capture-stable, pooled partials ⇒ capture-safe. Conservative `choose_split_k` only splits under ~2 waves so it never regresses the already-occupied M≥32 / large-N cases. Measured (H200, split-K vs direct): N=5120 M=1 **2.81×**, M=2 2.81×, M=8 2.46×; N=13824 M=1 2.38×, M=2 1.57×. Gates intact: f64-oracle parity + determinism tests pass; e2e greedy tokens **byte-identical** vs tiled with SPLITK=1 on glm-4-9b and qwen2.5-14b (zp); SM80 guard + fallback unchanged. **cp.async multistage: rejected** (measured slower — small-M is occupancy-bound, not latency-bound; hiding load latency only added barrier overhead). **M=1 DRAM precondition — honest re-run: still fails.** Best split-K lifts M=1 Marlin weight-DRAM to only ~8.5% peak, far below the ≥40% floor ⇒ **M=1 stays on the existing GEMV**. Remaining: gate_up SwiGLU M>1 stays on the direct kernel (split-K reduce for the fused SiluMul epilogue is follow-up). --- ## Stage 2 of spec-capture (#957): last two M>1 capture barriers removed With Marlin clearing all 240 MatMulNBits M>1 seams, the only ops still declaring `KernelCaptureUnsupported` at query-width>1 were **GroupQueryAttention (×40)** and **SkipSimplifiedLayerNormalization (×80)**. Both were conservative gates, not real hazards; now both advertise capture support for a warmed M=K signature. - **SkipSimplifiedLayerNorm** (signature relaxation): static grid, shape-keyed metadata cache already the pre-warm valve; relaxed the `num_groups == 1` latch. - **GQA** (pre-warmed fixed-M flash): admit an fp16/bf16 `q_seq>1` fused-flash signature. `flash_attention::run` is already capture-safe (static grid, device `total_lengths`/`past_lengths`, no host read-back); on-device metadata + `capture_error_ptr` replace host validation; capture-safe branch uses `present_capacity` (device-masked tail) like the M=1 path. f32 M>1 stays eager. Gates: greedy tokens byte-identical (glm-4-9b-int4 M>1 prefill); GQA + norm unit tests pass; fmt + lib clippy + native-backend clippy clean. Capture still declared only for a warmed signature (pre-warm M=K before capture). Sebastian re-probes for the decisive segments→~1 / B*≤~2 (speculative-capture GO). --- ## Segment census confirmed (Sebastian) + last dense seam closed Sebastian's Increment-0 re-probe (glm-4-9b, M=8, `MARLIN_M_GT_1=1`) confirmed the GQA + SkipSimplifiedLayerNorm capture-safety: **segments 120 → 1**, all 40 GQA + 80 SkipSimplifiedLayerNorm + 240 MatMulNBits M>1 seams gone, byte-identical, no NaN under capture. Two follow-ups (commit `29714037`): - **Split-K default-ON** within the opt-in Marlin M>1 path (opt out `ONNX_GENAI_MARLIN_SPLITK=0`). It is the measured B\* lever — Sebastian's probe: B\* **5.10× → 2.69×** at M=8 (wall 51.8→27.3ms). `choose_split_k` keeps large-M prefill on the byte-identical direct kernel; greedy tokens byte-identical. - **Last M>1 capture seam closed:** the lone residual node was `lm_head/MatMul` → `logits` (plain fp16 dense `MatMul`, not MatMulNBits), which fell to cuBLASLt's per-call heuristic at M>1. Added a cached-plan fast path (`blas::CaptureGemmPlan` + `DenseGemmPlan`): algorithm + persistent workspace selected once at warmup, replayed with no heuristic/alloc/sync — the M>1 analogue of `F32GemvPlan`, bit-for-bit identical to `governed_gemm` at a fixed shape. Batched/broadcast GEMMs keep the per-call path. Gates: greedy tokens byte-identical on glm-4-9b-int4 and qwen2.5-14b-int4-zp; fmt + lib clippy + native-backend clippy clean; SM80 guard + fallback unchanged. --- ## Small-M split-K retune for the M=8 verify shape Profiled glm/qwen decode-GEMM dims at M=8 and found `choose_split_k` was under-splitting: it picked `split_k=2-3` (or `1` for large-N gate_up) where the optimum is `8`. Root cause — block count is a poor fill proxy at small M, where each block is latency-bound; deeper K-splitting raises weight-DRAM throughput. Retuned (split for `m<=32`, oversubscribe ~8 waves, floor toward the measured small-M optimum, cap 8). M=8 auto factor before→after: | GEMM | K,N | old | new | |---|---|---|---| | o/q proj | 4096,4096 | sk3 2.34× | **sk8 2.93×** | | down_proj | 13696,4096 | sk3 2.45× | **sk8 3.18×** | | gate_up (glm) | 4096,27392 | none | **sk8 1.11×** | | gate_up (qwen) | 5120,13824 | none | **sk8 1.35×** | The dominant MLP verify cost now splits; prefill (`m>32`) stays on the byte-identical direct kernel. Gates: split-K f64-oracle parity + determinism at sk=8; greedy tokens byte-identical on glm-4-9b-int4; fmt + lib clippy clean. ### gate_up (SwiGLU MLP) routed through split-K (3735d57) Found the residual B\* gap: `try_launch_marlin_gate_up_prefill` launched the two largest MLP GEMMs (gate/up, glm K=4096 N=13696 each ×40 layers) as **direct** `launch_marlin_gemm`, bypassing split-K. The update-7 retune therefore never reached gate_up. Both projections now go through `maybe_launch_marlin_splitk` (sk=8 at M=8); gate/up reuse the slot-4 partials scratch sequentially (correct under single-stream ordering), and both split-K warm flags fold into the returned capture warmth. Gates: `marlin_gate_up_swiglu_matches_tiled_*` PASS at M=8 (split-K now active; warm replay byte-identical, cold-miss valve intact); glm-4-9b-int4 **and** qwen2.5-14b-int4-zp e2e greedy tokens byte-identical; fmt + lib clippy + native-backend clippy clean. --- ## FINAL STATUS — frozen for merge review (supersedes the "Next"/"Draft" notes above) Capture-safety is **complete** and the M>1 Marlin path is **landed opt-in**. This branch is frozen for numerics (Chew) + quality (Gaff) review. **Frozen head:** `a11facbc` (code frozen at `3735d57e`; later commits are decision-drops only). Sebastian's decisive perf A/B was re-probed **on the frozen code commit `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 cold-miss. gate_up split-K routing is thus both the most complete config and the independently-measured one. **Capture / perf arc (glm-4-9b, M=8 speculative verify, Sebastian's Increment-0 probe):** - Segment census: **41 → 1** — the whole M=8 verify (all 240 MatMulNBits + 40 GQA + 80 SkipSimplifiedLayerNorm + the lone dense lm_head MatMul) captures as **one whole-graph replay, zero `KernelCaptureUnsupported` nodes**. - Capture **B\* 8.76 → 4.99 (Marlin) → 2.71 (split-K) → 2.63 (lm_head dense-plan) → 2.16× (small-M split-K retune)** — an honest plateau at the ≤2 practical-GO line (residual is intrinsic M=8 `mma.m16n8k16` half-occupancy, not a capture blocker). Prefill ~2× (glm 218→426, qwen 154→300). Marlin-OFF == BEFORE (no <SM80 / consumer-tier regression). **Per-model note (block size).** glm-4-9b is **block-128**: its gate/up run as separate MatMulNBits through the general Marlin split-K path (covered by the retune), so the fused gate_up split-K commit (`3735d57e`) is a **no-op for glm** and its 2.16× is the retune alone. The fused gate_up SwiGLU path requires `block_size==32`, so `3735d57e` benefits **block-32** models (qwen2.5-14b), where it is confirmed firing under capture (qwen census: segments=1, zero unsupported — a fall-through would hit the non-capture-safe tiled path and fragment). qwen M=8 verify lands at **B\*≈4.7×**: capture-safety fully generalizes (whole-graph, zero seams, byte-identical eager parity), but the ratio is inflated by a denominator effect (capture speeds qwen's fast block-32 M=1 GEMV 34% vs the heavier 14B M=8 verify 18%). Closing qwen's B\* is a drafting-depth story, not a kernel fix. **glm-4-9b (block-128) is the canonical practical GO at 2.16×.** **Env-var contract (both implemented — verified in `marlin_gemm.rs`):** - `ONNX_GENAI_MARLIN_M_GT_1` — opt-in, **default OFF**, tier-scoped (Rule 11). `marlin_m_gt_1_enabled()` (`1`/`true`/`on`). SM80 runtime guard (`device_supports_marlin`) + byte-identical fallback to the portable tiled GEMM on <SM80 / CPU EP. Two weight layouts coexist; repack is an added packaging step. - `ONNX_GENAI_MARLIN_SPLITK` — **default ON** within the opt-in M>1 path; opt out `=0`/`false`/`off` for the byte-identical direct kernel. `choose_split_k` only elects a split for small-M / low-wave shapes (returns 1 for large-M prefill), so prefill/decode stay on the byte-identical direct kernel. **Coverage:** zero M>1 tiled-fallback for the hot MatMulNBits nodes on glm-4-9b + qwen2.5-14b (plain GEMM, rmsnorm-prologue, swiglu/gate_up, down_proj all fused through Marlin). GQA + SkipSimplifiedLayerNorm made M>1 capture-safe; lone dense lm_head MatMul made M>1 capture-safe via a cached cuBLASLt plan (`CaptureGemmPlan` / `DenseGemmPlan`). **Gates at frozen head (H200 GPU6):** f64-oracle split-K parity PASS; `marlin_gate_up_swiglu_matches_tiled_{plain,rmsnorm,decomposed}` PASS (split-K active at M=8, warm replay byte-identical, cold-miss valve intact); glm-4-9b-int4 **and** qwen2.5-14b-int4-zp e2e greedy tokens **byte-identical** to the tiled path; `cargo fmt --all -- --check` clean; `cargo clippy` lib + native-backend (warnings-as-errors) clean. --------- Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

…961) ## Numerics validation gate for Marlin int4 GEMM Builds the reusable **f64 dequant→GEMM oracle + parity gate** that Chew (numerics reviewer) will use to sign off Deckard's Marlin int4 GEMM (`squad/marlin-kernel`). Refs #957 (Lever A — Marlin int4 relayout, the primary decode lever). Marlin's weight relayout reorders partial sums, so its output is **not byte-exact** vs the current tiled kernel and cannot be validated by a bit-for-bit diff — the only defensible reference is a high-precision f64 oracle with a justified tolerance. ### What's here New self-contained integration test `crates/onnx-runtime-ep-cuda/tests/matmul_nbits_marlin_numerics.rs` (no edits to `matmul_nbits.rs` → no conflict with the Marlin kernel branch): - **f64 oracle:** dequantize packed int4 to f64 as `(code − zp)·scale` (scale pre-rounded to storage dtype) and accumulate `Σ a·w` in f64. Candidate and oracle share the same fp16-rounded activations + rounded scales, so the residual isolates only accumulation precision + fp16 output rounding. - **Coverage:** group sizes {16,32,64,128}; symmetric **and** asymmetric-zp; M ∈ {1,2,4,8,16,32}; realistic **glm-4-9b** (4096/4096, 4096/256, 4096/13696, 13696/4096) + **Qwen2.5-1.5B** (1536/8960, 8960/1536) attention+MLP projections; fp16 & fp32 scales. - **Marlin-ready:** drives the op purely via **op semantics** (ONNX `MatMulNBits` node + `get_kernel`/`execute`), never internal layout. Once Marlin is wired into the same dispatch, `run_matmul_nbits_f16` validates it with **zero changes**; or feed any candidate output slice into `Int4Problem::parity`. ### Justified tolerance envelope - **Absolute (primary):** `max(max_out·4e-3, 4e-3)` — ~8 fp16 ULP; covers a Marlin re-associated fp32 reduction (`~K·eps_f32 < 2e-3` at K≈13696). - **Relative (secondary):** `5e-2` against a **conditioning-aware** denominator floor `max(1e-1, 3e-2·max_out)`, so cancellation-dominated small-output columns stay absolute-governed. ### Current-path baseline (H200) | sweep | max_abs | max_rel | max_out | |---|---|---|---| | group-size (K=4096,N=896) | 6.50e-2 | 2.71e-3 | 1.67e2 | | projections (bs=32, M∈{1,8}) | 2.53e-1 | 8.14e-3 | 5.27e2 | Both ≈ **1 fp16 ULP** (abs ≈ `4e-4·max_out`) — ~8–10× inside the abs bound, ~6× inside the rel bound. **Marlin must land in the same envelope.** ### How to run ``` source .cudaenv.sh CUDA_VISIBLE_DEVICES=<idle-gpu> cargo test -p onnx-runtime-ep-cuda \ --features gpu-tests --test matmul_nbits_marlin_numerics -- --nocapture --test-threads=1 ``` 5 pure-CPU oracle self-consistency tests run without a GPU / feature; GPU baseline tests are `#[ignore]` unless `gpu-tests` is set (CPU-only CI stays green). `cargo fmt` + `clippy` clean. ⚠️ Draft — multi-turn; coordinator will send follow-ups. Numerics reviewer (Chew) sign-off pending; validates against Marlin (Deckard) once it lands. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

…957 Lever A) (#962) ## What New reproducible perf/capture harness **`marlin_bench`** (`crates/onnx-genai-bench/src/bin/marlin_bench.rs`, `--features bench-native,cuda`) that gates the full Marlin int4 GEMM build (**Lever A**, greenlit in #957 / `docs/research/decode-remaining-levers-feasibility.md`), plus the recorded **BEFORE baseline** on today's code so the moment Marlin lands the *same commands* produce a decisive delta. New file only — **zero overlap with Deckard's kernel branch** (`squad/marlin-kernel`). Uses only the public `NativeDecodeSession::load_with_resolved_io` API. ## The four gate axes (BEFORE baseline, H200 SM90) Full table, per-shape ncu rows, and exact reproduce commands are in the decision drop `.squad/decisions/inbox/sebastian-marlin-bench.md`. **Axis 1 — M=1 int4 GEMV DRAM%** (`ncu --graph-profiling node`): qwen MLP gate_up **~29.5%** (reproduces the doc §3 "~29% DRAM on an MLP shape" precondition), down_proj ~40%, attn ~24%; glm untuned `gemv_f16_general_bs` 11–17%. All far below the ≥55% Marlin M=1 pass line → GEMV is latency/dispatch-bound, as diagnosed. **Axis 2 — M-scaling wall (the cliff)** (`marlin_bench`, M∈{1,2,4,8,16}, min ms): glm past=32 **cliff 67.7 ms, M=8 91 ms, tail 1.65 ms/row** — reproduces Deckard's Increment-0 (~67 ms / ~87 ms / ~1.85). glm past=512 cliff 88.7 ms; qwen 120.9 ms. **Axis 3 — e2e tok/s native vs ORT** (`profile_native --steady`): qwen native **127.8** vs ORT **89.2 tok/s** (**1.43×**); glm native **98.6 tok/s** (matches the 96.4 doc reference). ORT can't load the native-authored glm ONNX (`GroupQueryAttention` `rotary_embedding_dim`) — foundry-local ORT-genai head-to-head is a follow-up. **Axis 4 — capture-safety re-probe** (existing `leverb_phase0` test, re-run): **segments=41**, captured M=8/M=1 = **8.77×**, NO-GO for Lever B today. Wired to re-run against Marlin (target: segments→1, ratio→~1). Its Part-C eager sweep independently reproduces Axis 2's cliff (67.2 ms) / tail (1.88 ms/row) — two harnesses agree. ## Measurement discipline / portability - Shared 8× H200 box: GPU pinned to a verified-idle high index, `nvidia-smi` re-checked before runs; harness reports median + p10/p90/max so contention is visible; **min is the contention-invariant estimator** used in the table. - **Rule 11:** every number is the current split-K/`general_bs` fallback layout (pre-Marlin), so re-running the identical harness proves <SM80 does not regress. `marlin_bench` prints device `compute_cap`. ## Validation `cargo fmt --all -- --check` clean · `cargo clippy` clean · `marlin_bench` + `profile_native` build under `bench-native,(bench-ort,)cuda`. Refs #957 · `docs/research/decode-remaining-levers-feasibility.md` ⚠️ Draft: this is the BEFORE baseline + harness. Multi-turn — AFTER numbers land once Marlin (Deckard) merges. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

Scribe consolidation for the Marlin int4 GEMM program landing (#960/#961/#962). - Merge 5 decision inbox drops (deckard/pris/sebastian kernel+bench+numerics + chew/gaff reviews) into one consolidated Marlin program entry; corrected per-model attribution (glm block-128 → gate_up fusion no-op; 2.16× was the general split-K retune). - Archive resolved #957/#948/#949 spec-capture + Lever B entries to decisions-archive/2026-08.md (decisions.md 45,991 → 42,104 B, under the 50 KB gate). - Orchestration logs (deckard/sebastian/pris/chew/gaff), session log, cross-agent history updates; summarized sebastian history (>8 chronicle threshold). State-only change. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

…wen 2×2 (#975) Follow-up to the Marlin int4 M>1 capture arc (kernel #960, bench #962, numerics gate #961, flaky-parity hardening #974). Refs #957 and `docs/research/decode-remaining-levers-feasibility.md`. ## What Deckard's final ask before Chew's #960 numerics review: an explicit **qwen captured-M=8 vs eager-M=8 token-equality** data point to fill the last cell of the captured/eager × glm/qwen 2×2. The shipping speculative path uses the *eager* verify path (`native_decode/backend.rs`), and the existing leverb INC0 probe writes a constant token (timing-only), so no existing harness answered this. This adds a **test-only** parity variant that feeds real tokens and compares captured-replay vs eager warm-forward per-row argmax **and** raw logits bytes. - `native_decode/cuda.rs`: `leverb_increment0_token_parity_attempt` (gated `#[cfg(all(test, feature="cuda"))]`). - `native_decode/leverb_phase0_probe.rs`: **Part E** prints per-row eager vs capture argmax + `logits_byte_identical` + PASS/FAIL. ## Result — full 2×2 GREEN, captured == eager **byte-identical** on both models | Model (block) | M | argmax_match | byte_identical | segments | B* | |---|---|---|---|---|---| | qwen2.5-14b (block-32) | 8 | ✅ | ✅ | 1 | 4.45× | | qwen2.5-14b | 1 | ✅ | ✅ | 1 | — | | glm-4-9b (block-128, canonical) | 8 | ✅ | ✅ | 1 | 2.19× | | glm-4-9b | 1 | ✅ | ✅ | 1 | — | B* refresh matches the frozen plateau (glm ≈2.16×, qwen ≈4.45×). No shipping path touched; fmt + engine clippy clean. Reproduce commands in `.squad/decisions/inbox/sebastian-captured-eager-parity.md`. ⚠️ Bench/validation follow-up — a squad member (Chew/Deckard) should confirm before merge. Co-authored-by: Sebastian <sebastian@squad.local> Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

…e decode Add the CUDA-graph-captured verify path (WP4 of #957): a persistent padded [1, width] token/position/logits binding set and a warm->arm->capture->replay state machine (DecodeCudaState::verify_capture + decode_verify_captured), modelled on run_one_token and the Increment-0 capture probe. Graceful eager fallback on NotCapturable keeps correctness independent of capture. Session surface: set_retain_graph_on_rewind / is_cuda / reset_captured_verify; rewind(0) tears down the retained verify graph before freeing its buffers. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

justinchuby merged commit 2f0b62b into main Aug 14, 2026
12 of 13 checks passed

justinchuby deleted the squad/spec-capture-design branch August 14, 2026 18:00

justinchuby mentioned this pull request Aug 14, 2026

chore(scribe): consolidate spec-capture feasibility decision #958

Merged

This was referenced Aug 14, 2026

Marlin int4 tensor-core GEMM for MatMulNBits M>1 (capture-safe, WIP) #960

Merged

test(cuda): f64 numerics gate for MatMulNBits int4 (Marlin sign-off) #961

Merged

perf(bench): Marlin int4 GEMM BEFORE/AFTER harness + baseline (gates #957 Lever A) #962

Merged

justinchuby mentioned this pull request Aug 14, 2026

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

Merged

justinchuby mentioned this pull request Aug 14, 2026

bench(capture): captured-vs-eager M=8 token parity — closes the glm/qwen 2×2 #975

Merged

Navigation Menu

Repository navigation

docs(research): speculative-capture feasibility — CONDITIONAL-GO gated behind Marlin - #957

docs(research): speculative-capture feasibility — CONDITIONAL-GO gated behind Marlin#957
justinchuby merged 1 commit into
mainfrom
squad/spec-capture-design

justinchuby commented Aug 14, 2026

Uh oh!

Uh oh!

codecov Bot commented Aug 14, 2026 •

edited

Loading

Uh oh!

github-actions Bot commented Aug 14, 2026

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

1 participant

Conversation

justinchuby commented Aug 14, 2026

Speculative decoding under CUDA-graph capture — feasibility (DESIGN-ONLY)

Verdict: CONDITIONAL-GO, and the single condition is Marlin (Lever A) — which we are already building as the primary decode lever.

Staged plan with hard gates

Context

Uh oh!

Uh oh!

codecov Bot commented Aug 14, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Codecov Report

Uh oh!

github-actions Bot commented Aug 14, 2026

🔴 Benchmark Regression Detected

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

1 participant

codecov Bot commented Aug 14, 2026 •

edited

Loading