[justinchuby]

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.

---

Test hardening (4803b4fc, test-only — kernel frozen at 3735d57e)

marlin_m_gt_1_matches_tiled_on_qwen2_5_14b_int4 was ~25% flaky: it asserted
exact greedy-token equality against the portable tiled GEMM, whose fp32
atomic reduction order is nondeterministic. At a near-degenerate argmax the
tiled reference flips run-to-run (Marlin is the deterministic side); greedy
decode is autoregressive so one early flip cascades. Not a Marlin bug — a test
oracle bug (glm's prompt has no such tie, so its test was already clean).

Replaced exact-equality with a cascade-robust classifier: compare one tiled vs
one Marlin stream; if they first diverge at position d (identical prefix), a
re-run tiled probe that reproduces the prefix but flips at d proves d is a
near-tie → not a regression. Extracted into a pure GPU-free classify_parity()
with deterministic unit tests for every outcome, so the branch is CI-verified
without needing to reproduce a live tie. Gates: unit tests + glm/qwen e2e pass,
fmt + clippy (lib/native-backend/cuda) clean.

[codecov]

Codecov Report

✅ All modified and coverable lines are covered by tests.
✅ Project coverage is 79.00%. Comparing base (94c69fd) to head (41e39a6).
⚠️ Report is 8 commits behind head on main.

Additional details and impacted files

@@            Coverage Diff             @@
##             main     #960      +/-   ##
==========================================
- Coverage   79.01%   79.00%   -0.01%     
==========================================
  Files         356      356              
  Lines      150096   150158      +62     
  Branches   150096   150158      +62     
==========================================
+ Hits       118596   118635      +39     
- Misses      26849    26871      +22     
- Partials     4651     4652       +1     

Flag	Coverage Δ
cli-ort-linux	83.52% <ø> (ø)
cli-ort-windows	83.11% <ø> (+0.09%)	⬆️
mlas	80.92% <ø> (-0.05%)	⬇️
offline	78.80% <ø> (-0.01%)	⬇️

Flags with carried forward coverage won't be shown. Click here to find out more.
see 5 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]

🔴 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
🔴	matmul/medium_generic_bf16_threads=8/32x512x512	371.45 µs	699.07 µs	+88.2%
🔴	matmul/medium_generic_f32_threads=8/32x512x512	962.18 µs	1.53 ms	+59.4%
🔴	matmul/large_generic_f32_threads=1/32x1024x1024	8.93 ms	14.15 ms	+58.4%
🔴	block_quantized_matmul_cached_dense/mxfp4_preexpanded_dense_oncelock_like_proxy/1x1024x1024	52.59 µs	68.50 µs	+30.3%
⚠️	matmul/small_generic_bf16_threads=8/1x256x256	29.55 µs	36.40 µs	+23.2%
⚠️	gather/large_f16_threads=1-internal/131072	10.85 µs	13.12 µs	+21.0%
⚠️	matmul/large_generic_f16_threads=8/32x1024x1024	86.61 µs	104.68 µs	+20.9%
⚠️	matmul/large_generic_f32_threads=8/32x1024x1024	3.94 ms	4.72 ms	+19.7%
⚠️	grammar_masking/llguidance_compute_mask/32	69.33 µs	82.73 µs	+19.3%
⚠️	matmul/medium_generic_bf16_threads=1/32x512x512	512.80 µs	605.87 µs	+18.1%
⚠️	matmul/small_generic_bf16_threads=1/1x256x256	29.14 µs	34.05 µs	+16.9%
⚠️	sampling_latency/min_p_per_token	191.57 µs	223.04 µs	+16.4%
⚠️	matmul/medium_generic_f16_threads=1/32x512x512	28.41 µs	32.93 µs	+15.9%
⚠️	block_quantized_matmul_cached_dense/mxfp4_cached_dense_repeated_call/1x1024x1024	47.27 µs	54.39 µs	+15.1%
✅	sampling_latency/greedy_per_token	2.97 µs	3.33 µs	+12.2%
✅	matmul/medium_generic_f32_threads=1/32x512x512	2.31 ms	2.59 ms	+12.2%
✅	tokenization/encode_tokens_per_second	351.94 µs	389.97 µs	+10.8%
✅	logit_processing/seven_processor_chain_per_step	299.68 µs	331.22 µs	+10.5%
✅	sampling_latency/top_p_per_token	354.53 µs	389.57 µs	+9.9%
✅	tokenization/decode_tokens_per_second	5.85 ms	6.41 ms	+9.6%
✅	qwen3_sampling_processors/top_k_partial_selection	129.02 µs	141.00 µs	+9.3%
✅	gather/small_f32_threads=1-internal/4096	631.1 ns	686.8 ns	+8.8%
✅	matmul/medium_generic_f16_threads=8/32x512x512	28.22 µs	30.64 µs	+8.6%
✅	matmul/large_generic_f16_threads=1/32x1024x1024	81.98 µs	88.98 µs	+8.5%
✅	qwen3_sampling_processors/top_k_top_p_fast	626.52 µs	675.64 µs	+7.8%
✅	qwen3_sampling_processors/top_k_full_sort_baseline	2.01 ms	2.16 ms	+7.5%
✅	qwen3_sampling_processors/top_p_full_sort_after_top_k_baseline	3.37 ms	3.61 ms	+7.4%
✅	kv_cache/alloc_dealloc_pages	36.10 µs	38.47 µs	+6.6%
✅	qwen3_sampling_processors/top_k_top_p_full_sort_baseline	5.59 ms	5.89 ms	+5.4%
✅	matmul/large_generic_bf16_threads=1/32x1024x1024	2.00 ms	2.09 ms	+4.9%
✅	sampling_latency/top_k_per_token	52.19 µs	54.71 µs	+4.8%
✅	gather/small_f16_threads=1-internal/4096	462.2 ns	479.4 ns	+3.7%
✅	matmul/large_generic_bf16_threads=8/32x1024x1024	1.29 ms	1.33 ms	+3.0%
✅	matmul/small_generic_f16_threads=1/1x256x256	28.56 µs	29.26 µs	+2.5%
✅	block_quantized_matmul_cached_dense/mxfp4_uncached_dequant_each_call/1x1024x1024	532.41 µs	536.34 µs	+0.7%
✅	gather/small_bf16_threads=1-internal/4096	446.0 ns	446.5 ns	+0.1%
✅	reduce_mean/large_f32_threads=1-internal/262144	949.97 µs	946.45 µs	-0.4%
✅	gather/medium_f16_threads=1-internal/32768	2.24 µs	2.21 µs	-1.5%
✅	matmul/small_generic_f32_threads=1/1x256x256	35.17 µs	34.43 µs	-2.1%
✅	matmul/small_generic_f16_threads=8/1x256x256	28.48 µs	27.79 µs	-2.4%
✅	reduce_mean/small_f32_threads=1-internal/4096	14.30 µs	13.94 µs	-2.6%
✅	gather/medium_bf16_threads=1-internal/32768	2.29 µs	2.21 µs	-3.6%
✅	qwen3_sampling_processors/top_p_fast_after_top_k	531.11 µs	511.13 µs	-3.8%
✅	matmul/small_generic_f32_threads=8/1x256x256	33.86 µs	32.26 µs	-4.7%
✅	add/large_f16_threads=1-internal/4194304	40.54 ms	38.22 ms	-5.7%
✅	reduce_mean/medium_f32_threads=1-internal/65536	242.90 µs	225.11 µs	-7.3%
✅	gather/medium_f32_threads=1-internal/32768	3.62 µs	3.33 µs	-8.0%
✅	gather/large_f32_threads=1-internal/131072	26.29 µs	23.16 µs	-11.9%
✅	add/large_bf16_threads=1-internal/4194304	44.67 ms	39.17 ms	-12.3%
✅	add/small_bf16_threads=1-internal/1024	13.79 µs	11.98 µs	-13.1%
🟢	gather/large_bf16_threads=1-internal/131072	12.92 µs	10.97 µs	-15.1%
🟢	add/medium_f32_threads=1-internal/262144	2.74 ms	2.32 ms	-15.2%
🟢	add/medium_bf16_threads=1-internal/262144	3.06 ms	2.58 ms	-15.9%
🟢	block_quantized_moe_cached_dense/mxfp4_cached_dense_expert_repeated_call/rows=1,H=256,I=256,E=4,top_k=1	96.34 µs	77.70 µs	-19.3%
🟢	block_quantized_moe_cached_dense/mxfp4_uncached_expert_dequant_each_call/rows=1,H=256,I=256,E=4,top_k=1	564.74 µs	447.40 µs	-20.8%
🟢	add/large_f32_threads=1-internal/4194304	48.77 ms	37.55 ms	-23.0%
🟢	add/medium_f16_threads=1-internal/262144	3.23 ms	2.37 ms	-26.7%
🟢	add/small_f16_threads=1-internal/1024	16.88 µs	11.83 µs	-30.0%
🟢	add/small_f32_threads=1-internal/1024	319.3 ns	183.3 ns	-42.6%

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: { 4.63 3.80 6.56 }

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)