perf(cpu-ep): fuse the 8-bit dequant into the prefill GEMM pack step (17x kernel; every ORT prefill cell loss -> win) by justinchuby · Pull Request #1403 · justinchuby/onnx-genai

justinchuby · 2026-08-19T05:07:46Z

Stacked on #1356 (squad/roy-int4-prefill-gebp), which this generalizes. Review the last commit only.

The loss

On a native (non-mlas) build there was no borrowed route at all for bits == 8, m > 1. try_prefill_mlas_nt declines, and the kernel falls to:

let weight_kn = self.dequantize_weight(.., WeightLayout::Kn)?;
gemm(&activations, &weight_kn, result, m, self.k, self.n)?;

Kn is the transposed layout the dense GEMM wants, so the dequant writes four bytes of f32 for every one byte of weight read, at stride n, into a buffer far larger than any cache — and does it on every call, because nothing caches Kn (only Nk is cached, and only on the MLAS path). A 3584x3584 8-bit node materialized 51 MB per prefill to multiply it once. That term does not shrink as rows are removed, which is why the loss was worst where it was least expected: m = 8 spent 41.7 ms to do 1.8 GFLOP.

The fix

Generalize #1117's fused-dequant GEBP to both bit widths. pack_b_quant expands each KC x NR panel straight into the L1-resident f32 panel the 6x16 microkernel already consumes; all m rows reuse it. The bit width is a BlockQuantWeight implementation (Int4Weight / Int8Weight), so 4-bit and 8-bit share packing, blocking, threading and microkernel and differ only in dequant_column. Each packed byte is read once per call; no f32 weight is ever resident.

The numbers are bit-identical, not merely close

Both routes compute the same (q - zero_point) * scale in f32 and accumulate along k in strictly increasing order into one accumulator per output element, so blocking cannot reorder the reduction. The A/B bench compares the full output of both arms in one process and reports bitexact for every cell of every shape and row count measured. That is also why the tests need INT8_PREFILL_GEBP_TEST_CALLS: values alone cannot tell which route ran.

Result — kernel level

k = n = 3584, block 32, symmetric 8-bit; steady-state ms, median of 3 interleaved arm-pairs, each the median of 7 timed calls:

m	dequant + GEMM (previous)	fused GEBP	speedup	GFLOP/s
1 (control, same route both arms)	1.054	1.054	1.00x	24
2	15.465	0.905	17.1x	57
4	15.572	0.864	18.0x	119
8	15.245	1.096	13.9x	188
64	16.520	2.833	5.83x	580
256	21.426	8.263	2.59x	796
512	28.731	15.602	1.84x	843

4096x11008 at m = 256: 67.392 → 22.561 ms, 1023 GFLOP/s. 2048x2048 at m = 4: 11.610 → 0.448 ms (25.9x). The m = 1 rows read 1.00x / 1.00x / 0.99x across the three shapes — that is the noise gauge, and it says the prefill ratios are not host drift.

Result — against ONNX Runtime

At the assignment matrix's own geometry (k = n = 3584), one build, arms differing only by ONNX_GENAI_CPU_MM_INT8_GEBP, interleaved trial by trial (7 trials; 11 at the noisier t = 8 rows), parity=PASS on every trial:

M	threads	before	after
128	2	2.061 [1.974-2.291]	0.655 [0.594-0.664]
128	4	2.012 [1.818-2.210]	0.745 [0.533-0.842]
128	8	2.321 [2.115-2.465]	0.681 [0.537-0.877]
256	2	1.700 [1.655-1.709]	0.767 [0.745-0.792]
256	4	1.694 [1.638-1.706]	0.834 [0.736-0.990]
256	8	1.887 [1.747-2.435]	0.879 [0.752-1.110]
512	2	1.412 [1.387-1.423]	0.864 [0.850-0.870]
512	4	1.421 [1.354-1.848]	0.877 [0.861-0.996]
512	8	1.542 [1.482-1.845]	1.011 [0.886-1.131]

Every measured 8-bit prefill cell crosses from a loss to a win except M = 512 at 8 threads, which reaches parity — 1.011 over 11 trials with a range straddling 1.0. By this repo's own bar (>= 5% repeatable win at every thread count) that row is not closed, and it is recorded as parity, not as a win.

A --features mlas research build of the same tree reads 0.563 / 0.503 at M = 128 and 0.824 / 0.763 at M = 512 (2 / 8 threads). The pure-native path went from 2.5-3.7x behind that build to within 1.05-1.35x of it, with no MLAS and no resident f32 weight.

Controls

control	arm A	arm B	reading
8-bit `M = 1` (decode; route untouched) at t=2/4/8	0.154 / 0.207 / 0.244	0.153 / 0.181 / 0.219	overlapping — unchanged, still a 4-6x win
4-bit 512 rows, same binary, switch it ignores	1.309 [1.198-1.877]	1.394 [1.167-1.681]	overlapping over 9 trials — wash
4-bit 1 / 128 / 512 rows, parent-branch binary vs this one	3.823 / 1.914 / 1.305	3.796 / 1.903 / 1.305	the `BlockQuantWeight` generalization costs the 4-bit path nothing (native p50 within 0.5%)

The last one is the control that matters for the refactor: int4_prefill_gebp became generic, so the 4-bit path had to be re-measured against a build of the parent branch, not merely against the same binary with a switch it ignores.

An earlier 5-trial run of the 4-bit same-binary control read 1.240 vs 1.375 with non-overlapping ranges. That was small-sample host drift — at 9 trials the ranges overlap heavily and the after arm is faster in absolute terms (58.5 vs 63.4 ms). It is written into the benchmark doc because the disjoint-range version is exactly the kind of number that gets published by accident.

Row gate is derived, not inherited

INT8_PREFILL_GEBP_MIN_ROWS = 2. Unlike 4-bit — where the GEBP competes with a genuinely cheap row-serial borrowed kernel and only wins from m >= 4 — there is no competitor here: the replaced route pays its 51 MB whatever m is. The constant was temporarily lowered and m = 2, 3 measured on both arms before the gate moved; the fused route is 17.1x faster at m = 2 and the ratio only grows. m = 1 never reaches this branch (the 8-bit decode GEMV claims it, and its Nk weight is cached), so 2 is the smallest value the constant can express.

Because the gate no longer excludes any prefill this branch sees, ONNX_GENAI_CPU_MM_INT8_GEBP=0 is the only way back to the old route — hence matmulnbits_int8_gebp_kill_switch_restores_the_dequant_route, which asserts the switch actually switches instead of merely existing.

Scope

x86_64 + AVX2 + FMA only; other targets keep the previous path unchanged.
Tried after try_prefill_mlas_nt declines, so an mlas build keeps its cached-Nk + trans_b route and nothing measured there changes.
Declines on per-row g_idx; not gated on accuracy_level (the replaced route was full-f32 compute at every level, and so is this).
Explicit length checks on packed / scales / zero_points before the pack, because the pack indexes them directly — a short operand would be an OOB read, not a wrong answer.

Tests and gates

New: matmulnbits_int8_prefill_gebp_matches_reference_and_ran (symmetric + asymmetric, k not a multiple of block_size, n not a multiple of NR = 16, row counts either side of one MR = 6 A panel, counter asserted), matmulnbits_int8_gebp_kill_switch_restores_the_dequant_route, matmulnbits_int8_gebp_declines_decode_and_narrow_prefill.

Green locally: 1440 onnx-runtime-ep-cpu lib tests · cargo fmt --all · clippy --all-targets -D warnings on x86_64 and aarch64 · cargo check --target i686-unknown-linux-gnu (clean once the pre-existing #1391 literal is patched; the new code is x86_64-gated) · env-var honesty (109 documented) · workspace test packages.

Also

scripts/ort_ab/gen_gemm.py grew 8-bit cells and the square 3584 geometry the assignment matrix is quoted at — that matrix carried 8-bit rows no generator in the tree could reproduce.

Full record: docs/benchmarks/2026-08-19-int8-prefill-gebp.md.

One honest note on the matrix

CPU_MATMUL_ASSIGNMENT.md recorded M = 128 8-bit as a win (0.90 / 0.87 / 0.94) and M = 256 as 1.17 / 1.15 / 0.99. The paired before-arm here reads 2.06 / 2.01 / 2.32 and 1.70 / 1.69 / 1.89. The M = 512 rows do reproduce (1.41 → 1.41, 1.39 → 1.42), and the --features mlas build does not reproduce the old small-M rows either (0.563, not 0.90) — so "the old rows were an MLAS build" is not the explanation. Those rows come from a tree that cannot be reconstructed here; they are replaced with the paired re-measurement rather than defended, and the discrepancy is stated in both docs instead of being quietly overwritten.

…tractions (#1370) Scribe consolidation: merges 5 settled inbox notes into decisions.md, records 3 durable claim-integrity decisions (GLM 1.18x deficit, DeepSeek QMoE fair-A/B, fp32 decode floor), archives the 2026-08-18 batch (~38KB) to decisions-archive. Live ledger 43KB->14.9KB. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

… fp16-QMoE NO-GO (#1372) Deckard's fp16-QMoE scoping produced a decisive graph-config finding: native graph=0 ~65 -> graph=1 (production) ~137 tok/s (same binary, medians-of-5). Native's production CUDA-graph path (~137) is ~1.58x FASTER than ORT graphs-off (86.78) — the 'ORT 1.57x faster' number was native-eager-vs-ORT. Also: native QMoE is fp32-accumulate (not activation-locked); fp16 tensor-core QMoE = ~0 M=1 prize = NO-GO. Definitive both-graphs-on A/B pending (Sapper). Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

…ntime (#1363) ## The loss `gemm_nbits_qwen3_0p6b_qkv_t8` in a **pure-native** build (`bench-native`, no `mlas` feature — no MLAS at runtime, no ORT CPU fallback), native p50: | t=8 | t=16 | t=32 | | --- | --- | --- | | 1.9 ms | 1.6 ms | **37.7 ms** | §33.9 named this cell and recorded a 4.6× inversion. It is now 23×. Sixteen extra hardware threads make a 1.6 ms kernel take 37.7 ms. ## The site is not the one that looks like it The obvious suspect is `borrowed_affine_int4_matmul_prefill`, the `m > 1` native column fan-out — the twin of the MLAS tiling #1238 fixed. **Routing it changed nothing on any cell**, because it is dead in the default build: gated on `borrowed_int4_prefill_block_enabled()`, which reads `ONNX_GENAI_CPU_MM_INT4_PREFILL` and defaults to *false* (as does `ONNX_GENAI_CPU_MM_INT4_NBLK`). The route every int4 GEMM actually takes is `borrowed_affine_int4_matmul`, which has no `m > 1` fan-out at all — it loops over activation rows and calls **`parallel_output_rows` once per row**. An 8-token prefill is eight fork-joins; a 512-token prefill is 512. `parallel_output_rows` was a bare `result.par_chunks_mut(chunk)` on whichever Rayon pool happened to be installed. It is the single flat fan-out shared by eight call sites (`int4_matmul_m1`, `packed_nbits_output_row`, `borrowed_affine_int4_matmul`, `borrowed_affine_int4_matmul_nblock`, `int8_row`, `gemv_nk`, `gemv_nk_u8`, `gemv_nk_u8_i16`) — the §26 disease on the hottest path in the CPU EP, three phases after §35.6 declared that path clean. This PR corrects that claim in the document. ## Measuring anything at all on this host The reference host carried a load average of **17–35 on 32 logical CPUs** throughout, because several agents share it. A single cell can swing 5× between trials. So every number below comes from a **three-arm** invocation: `base`, `new`, and `null` — a byte-identical copy of the `base` binary. The null arm measures what the harness reports for a change that provably does not exist. t=32, 16 `gemm_nbits_*` cells, 5 trials × 7 runs, arms interleaved within each trial: | | geomean | worst cell | best cell | | ---- | ---- | ---- | ---- | | `null` (base against itself) | 1.00× | 0.72× | 1.20× | | `new` (this change) | **2.33×** | 0.65× | 16.6× | A per-cell reading inside 0.72×–1.20× is indistinguishable from nothing here. A second, quieter invocation of the same grid reproduced the geomean at **2.15×**. ## The change Route the flat fan-out through `task_runtime::for_each_range`, passing the existing `output_chunk_len` through unchanged as the **minimum grain**, so the partition can only get coarser, never finer, than the one Rayon used. The `numa-split` and SPMD branches are untouched. Three things gate it, and all three are load-bearing (each was falsified — flipping any one constant turns a test red on its own). ### 1. `HOT_FAN_OUT_MACS` + `HOT_FAN_OUT_CALLS` — the hot-prefill carve-out Unconditional routing regressed two cells: `llama3_8b_mlp` at t128 (0.93×) and t512 (0.86×). They run the *same* 56 Mi fan-out as `llama3_8b_mlp_t8`, which **improves 1.5×**. Identical per-call work, opposite result — so the discriminator is not the work, it is **how many fan-outs arrive back to back**. That is §33.4's park model read in reverse. A parked Rayon worker takes 67–226 µs to reach its first task, which is why the runtime wins wherever fan-outs are sparse. But one fan-out per row of a 512-token prefill never lets the pool sleep: the park cost is paid once and amortised to nothing, leaving only the runtime's SMT cap — 16 physical lanes against Rayon's 32. **No threshold on total operator work can express this:** | cell | fan-out MACs | calls | total | wants | | ---- | ---- | ---- | ---- | ---- | | `llama3_8b_mlp_t128` | 56 Mi | 128 | 7.2 Gi | Rayon | | `llama3_8b_qkv_t512` | 24 Mi | 512 | **12.3 Gi** | task runtime | The cell that wants Rayon has *less* total work. Reusing #1238's `WIDE_PREFILL_MACS` would have thrown away five measured wins to fix two losses. So the wide path is taken only when **both** hold: `fan_out_macs >= 32 Mi` (measured separation 24 Mi vs 56 Mi) and `calls >= 32` (measured separation 8 vs 128). `parallel_output_rows` keeps its one-shot spelling and delegates with `calls = 1`, so no decode step can reach the wide path however large its projection; only the two row-looping int4 kernels pass `m`. ### 2. `lanes` is a closure — a pool that should not have existed Routing was worth 2.2× at t=32 and 1.07× at t=16, but the same grid at **t=4 came back at 0.81×** with six of eight cells losing — outside the 0.95× the null arm shows at that width. The cause was not the routing. It was `task_runtime::width()`, which is documented to **build the pool** so a kernel can size a partition against it. The policy asked for the width, chose Rayon anyway, and left behind resident workers that spin before they park — on a configuration that has three other threads doing the arithmetic. `flat_fan_out` now takes `lanes` as `impl FnOnce() -> usize`: the width test that needs no pool runs first, and a fan-out that will stay on Rayon never asks. The type makes it impossible to reintroduce, and `flat_fan_out_does_not_ask_for_a_width_it_will_not_use` pins it. ### 3. `MIN_ROUTED_FAN_OUT_WIDTH` — the narrow half is left alone Below 16 Rayon threads there are too few wake-ups for their cost to be worth re-homing a fan-out for. The crossover is measured between 8 and 16: | width | geomean, routed unconditionally | null arm | | --- | --- | --- | | t=1 | 1.01× | — | | t=2 | 1.01× | — | | t=4 | **0.81×** | 0.95× | | t=8 | **0.94×** | 0.97× | | t=16 | 1.07× | — | | t=32 | **2.22×** | 1.00× | So the whole `t ≤ 8` half of the grid is left exactly as it was. t=1/t=2 are flat because the partition policy declines to split and both arms run identical serial code. ## The matrix after the fix Native-only p50 at t=32 — the only width at which the policy can choose: | cell | fan-out | calls | path | Rayon | new | speedup | | ---- | ---- | ---- | ---- | ---- | ---- | ---- | | `qwen3_0p6b_mlp_t1` | 6 Mi | 1 | runtime | 10.68 ms | 1.24 ms | **8.6×** | | `qwen3_0p6b_qkv_t1` | 3 Mi | 1 | runtime | 4.84 ms | 0.70 ms | **6.9×** | | `qwen3_0p6b_mlp_t8` | 6 Mi | 8 | runtime | 60.20 ms | 9.98 ms | **6.0×** | | `qwen3_0p6b_qkv_t8` | 3 Mi | 8 | runtime | 48.25 ms | 8.40 ms | **5.8×** | | `llama3_8b_qkv_t1` | 24 Mi | 1 | runtime | 14.01 ms | 2.45 ms | **5.7×** | | `llama3_8b_qkv_t8` | 24 Mi | 8 | runtime | 65.29 ms | 29.40 ms | 2.2× | | `llama3_8b_mlp_t1` | 56 Mi | 1 | runtime | 16.92 ms | 8.26 ms | 2.1× | | `llama3_8b_mlp_t8` | 56 Mi | 8 | runtime | 67.34 ms | 44.74 ms | 1.5× | | `qwen3_0p6b_qkv_t128` | 3 Mi | 128 | runtime | 80.11 ms | 53.13 ms | 1.5× | | `qwen3_0p6b_qkv_t512` | 3 Mi | 512 | runtime | 144.62 ms | 117.31 ms | 1.2× | | `llama3_8b_qkv_t128` | 24 Mi | 128 | runtime | 180.50 ms | 158.75 ms | 1.1× | | `qwen3_0p6b_mlp_t512` | 6 Mi | 512 | runtime | 183.28 ms | 167.18 ms | 1.1× | | `llama3_8b_qkv_t512` | 24 Mi | 512 | runtime | 528.29 ms | 490.51 ms | 1.1× | | `qwen3_0p6b_mlp_t128` | 6 Mi | 128 | runtime | 87.71 ms | 82.94 ms | 1.1× | | `llama3_8b_mlp_t512` | 56 Mi | 512 | **wide** | 899.74 ms | 891.01 ms | 1.0× | | `llama3_8b_mlp_t128` | 56 Mi | 128 | **wide** | 212.54 ms | 213.85 ms | 1.0× | The two diverted cells were then re-measured on their own at **9 trials, three arms**, to settle them past the noise: | cell | base | `null` | `new` | | ---- | ---- | ---- | ---- | | `llama3_8b_mlp_t128` | 270.1 ms | 262.5 ms (1.03×) | 264.0 ms (**1.02×**) | | `llama3_8b_mlp_t512` | 906.5 ms | 898.0 ms (1.01×) | 908.0 ms (**1.00×**) | `new` tracks `null` on both — the expected result for a policy that hands those cells back to the code the base binary runs. Both regressions are gone and no win was traded for them. The four dense f32 cells are the other control (unreachable by this change): 0.96×–0.99×. **Every cell reported `parity=PASS`.** Row-sharding a GEMV is exactly associative — no cross-row reduction — so results are bit-identical however the rows are partitioned. ## Tests Ten new tests. The policy ones run against a synthetic 16-lane/32-thread host, so they do not need SMT to be meaningful: - `flat_fan_out_sends_hot_repeated_fan_outs_to_the_wide_pool` — pins every row of the measured table. - `flat_fan_out_leaves_narrow_fan_outs_where_they_are` — the `t ≤ 8` half, plus the 16-thread crossover. - `flat_fan_out_does_not_ask_for_a_width_it_will_not_use` — a `Cell<bool>` proving the pool is not built on the declining path. - `flat_fan_out_ignores_a_wide_pool_that_is_not_wider`, `flat_fan_out_keeps_one_shot_fan_outs_on_the_runtime`, `parallel_output_rows_is_the_one_shot_spelling` — a decode step can never be diverted. - `parallel_output_rows_covers_every_output_exactly_once` and `..._repeated_covers_every_output_on_the_wide_path` — per-element `AtomicU32` write counters plus `output_start` correctness, so both executors owe the same coverage guarantee. - `parallel_output_rows_dispatches_to_the_task_runtime` — asserts a real dispatch, so the routed path cannot silently degrade to serial. - `prefill_column_fan_out_matches_its_serial_self` — bit-identity against `testing::force_serial()`; falsified during development (truncating one task's range turns it red at index 95) and it asserts `counters().tasks` moved, so it cannot pass by comparing serial against serial. All three policy constants were falsified before being trusted: `HOT_FAN_OUT_MACS` → 128 Mi, `HOT_FAN_OUT_CALLS` → 1024, and `MIN_ROUTED_FAN_OUT_WIDTH` → 8 each turn a test red on its own. ## Validation Local, rustc 1.97.1 (8bab26f4f 2026-07-14) — the toolchain `Rust quality` pins: - `cargo test -p onnx-runtime-ep-cpu` — **1446 passed, 0 failed, 17 ignored**, all integration binaries green - `cargo clippy -p onnx-runtime-ep-cpu --all-targets -- -D warnings` — clean - `cargo fmt --all -- --check` — clean - all 8 `Rust quality` guard scripts — pass Rebased on `main` at `6a855d5e0`. ## What is still lost `gemm_nbits_qwen3_0p6b_qkv_t8` is 8.4 ms at t=32 against 1.6 ms at t=16. The inversion is 23× smaller but not zero, and *more threads, more time* is still a scheduler signature. Leading hypothesis, untested: `with_decode_pool`'s `pool.install(operation)` is itself a fork-join onto a possibly-parked pool, paid once per `MatMulNBits` node call. Probing it needs care — skipping the install changes what `output_chunk_len` sees from `rayon::current_num_threads()`. Nesting is the second open item. `packed_nbits_output_row` and `int8_row` are called from inside a `par_chunks_mut` for `m > 1` with `parallel_columns` set, so several Rayon workers can now dispatch into the task pool at once. The pool's eight job slots make surplus dispatches fall back to inline execution, which is correct and bounded — but that is an argument, not a measurement. The third is this document's own method. §36.3 is the campaign's first null arm, and it says a large fraction of the per-cell readings in earlier phases — anything inside 0.72×–1.20× at t=32 — carried no information about the change being measured. The 5–15% movements tabulated in several earlier sections should be read with that in mind. Documented as **§36 / phase 16**, including the correction to §35.6. Co-authored-by: Sebastian <sebastian@squad.local> Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

…1364) ## What A packed-panel `6x16` microkernel for `gemm_bt`, the kernel that computes `C = A·Btᵀ` directly on the `[out][in]`-major MoE expert weights (#1241). It is entered only from the **single-threaded** driver on prefill-shaped row extents; decode, small-token MoE and every multi-threaded shape keep the existing `tile_4x2` path unchanged. ## Why MoE expert GEMMs were the largest remaining single-thread loss on the production-native matrix — §31.5/§31.8 of the CPU-EP ledger put MoE prefill at **1.22–1.56x ORT** at `t=512` and attributed it to "microkernel efficiency against MLAS's packed panels and wider register tiles". That attribution is right, and it is fixable without handing anything to ORT's CPU EP. `tile_4x2` keeps **k-partials in the SIMD lanes**. Two consequences: every output cell ends in a horizontal reduction, and the `b` side of the inner loop is strided by `k` — six loads per eight FMAs. ## How | | `tile_4x2` (unchanged, still the default) | `micro_6x16` (new) | |---|---|---| | lanes hold | k-partials | **c columns** | | horizontal reductions | one per output cell | **none** | | per k step | 6 loads / 8 FMAs = 1.33 | 2 contiguous loads + 6 broadcasts / **12 FMAs = 1.5** | | `b` access | strided by `k` | two sequential cache lines | * `pack_panel` transposes one column panel of `bt` into micro-panel-major order (`packed[s*k*PNR + p*PNR + t]`). This is the only transpose in the path, it is bounded by the panel rather than the whole expert bank, and every row band in the sweep reuses it. * `block_packed` drives column panel → row band → micro-panel. Both remainders (`rows % 6`, `n % 16`) fall through to the existing full-`k` `dot`, which is the same code the unpacked path already uses for its edges. * Single-threaded `gemm_bt` hands the **whole row extent** to one block when the packed kernel is selected, so the pack is amortised once rather than once per 64-row `MAX_MC` block. ### The gate is structural, not tuned `PACK_MIN_ROWS = PMR * PACK_MIN_BANDS = 72`, deliberately **above `MAX_MC` (64)**, pinned by a `const` assertion: ```rust const _: () = assert!( PACK_MIN_ROWS > super::MAX_MC, "the packed kernel must stay out of the multi-threaded driver's row \ blocks, which are capped at MAX_MC and would each re-pack the panel" ); ``` The multi-threaded driver slices `m` into `2*threads` blocks capped at `MAX_MC`, and each block re-packs the same panel. Keeping the bar above that cap makes "the panel is packed once per sweep" a property of the code rather than of the shape, and means **multi-threaded behaviour is byte-identical to `main` for every shape**. The assertion stops a future `MAX_MC` bump from silently undoing that. ### Why I will not make a multi-threaded claim I tried to. It does not survive its own control. A 2-thread A/B of two binaries that I *traced* to be on the identical code path (every row block `rows=56..64`, `packed=false`) still came out **~40% apart on the median**. Multi-thread measurement on this shared 32-core box is not usable at this effect size, in either direction, so this PR takes the position that costs nothing: MT is left exactly as it was. Sharing one pack across threads (or making `gemm_bt_col_parallel` the packed path, since each of its tasks already owns all `m` rows) is the natural follow-up, and it needs a quiet machine. ## Evidence Both binaries built `--no-default-features --features bench-native,cuda-13000`, differing **only** in `matmul.rs`. The arm is proved by the linker, not the flag: `nm -C <bin> | grep -ci mlas` = **0** on both. Arms interleaved innermost in one driver process, arm order alternating per rep, thread-matched `--native-threads 1 --ort-intra-threads 1`. Metric is `native_min / ort_min` (lower is better; <1.000 beats ORT). ### Single thread, `t=512` — the shapes that reach the new kernel Final gate, shipped binary (6 reps): | model | before | after | Δ min | Δ median | parity | |---|---|---|---|---|---| | `moe_phi35moe_h2048_i6400_e4_t512` | 1.356 | **1.009** | −25.61% | −16.50% | PASS | | `moe_qwen3moe_h2048_i768_e16_t512` | 1.297 | **1.101** | −15.12% | −15.92% | PASS | | `moe_mixtral_h1024_i3584_e8_t512` | 1.226 | **1.173** | −4.27% | −2.85% | PASS | Four independent passes, Δ on the min ratio: | model | pass 1 | pass 2 | pass 3 | pass 4 (shipped) | |---|---|---|---|---| | qwen3moe t512 | −16.7% | −23.4% | −19.6% | −15.1% | | phi35moe t512 | −16.9% | −24.9% | −27.6% | −25.6% | | mixtral t512 | −1.3% | −6.6% | −3.0% | −4.3% | **Phi-3.5-MoE prefill crosses 1.000 in the shipped pass and Qwen3-MoE lands at 1.10**, from a 1.30–1.54 loss. Mixtral moves least, and I believe that is real rather than measurement: its expert bank is ~352 MB of fp32, so that shape is closer to DRAM-bound than microkernel-bound and a better inner loop has less to win. ### `t=1` / `t=32` — unchanged by construction, and verified Traced with a temporary instrumented build: every `gemm_bt_block` at `t=32` has `rows` in **14–19**, far below the gate, so `packed=false` and the executed code is the unpacked path. Their run-to-run spread (±10% on min, ±3% on median, sign flipping between passes) is therefore a direct read of this box's **noise floor**. At `t=512` the traced extents are 74–365 rows, all packed. ### Correctness * `cargo test -p onnx-runtime-ep-cpu --lib` → **1434 passed, 0 failed, 17 ignored**. * New test `gemm_bt_packed_kernel_matches_naive_single_threaded`: 10 shapes against `naive_bt` on a 1-thread pool, covering both remainders, `k` tails, the shortest admitted `k`, and multi-column-panel sweeps. It carries a **non-vacuity assertion** — if a policy change stops these shapes from reaching the packed kernel, the test fails loudly rather than silently measuring the old path. * **Falsified twice**: swapping the two stored halves of the micro-tile, and corrupting the pack layout to `dst[t*k+p]`. Each mutation fails **only** the new test; the other five `gemm_bt` tests stay green. * Clippy clean, `cargo fmt --check` clean, `cargo check --features mlas` compiles. * Parity PASS in every benchmark cell above. ### Review Reviewed by Opus (`claude-opus-4.8`), verdict **APPROVE-WITH-NITS**, no blocking findings. It independently re-ran the packed test with 7 additional adversarial shapes (the `nc`-floor branch at `k>8192`, single-micro-panel wide sweeps, whole-`m` widening at `m=1000`, maximal row/column remainders) **under AddressSanitizer** — all clean — and verified the panel-index bounds, the `jend-j0 ≡ 0 (mod PNR)` contract, remainder exhaustiveness/disjointness, and that no consumer of `gemm_bt` depends on its accumulation order. Its one substantive nit — that the original `rows >= 48` gate did *not* keep large-`m` multi-threaded blocks out of the packed path, contrary to the comment — is what produced the structural `> MAX_MC` gate and the `const` assertion above. ## Scope * No ORT CPU fallback, no deferral — assigned == executed is unchanged. * Multi-threaded blocking policy untouched; MT is on the identical code path. * `tile_4x2`, `block`, `dot`, `gemm_bt_block_scalar` and the `mlas` arm untouched. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

… that failed (#1373) Documentation only — records Phase 17 in the CPU-EP ledger. ## What lands here * **§37.1–37.2 — #1245**, the pure-Horner `exp8`. Full 7-model A/B across two independent passes, plus the accuracy contract: 1 ULP over **1,118,743,631** `f32` values by exhaustive sweep, and the story of the sweep that first reported `max_ulp = 8388646` because it included the deliberate flush-to-zero threshold. * **§37.3 — #1364**, the packed-panel `6x16` `gemm_bt` microkernel. Four independent passes. Phi-3.5-MoE prefill reaches **1.009** (parity with ORT) from 1.356; Qwen3-MoE **1.101** from 1.297. * **§37.5** — the unmerged two-pass online softmax, kept as a negative result (2.4× worse) with the reason it loses: §29 already removed the copy that would have made the extra pass expensive. ## §37.4 is the part that matters Choosing the packed kernel's row gate produced two negative results, and the second one is a methodology finding, not a perf finding. A four-band gate looked fine single-threaded and lost ~29% on the 8-thread Phi-3.5-MoE median. Raising it to eight bands removed that, but then the 2-thread cells showed +6% to +11% — which reads as a real multi-threaded regression. It is not usable data. After raising the gate above `MAX_MC`, the two binaries are **traced to be on the identical code path at two threads** (every row block `rows=56..64`, `packed=false`, verified with an instrumented build) — and they still measured **~40% apart on the median**, in the favourable direction for the arm that had just been slower. That is a worse null arm than §36.3's: ~40% at **two** threads on a code path proved identical, against §36.3's ±20–30% at thirty-two. Low thread counts are not the quiet regime on a shared box; they are the regime where one noisy neighbour is a larger fraction of what is running. The consequence is stated plainly in the section: this document now has two independent measurements of its own instrument disagreeing with itself by 20–40%, and every cross-arm claim in it that rests on a single pass without a null control should be read as provisional. ## Scope No code. No behaviour change. Both PRs being documented are already merged (#1245 as `512885a7b`, #1364 as `307c7277d`). Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

…ent ORT pool (#1374) Documentation only — no code change, no behaviour change. §36.8 left the t=32 residual open and called it "a scheduler signature". It is — but it is mostly not *our* scheduler. This phase spent its budget on measurement. ## What it found `bench_generic --native-only` / `--ort-only` build only the arm being timed, so running the same binary with and without a co-resident ORT session isolates contention from arithmetic. At t=32: | cell | native alone | native paired | ratio | | --- | --- | --- | --- | | `gemm_nbits_llama3_8b_qkv_t8` | 6.90 ms | 28.71 ms | **4.18×** | | `gemm_dense_tall_128x4096` (f32 dense) | 5.62 ms | 26.92 ms | **4.79×** | | `gemm_nbits_llama3_8b_mlp_t8` | 15.99 ms | 42.56 ms | **2.66×** | The first row reproduced to three significant figures across two independent sessions an hour apart (28.698 / 28.713 paired, 6.868 / 6.900 alone) — four times outside §36.3's measured noise band. Measured **alone**, `llama3_8b_qkv_t8` is 7.15 / 6.22 / 6.87 ms at t=8/16/32: flat, no inversion. The inversion only exists in the paired numbers. ## Why The tax is asymmetric — ORT pays only 1.2–1.3× for our co-residency, we pay up to 4.2× for its. ORT's intra-op pool spin-waits long after its last op; our task runtime spins briefly and parks. A pool that parks quickly is invisible to its neighbours; a pool that spins is not. It hits the **dense f32 control hardest**, so it is contention, not anything about int4. Lane width is ruled out first (§38.1): t=16 and t=32 both run 16 lanes and still differ ~3×, and the default width is at or near optimal on 3 of 4 cells. ## What it changes - Adds the rule to the method: long cells above 16 threads must be measured with solo arms. The paired harness stays correct for parity, short cells, and base-vs-new comparisons of our own binaries. - Corrects the overstated ratios: the honest t=32 figure for `llama3_8b_qkv_t8` is **10× behind ORT, not 41×**. Still a real gap, still a kernel problem. - **#1363 / §36 is unaffected** — both arms there were native binaries paying the same tax. - Records an affinity probe that failed its own noise check (§38.5), so the next attempt knows it has been run once on a loaded box for nothing. - Downgrades §36.8's `with_decode_pool` hypothesis: it early-returns inline when `with_decode_pool_scope` is active, so the `install` is per-pass, not per-node. ## Validation Docs only. Full local `Rust quality` proxy run anyway: `cargo fmt --all --check` clean, all 8 guard scripts pass, `workspace_test_packages verify` OK. Co-authored-by: Sebastian <sebastian@squad.local> Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

…, byte-identical) (#1375) ## Summary Ports `wide_multicol`'s register-blocking to the **grid-starved split-K int4 decode GEMV** shapes (GLM down_proj / qkv / attn-out), the #1 GLM decode kernel (41% of decode) that Wallace's ncu diagnosis pinned as latency-bound at M=1 (16.6% DRAM peak, No-Eligible 30.8%). New kernel `matmul_nbits_gemv_f16_general_bs_splitk_wide_multicol` (+ symmetric-interleaved sibling): K_SPLIT=4 warps cooperate per column group (shared-mem fp32 partial reduction, exactly like `splitk_wide`) **while each warp register-blocks WIDE_NC=4 output columns** — issuing 4 independent 128-bit weight loads/chunk for the memory-level parallelism the single-column split-K lacked. **Byte-identical** to the old `splitk_wide` (same per-lane depth0/stride, per-column accumulation order, and K_SPLIT reduction order). Default-on, reversible via `ONNX_GENAI_GEMV_SPLITK_MULTICOL=0`. ## Correctness (byte-identical — no reduction-order change) - ✅ DeepSeek-V2-Lite golden decode lock: **PASS byte-identical** (run alone, `--test-threads=1`). - ✅ GLM-4-9B decode token IDs byte-identical to baseline. - ✅ `matmul_nbits_gpu` decode gates in isolation: `fp16_symmetric_splitk` (split-K path), `accuracy4_block128`. - ℹ️ 2 unrelated `matmul_nbits_gpu` failures (`block128_batched_asymmetric_non_multiple_k`, `int4_fp16_prefill`) are **M>1 prefill** tests that fail identically with this kernel disabled → pre-existing on `main`. ## Performance (H200, greedy, CUDA-graph, `glm-4-9b-int4-cuda-ortfair`, medians) | Config | before | after | Δ | ORT ref | after vs ORT | |---|---|---|---|---|---| | short 128 | 211.8 | **218.7** (best 227.8) | **+3.3%** | 249.83 | 1.14× (was 1.18×) | | deep 2600 | 194.4 | **207.7** | **+6.8%** | 243.04 | 1.17× (was 1.25×) | **Mechanism (ncu):** #1 kernel DRAM **16.62% → 19.88%** peak (register-blocked MLP; higher DRAM at lower occupancy — the `wide_multicol` signature). ## Honest scope Real, correct, mechanism-proven improvement that **narrows but does not close** the ORT gap. The down_proj shape can't get both WIDE_NC=4 MLP and >1.5 waves; **K_SPLIT=8 and a depth-2 register pipeline both measured regressions**, so K_SPLIT=4/WIDE_NC=4 is the measured optimum (~20% DRAM, not the 37% full multicol regime). Reaching/passing ORT needs a further lever + Wallace's independent attention-depth-scaling fix. ⚠️ Do not self-merge — coordinator validates the golden lock alone on a pinned GPU and admin-merges. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> Co-authored-by: Sebastian <sebastian@squad.local> Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

…1377) Closes the last open item from phases 16 and 18. Test-only — no kernel or runtime code changes. ## The gap §36.8 and §38.7 both said the same thing: `packed_nbits_output_row` and `int8_row` dispatch into the task pool from inside a `par_chunks_mut` for `m > 1` with `parallel_columns` set, so several Rayon workers can dispatch concurrently. The pool's eight job slots bound the surplus by declining it back to the caller to run inline — *"correct and bounded", but that is an argument, not a measurement, and it has not been measured.* ## The measurement `nested_dispatch_slot_pressure` joins the existing `task_runtime_latency` harness and reproduces the shape directly: an outer `par_chunks_mut` over rows, an inner `task_runtime::for_each_range` in each, with the pool counters read across it. `#[ignore]`d like its neighbour, since it is a measurement. Four runs against a 16-lane pool: | outer dispatchers | wall | dispatches | declined | | --- | --- | --- | --- | | 1 | ~1.9–2.4 ms | 1 | 0 | | 2 | 0.29–0.44 ms | 2 | 0 | | 4 | 0.45–0.74 ms | 4 | 0 | | 8 | 0.84–1.51 ms | 8 | 0 | | 16 | 1.35–1.84 ms | 13–16 | **0** (3 once) | **Slot exhaustion essentially does not happen.** Three of the four runs declined nothing at any width; the single run that declined 3 of 16 was a cold process's first dispatch. The slots turn around faster than sixteen Rayon workers can collide on them, so the inline fallback is a real safety net that is almost never used. Per-row cost also improves monotonically (175 → 99 µs/row from 2 to 16 dispatchers), so the nesting is not serialising. The ~2 ms one-dispatcher row is pool construction on the process's first dispatch. ## What it asserts Timing is reported, not asserted. What *is* asserted is the property that matters under nesting: every element covered exactly once (so the disjoint `start..end` reconstruction never overlaps), and no task body panicking. ## Validation - `cargo test -p onnx-runtime-ep-cpu --release`: **1447 passed, 0 failed**, 17 ignored - `cargo clippy -p onnx-runtime-ep-cpu --release --all-targets -- -D warnings`: clean - `cargo fmt --all -- --check`: clean - all 8 `Rust quality` guard scripts: pass Co-authored-by: Sebastian <sebastian@squad.local> Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

) Updates the campaign brain to reflect merged #1375 (6cde53b): GLM multicol × split-K register-blocked int4 GEMV, byte-identical golden-locked (coordinator-validated GPU3). GLM deficit narrowed 1.18x->1.14x short / 1.25x->1.17x deep. GEMV now at empirical optimum; remaining deep deficit is attention-depth-scaling (Sebastian on 2nd lever). Also adds the top-ledger #1375 entry. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

…h it (#1380) ## The problem Two phases in a row reported a delta and then had to reconstruct, *after the fact*, whether the instrument could resolve it — §36.3 from cells the change could not reach, §37.4 from a code path traced to be identical. That reconstruction is an argument, and it arrives too late to change what was measured. ## The change `ab.py --null-control` adds a third arm that is **the first arm's own binary, with the first arm's `--arm-env`, under a second name**, interleaved and order-alternated exactly like the real arms. It cannot measure the change. Its delta is the host's noise floor for that cell, *in that invocation*. The driver also gained a deltas table (it previously printed only per-arm medians, leaving the subtraction to the reader): ``` === deltas vs 'before' (median ratio; negative = arm is faster) === The 'null' column is the same binary as 'before'. Its delta is this host's noise floor for the cell, so a real delta no larger than it is not a result. moe_phi35moe_h2048_i6400_e4_t512 t=1 after: -25.48% > noise (0.19%) null: -0.19% moe_mixtral_h1024_i3584_e8_t512 t=1 after: -2.38% > noise (0.63%) null: -0.63% ``` Anything inside the control prints `WITHIN NOISE` instead of a number that looks publishable. ## What it measured (§38) Same-binary control arm, `|Δ|` of the median native/ort ratio against the identical binary: | cell | 1 thread | 2 threads | 8 threads | | --- | --- | --- | --- | | `moe_mixtral_h1024_i3584_e8_t512` | 0.63% | 0.22% | **19.54%** | | `moe_qwen3moe_h2048_i768_e16_t512` | 4.75% | 3.97% | **12.68%** | | `moe_phi35moe_h2048_i6400_e4_t512` | 0.19% | 4.28% | 3.03% | | `sm_bert_b8_s128` | 0.21% | — | — | | `sm_decode_h32_kv8192` | 1.32% | — | — | | `sm_prefill_h32_s512` | 0.22% | — | — | **Single-thread cells are tight** — under 5%, mostly under 1.5% — so the single-thread ratios this ledger publishes are resolvable at the effect sizes it claims. **Eight-thread cells are not**: a 12–20% floor retires the 5–15% movements several earlier phases tabulated at `t=8`, exactly as §36.3 warned. Both phase-17 merges are re-measured against their own control, in the same invocation — the first self-controlled cells in the document: | change | cell (1 thread) | Δ median ratio | floor | verdict | | --- | --- | --- | --- | --- | | #1364 packed panel | `moe_phi35moe…t512` | **−25.48%** | 0.19% | 134× the floor | | #1364 packed panel | `moe_qwen3moe…t512` | **−21.56%** | 4.75% | 4.5× the floor | | #1364 packed panel | `moe_mixtral…t512` | −2.38% | 0.63% | 3.8× the floor | | #1245 Horner `exp8` | `sm_bert_b8_s128` | **−10.28%** | 0.21% | 49× the floor | | #1245 Horner `exp8` | `sm_decode_h32_kv8192` | **−11.73%** | 1.32% | 8.9× the floor | | #1245 Horner `exp8` | `sm_prefill_h32_s512` | **−10.33%** | 0.22% | 47× the floor | ## It corrects my own §37.4 §37.4 reported ~40% between two distinct binaries traced to identical code paths at two threads, and read that as the two-thread noise floor. **That reading was wrong.** The same-binary control puts the two-thread floor at 0.22–4.28%, an order of magnitude tighter, and the absolute ratios differ between the sessions (mixtral `t=2` reads 1.35 here, 2.23 there, same fixture, same thread count). The 40% was a *session*: a load episode long enough to outlast the arm alternation. That is a more useful finding than a floor, because it is not something interleaving can fix — a whole invocation can be poisoned, and only a control arm inside that invocation exposes it. §38.3 states the narrower rule that follows: run the control in every invocation whose result will be published, and discard the **invocation**, not the cell, when the control moves more than the effect. ## Scope * Additive: without `--null-control` the driver behaves as before, except that multi-arm runs now also print a deltas table. * `ruff check` clean. (`ruff format` does not pass on this file before or after this change, so it is left alone rather than reformatted in a perf-tooling PR.) * No Rust, no behaviour change in the EP. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

…space abort (Bug 1) (#1379) ## Headline correction (read first) **Bug 1's default greedy-decode path is already fixed on `main` by #1189** ("Fix Engine long-context Attention workspace", `b416a3e08`), which added `prepare_decode_workspace_after_capacity_growth(grew)` — the exact re-prepare of the persistent `::Attention` fp32 score workspace when the KV bucket grows. The original abort repro was measured on a pre-#1189 binary. On current `main` the repro (**generate 320 tokens from a 1-token prompt with default KV**) completes cleanly: the bucket grows 256→512 and #1189 handles it. This PR therefore does **not** re-implement the main-path fix. It adds the pieces #1189 left open. ## What this PR adds 1. **End-to-end GPU regression guard** (`tests/decode_workspace_bucket_growth.rs`, ignored/feature-gated) locking the generate-past-256 repro so #1189 can't silently regress: - `deepseek_v2_lite_native_cuda_generates_past_kv_bucket_growth` — MoE, 320 tok, default KV. - `qwen2_5_0_5b_native_cuda_generates_past_kv_bucket_growth` — dense, 320 tok (non-MoE past-256). - `qwen2_5_0_5b_native_cuda_speculative_generates_past_kv_bucket_growth` — prompt-lookup spec-decode liveness. 2. **Two residual eager-path guards** #1189 missed (`native_decode/cuda.rs`): the spec-decode verify path `decode_cuda_eager` and the routed-pipeline path `decode_cuda_eager_step_inputs` called `ensure_capacity` without the re-prepare. Same invariant, same fix, only fires on `grew`. ## Root cause KV capacity grows in powers of two (256, 512, 1024…). A decode step that grows the bucket but does not re-prepare the persistent governed workspace leaves the `::Attention` fp32 score scratch (`batch·q_heads·q_seq·total_seq·4`) sized for the old bucket; the next execute needs exactly 2× and trips the strict prepared-workspace invariant (`required > prepared`). ## Graph-capture safety Both patched sites are **eager (non-captured)**: `decode_cuda_eager` invalidates any captured graph up front; `decode_cuda_eager_step_inputs` is the explicit non-capture routed path. The re-prepare fires only on the `grew` transition (which already invalidates + forces re-capture), so no device pointer/layout changes mid-replay. The #1366 golden lock runs under graphs and stays byte-identical. ## Verification (GPU2, `--test-threads=1`, ORT 1.27.0) - **DeepSeek 24-tok golden decode lock: byte-identical (pass).** - All 3 new regression tests pass. - `native_prompt_lookup_matches_plain_greedy_cuda` (exercises `decode_cuda_eager`) still token-identical (no regression). ### Honest caveat I could not trigger a live abort *without* the eager guards in this environment: in the prompt-lookup driver the guaranteed greedy token (already guarded) grows the bucket first, and the low-level `decode_verify` path that would force it can't load the deployed Qwen model here (the pre-existing `native_cuda_verify_rewind_no_kv_corruption` test also fails at load with a `model.io.token_input` ambiguity — an env/metadata issue, not this change). The eager guards are justified by code inspection + correct-by-construction (identical to the proven #1189 sites) and verified no-regression. ## Relationship to #1223 (Bishop) Sibling PR #1223 (`squad/1223-reprepare-on-rebucket`) fixes this **generally** at the dispatch layer via an `eager_workspace_growth` flag, a **superset** that also covers both eager paths. **These overlap — merge one, not both.** #1223 is the cleaner general fix; my per-path guards become redundant (harmless) if #1223 lands. I did not duplicate #1223's dispatch change. ## Not in scope **Bug 2** (long single-shot prefill ≥ ~500 tok requiring a constant ~42 MB node-38 workspace) is deliberately untouched — separate follow-up. Do **not** self-merge — coordinator validates + admin-merges after running the golden lock across int4 models. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

…8%, crosses ORT +3.5%) (#1383) ## Summary Native **eager** decode was ~25% behind ORT eager on DeepSeek-V2-Lite short-ctx (**66 vs ORT 86.78 tok/s**) — a pure per-token **host-dispatch** deficit, not GPU compute (CUDA-graph capture already hid it: graph=1 ~143 tok/s). This PR closes and reverses that gap, **ON by default**. ## Root cause Two numerically-inert host serializers per op on the eager (`!capturing`) branch: 1. **Redundant trailing per-op stream drain** — every kernel ended with `synchronize()`. It buys nothing: the single in-order EP stream guarantees kernel to kernel ordering, and `dtoh`/`dtod` self-synchronize before their synchronous copy. 2. **Eager-only validation D2H readbacks** — rotary/gather/structural/GatherElements copy 24-byte scalars to host purely to bounds-check indices. A correct model never trips them; the captured (graph=1) path already uses a device error-latch instead. Per-token cost (nsys): DeepSeek MoE ~1187 launches, **416 cuStreamSynchronize**, **147 D2H**/token; ~54% of eager wall is recoverable host overhead. ## Fix (default ON; `ONNX_GENAI_DEFER_EAGER_SYNC=0` escape hatch) - `runtime.rs`: `defer_eager_sync` defaults **true** unless the env is explicitly falsey (`0`/`false`/`off`/`no`). Public `synchronize()` no-ops when deferred; new private `force_synchronize()` keeps the real drain for `dtoh`/`dtod` host reads (read correctness preserved). - `rotary_embedding.rs`/`gather.rs`/`structural.rs`/`indexing.rs`: skip eager validation D2H when deferred. Default-ON makes eager **consistent with the already-shipped captured path**; the env var is a one-flag rollback for debugging. ## Results (GPU3, greedy, medians-of-5, pinned idle GPU, flag UNSET = default ON) | Config | Native tok/s | vs ORT eager 86.78 | |---|---|---| | DeepSeek eager (old path, `=0`) | 66.06 | -24% | | **DeepSeek eager (default ON)** | **91.64** | **+5.6% WINS** | | GLM-4-9B dense eager (off to on) | 142.42 to 145.31 | +2.0% | | DeepSeek graph=1 (default ON) | 143.12 | fallbacks=0, still valid | ## Validation (flag UNSET, default ON) - **DeepSeek 24-tok golden lock BYTE-IDENTICAL:** yes (run alone, 21.18s, 1 passed). - **Graph=1 still valid:** yes captures=6/replays=372/**fallbacks=0**. - **Dense qwen2.5-0.5b logprobs byte-identical** vs old path: yes (md5 `d2e4bc2e...`). - **Escape hatch `=0` restores old path:** yes (66.06 tok/s). ## Recommendation Ship default-ON (this PR). Remaining headroom (DeepSeek eager 91.6 vs graph 143) is the ~1187 launches/token to **kernel fusion**, a separately-scoped follow-up. Findings: `.squad/decisions/inbox/wallace-eager-launch-overhead.md` Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> --------- Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

Records the #1383 native-eager-beats-ORT-eager default-ON win (DeepSeek 91.64 vs ORT 86.78, +5.6%; GLM +2%) and the #1379 Bug-1 regression guards in now.md. Also flags the pre-#1383 eager-vs-eager table as stale for the eager column. Docs-only. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

…1386) fix(ci): scope Wiki Pages triggers and stop cancelling main deploys (#1386) Merged via admin override: the two required checks (Fast (Linux x86_64), Rust quality) were stuck queued for 35+ minutes behind the repo's severe CI backlog (the exact class of noise this fix reduces). The change is workflow-YAML only, touches no Rust/build code, and was independently validated with actionlint + YAML parse + a manual event/path/concurrency matrix before merge.

## What The executor's kernel cache is keyed by node **and input shapes** and never evicted. Every distinct prompt length therefore compiled a fresh kernel for every node in the graph, and each compiled kernel owns device workspaces (attention scratch, dequantization scratch, broadcast metadata). A server whose requests each bring a new prompt length accumulated those workspaces without limit until the device ran out of memory. This bounds the cache to a small number of shape variants per node, evicting the least recently used beyond it. ## Why it looked like something else The bytes are invisible to the resource governor, which accounts for weights, KV pools and bindings but not for memory a kernel allocates for itself. `/v1/resources` reported `vram.used` = weights only while `nvidia-smi` showed 70+ GiB. That is why the failure first read as "prefill activations scale with prompt length". What separated the two hypotheses: replaying the same prompt lengths in **descending** order still grew memory monotonically (43 → 72 GiB). A per-request peak would have shown its maximum on the first, largest request. Attributing live `cuMemAlloc` blocks to their allocation sites showed the growth as thousands of retained per-node workspaces (`MatMulNBitsKernel::run`, `GqaWorkspace::reserve`, `BroadcastMetadataCache::prepare`) growing in lockstep at one group per graph node per distinct shape. ## Design - **Per node, not global.** Eviction stays proportional to what actually varies: a graph keeps its hot prefill and decode variants, and only a node that has genuinely seen many shapes gives one up. A global bound would let one busy node evict every other node's steady-state kernel. - **LRU, not FIFO.** The single-token decode shape is the most frequently served and must not be evicted by a run of new prefill shapes. - **Capture safety.** The EP's captured device graph is reset before an eviction frees a kernel, because a capture can have baked the pointer of a workspace the eviction is about to release. This mirrors what the device-binding drop path already does. - **Escape hatch.** `ONNX_RUNTIME_KERNEL_CACHE_VARIANTS_PER_NODE` overrides the default of 4. ## Measured 30B int4 model, native CUDA backend, 12 requests with prompt lengths from 469 to 5502 tokens: | prompt tokens | before (MiB) | after (MiB) | |---|---|---| | 469 | 38281 | 38311 | | 1384 | 56803 | 49081 | | 2757 | 72757 | 52943 | | 3214 | **OOM** | 49829 | | 4587 | 69915 | 53585 | | 5045 | **OOM** | 50991 | | 5502 | 69171 | 47671 | Before: three out-of-memory failures. After: every request succeeds, memory is flat, and latency is unchanged (43.7s vs 48.1s on the longest prompt). ## Tests `kernel_cache_bounds_the_shape_variants_it_keeps_per_node` feeds one node 24 distinct shapes and asserts the cache evicts rather than keeping one entry per shape, while every run still returns the correct result. `cargo test -p onnx-runtime-session` (all green), `cargo test -p onnx-genai-engine --lib` (427 passed), fmt and clippy clean. Fixes #1362 Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> Copilot-Session: 0190e2eb-abe4-451f-b36d-44a035a99b7e

## What A model directory can declare a prefill chunk size, and the ORT backend splits a prompt into that many tokens per forward. The native backend read the metadata and then never applied it — a prompt arrived at the decode session as one tensor and ran as one forward. The two backends therefore disagreed about the shapes a given model produces. This applies the declared size in the native session. ## How A prefill longer than the chunk runs as a sequence of forwards, each appending to the KV cache, keeping only the final chunk's logits (the earlier ones describe tokens the caller already has). Chunking has to slice whatever the caller passed, not just token ids. A multimodal prompt arrives as `inputs_embeds`, where the sequence lives on the middle axis of a `[1, tokens, hidden]` tensor, so step inputs are sliced along that axis alongside the ids. A tensor that does not expose a sequence axis that way is left whole rather than guessed at — slicing the wrong axis would silently corrupt the prompt instead of failing. The size is injected on **both** paths that build a native decoder: the direct engine load, and the pipeline builder. The latter matters because it constructs its decoder lazily on the first request and would otherwise never see the metadata. ## Not a memory fix This came out of investigating #1362, but it is not the fix for it. The pipeline caller already chunked its own prefill before reaching the decode session, so the accumulating device memory on long prompts came from elsewhere (see #1388). This PR stands on its own as backend parity. ## Tests `prefill_chunk_tests` covers the slicing directly: which tensors expose a sequence axis, that a slice takes exactly its row range out of the middle axis, and that a slice past the end is refused rather than truncated. `cargo test -p onnx-genai-engine --lib --features native-backend` — 544 passed. fmt and clippy clean. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> Copilot-Session: 0190e2eb-abe4-451f-b36d-44a035a99b7e

CI's format lane had drifted again — `cargo fmt --all -- --check` was exiting 1 on `main` at `f72700a8`. One file, `crates/onnx-runtime-ep-cuda/src/runtime.rs`, +3/-1. **No logic changes**: `cargo fmt --all` only. Verified locally: `cargo fmt --all -- --check` exits 0 on this branch. Co-authored-by: justinchuby <223556219+Copilot@users.noreply.github.com> Copilot-Session: d60eb808-7cc6-4abc-b48d-2a6dd3841624

…7x kernel, every ORT prefill cell from loss to win) On a native (non-`mlas`) build there was no borrowed route for `bits == 8, m > 1` at all. `try_prefill_mlas_nt` declines, and the kernel fell through to `dequantize_weight(WeightLayout::Kn)` + `gemm`: a full `k * n` f32 materialization, written four bytes out per byte in, at stride `n`, into a buffer far larger than any cache -- and rebuilt on *every call*, because nothing caches the `Kn` layout. A 3584x3584 8-bit node materialized 51 MB per prefill to multiply it once. Generalize #1117's fused-dequant GEBP to both bit widths. `pack_b_quant` expands each `KC x NR` panel straight into the L1-resident f32 panel the `6x16` microkernel consumes, and all `m` rows reuse it. The bit width is now a `BlockQuantWeight` impl (`Int4Weight` / `Int8Weight`), so the two share packing, blocking, threading and microkernel and differ only in `dequant_column`. Kernel level, `k = n = 3584`, steady state, median of 3 interleaved arm-pairs: 17.1x at m=2, 13.9x at m=8, 5.8x at m=64, 2.6x at m=256, 1.8x at m=512 (843 GFLOP/s). At 4096x11008, m=256 reaches 1023 GFLOP/s. The outputs are bit-identical to the route replaced -- same `(q - zp) * scale` accumulated along `k` in the same order -- which the A/B bench asserts over the full output of both arms, and which is why the tests need a route counter. Against ONNX Runtime at the assignment matrix's own geometry (`k = n = 3584`, paired arms, 7-11 interleaved trials): M=128, t=2/4/8: 2.06/2.01/2.32 -> 0.66/0.75/0.68 M=256, t=2/4/8: 1.70/1.69/1.89 -> 0.77/0.83/0.88 M=512, t=2/4/8: 1.41/1.42/1.54 -> 0.86/0.88/1.01 Every measured 8-bit prefill cell crosses from a loss to a win except `M = 512` at 8 threads, which reaches parity and is reported as parity. Decode (`M = 1`) is untouched and unchanged. The pure-native path is now within 1.05-1.35x of a `--features mlas` research build, where it was 2.5-3.7x behind. `INT8_PREFILL_GEBP_MIN_ROWS = 2` is derived, not inherited: the constant was temporarily lowered and m=2,3 measured on both arms. There is no crossover to find here -- the replaced route pays its 51 MB whatever `m` is -- so the gate is the GEBP's own lower bound. Also teaches `scripts/ort_ab/gen_gemm.py` to emit 8-bit cells and the square `3584` geometry, because the assignment matrix carried 8-bit rows no generator in the tree could reproduce. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

Per user directive ("对dual head size的支持也要general 万一以后有更多情况呢") and the standing generality principle in RULES.md §2, make head size (head_dim) an explicit member of the no-hardcoding prohibition and add a clarifying bullet: - head size is a fully runtime, per-attention-op parameter (loader → attention kernel → GEMV → KV-cache alloc) - no fixed values (128/256), no fixed count of distinct sizes, no `dual`-specific branches - arbitrary/mixed per-layer head sizes supported generally; an N-distinct-head-size model needs no new special case Docs-only clarification of existing Rule 2 intent; binds the in-flight gemma4-e2b dual-head-size work. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> Co-authored-by: Squad <squad@users.noreply.github.com> Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

…identical); hard-stop plain half at 2 ULP (#1404) ## What this does Lands the **byte-identical half** of the batch-decode CUDA-graph fragmentation fix (#1334): small-batch (`1 < M <= 8`) fused **gate/up-SwiGLU** decode is routed through a per-row capture-safe `M==1` GEMV loop instead of the tiled prefill GEMM, so the batch-decode graph captures as **one segment** instead of fragmenting into one eager seam per MLP layer. The routing is **gated to the RMS-norm-prologue path only** (`gamma.is_some()`); the plain (no-rmsnorm) path is hard-stopped per the decision rule below. Lives in the shared CUDA EP kernel, so native and ORT both inherit it (DRY satisfied). ## The invariant this touched, and why it was safe to change The blocker was the assertion **`"only M=1 decode may be advertised capture-safe"`** in `matmul_nbits.rs`. That invariant is **already stale on `main`**: default-on Marlin (`marlin_m_gt_1_enabled()`, `ONNX_GENAI_MARLIN_M_GT_1` defaults on) makes `run_f16_gate_up_swiglu` advertise `capture_safe = warm` after the **Marlin M>1** path, contradicting the `m == 1` rule. On pristine `main` this already reds `fused_gate_up_swiglu_rmsnorm_is_bit_exact_to_two_step_path` and `..._zero_points...` under the default config (see table). Filed separately as #1405. Rather than delete or widen the invariant, this PR **narrows it to exactly the capture-safe routing** on the rmsnorm path and **keeps it strict (`m == 1`) on the plain path**, with the reasoning written at each assertion. ## The measurement (ULP sweep) — the whole reason this PR exists The prior work claimed the fused decode GEMV is within 1 ULP of the two-op reference. **Measured instead of argued**, sweeping the fused decode GEMV vs the two-op reference (two standalone `MatMulNBits` projections + reference `silu_mul`) across activations (many seeds), shapes, dtypes, at **M==1 and M>1**: | path | M==1 (apples-to-apples decode) | M>1 | |---|---|---| | **rmsnorm-prologue** (`fp16_gate_up_swiglu_rmsnorm_two_op_ulp_bound_sweep`) | **0 ULP** (0/60 cases diverge) — byte-identical | ≤3 ULP | | **plain** (`fp16_gate_up_swiglu_two_op_ulp_bound_sweep`) | **2 ULP** (8/56 cases diverge) | — | The M>1 ≤3 ULP is a pure GEMV-vs-GEMM (decode-vs-prefill) reduction-order artifact, **not** a batching-contract violation. **M==1 is genuinely equivalent for the rmsnorm path (0 ULP) and merely under-sampled for the plain path (2 ULP).** The prior "1 ULP" claim was refuted: the plain fused decode GEMV is up to 2 ULP off the two-op reference even at M==1. ## Decision rule applied > If any configuration exceeds 1 ULP: hard-stop the non-rmsnorm half; land only the byte-identical rmsnorm half. The plain path exceeds 1 ULP at M==1, so: - **Landed:** the rmsnorm-prologue path (0 ULP at M==1, byte-identical). This is the path the production skip-rmsnorm fusion (`CudaGateUpSwiGluFusion` + `MATMUL_NBITS_RMSNORM_PROLOGUE_ATTR`, tested by `folds_skip_rmsnorm_into_gate_up_swiglu_node`) folds the pre-MLP RMS norm into for **every RMSNorm architecture (Qwen/Llama/Phi)** — i.e. the resident-model decode path. The production loop is gated to `gamma.is_some()`. - **Hard-stopped:** the plain path. Its capture-safe invariant stays `m == 1`; plain M>1 stays on the prefill/Marlin GEMM. Recorded in #1334. **Reproducibility argument:** under the default config, plain M>1 decode already runs through Marlin's split-K, which is documented non-byte-identical and possibly non-deterministic. The rmsnorm decode GEMV that *is* landed is a **deterministic 0-ULP** match to the two-op reference at M==1 — an improvement in reproducibility, not a relaxation. ## Invariants kept (the batching contract) `gate_up_swiglu_rmsnorm_loop_is_byte_identical_to_per_row_singlestream` (new) asserts the strong contract: row `r` of an `M>1` rmsnorm batch decode is **bit-for-bit equal** to running that row alone as `M==1` — *a request batched with others produces the same tokens as if it had run alone*. This holds by construction (each row dispatches the identical `launch_gate_up_swiglu_rmsnorm` over a 1×K sub-view). ## Performance (from the prior branch work; preserved here) The fragmentation fix was measured on the branch at **25 → 1 graph segments** and **5.15 ms saved at batch-2 (4.9×)** on `qwen05b-q4` (resident, RTX 4060), vs a ~5 ms prediction, with the full +10.4 ms M=1→M=2 delta attributing to `kernel_host_dispatch`. Streaming-neutral: `htod_bytes_per_token` unchanged on `qwen14b-zp`. This PR **preserves** that win: qwen05b's gate/up node carries the rmsnorm prologue, which is the only path still looped. (Plain, no-rmsnorm gate/up decode is not the resident-decode path and is the hard-stopped half.) ## Test results — full `onnx-runtime-ep-cuda` lib suite (`--features cuda --lib -- --test-threads=1`), RTX 4060, all on base `f72700a8` | | Marlin default-**ON** | Marlin M>1 **OFF** (`ONNX_GENAI_MARLIN_M_GT_1=0`) | |---|---|---| | pristine `main` (f72700a) | **404 passed / 6 failed / 21 ignored** | **407 passed / 3 failed / 21 ignored** | | this PR | **409 passed / 4 failed / 21 ignored** | **410 passed / 3 failed / 21 ignored** | (This PR adds 3 tests: the two ULP sweeps + the batching-contract identity, so passed counts sit higher.) **Deterministically fixed by this PR** (pristine main → PR, Marlin-on, 6 → the 2 rmsnorm reds gone): `fused_gate_up_swiglu_rmsnorm_is_bit_exact_to_two_step_path` and `fused_gate_up_swiglu_rmsnorm_zero_points_is_bit_exact_to_two_step_path` go red→green (the rmsnorm loop is genuinely capture-safe at M>1 and the assertion is narrowed to that routing). **Remaining failures are pre-existing and unrelated to this change:** - `fp16_gate_up_swiglu_is_bit_exact_to_two_op_path` — **Marlin split-K** artifact (plain path). Reds under Marlin-**on**, **passes under Marlin-off** on both main and this PR. This is the plain half we hard-stopped; reverted to main's `m == 1`. - `reference_scores_path_gates_on_dtype_and_decode_support` (GQA) — pre-existing, fails under **both** configs on main and PR (not a Marlin artifact; #1305 failure #4). - `graph::host_excursion_is_capturable_as_a_seam...`, `graph::mid_segment_capture_failure_is_recoverable_via_abort` — **flaky** capture tests (identical `graph` code to `main`; they intermittently red under both configs on both main and this PR — not attributable to this change). Closes #1334 (rmsnorm half). Plain half hard-stopped and tracked in #1334. Working as part of the batching workstream. `target/` left in place; CI not awaited per repo convention. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> --------- Co-authored-by: justinchuby <223556219+Copilot@users.noreply.github.com>

Make the bit-exactness test non-vacuous: assert INT8_PREFILL_GEBP_TEST_CALLS advances on the fused arm and stays put on the dequant arm, so a silent decline by the fused route can no longer make the comparison pass trivially. Mark the `p90` column on the nine re-measured 8-bit rows in the assignment matrix with a footnote: it is the per-trial maximum, not a p90, so no cell in that table is compared against a p90 as though the two were the same statistic. Record the review-time reproduction in the benchmark doc: the m >= 64 rows and the ORT loss-to-win crossover reproduce closely under load, while m = 2 and m = 4 read 13.9x/14.7x against 17.1x/18.0x. The small-m ratios are the host-sensitive ones because the run is ~1 ms and the fixed pack cost is a large fraction of it; the dequant arm reproduces to within 0.7%. State why the two routes are bit-identical rather than only asserting that they are: on native x86_64 + AVX2 the replaced route's `gemm` resolves to the same `sgemm_simd` packing and `6x16` microkernel the fused route drives, and that is the only configuration in which the fused route runs. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

justinchuby · 2026-08-19T05:34:08Z

Nits addressed in 61f0467bd. Per-nit disposition:

1. matmulnbits_int8_gebp_declines_decode_and_narrow_prefill was vacuous — fixed, and the fix found a second hole.
The old test asserted the counter did not advance for m = 1 and m = 3, which it would also not do if the whole route were deleted. Replaced with two tests that can each fail for exactly one reason:

…leaves_decode_on_the_gemv — m = 1 goes to the 8-bit decode GEMV, whose Nk weight is cached, so the counter must stay put and the values must still be right.
…row_gate_is_the_kernels_own_lower_bound — calls try_int8_prefill_gebp directly (via a new test_kernel_8bit, since the method is not reachable through Box<dyn Kernel>) and asserts false at m = 1, true at m = 2. This is the only way to test the constant, because m == 1 is claimed earlier in dispatch and the gate is unreachable from the production entry point.

4. Bit-exactness now has a test, and the reasoning is stronger than the doc claimed.
Added …is_bit_identical_to_the_dequant_route: exact assert_eq! on the full f32 output of both arms across (k,n,block) = (100,7,32), (96,20,32), (130,33,128), symmetric and asymmetric, m ∈ {2,7,19} — so partial trailing blocks and partial NR = 16 panels are both covered. It asserts the counter on both arms, so a silent decline cannot make it pass trivially (this is what I mean by the fix to nit 1 finding a second hole — I had written the comparison before adding the counter assertions, and it would have passed with the route disabled).

On the claim itself: I had justified it by reduction order, which is the weaker argument. The real reason is that on native x86_64 + AVX2 the replaced route’s gemm resolves to x86_sgemm::sgemm_simd — the same packing and the same 6x16 microkernel quant_prefill_gebp drives — and pack_b_quant emits exactly the bytes pack_b would emit from the dequantized weight. Since that is also the only configuration in which the fused route runs, the claim holds wherever it is made. Both docs now say this instead. It is still not a general statement about any two GEMMs, and the caveat list says so.

2. The p90 column is marked.
The nine re-measured rows now carry † and a footnote directly under the table: that column is the per-trial maximum (worst of 7–11 interleaved trials), because ab.py reports p50 [min-max]. Marked rather than converted, so the table stands alone and nobody compares a max to a p90. The prose bullet is trimmed to avoid repeating it.

3. Your reproduction is in the doc, under its own heading.
New section “The small-m ratios are the host-sensitive ones”: your 5.55x / 2.54x / 1.75x and 790 GFLOP/s against my 5.83x / 2.59x / 1.84x and 843, and your 13.9x / 14.7x at m = 2,4 against my 17.1x / 18.0x. The whole difference is in the fused arm’s absolute time (1.12 ms vs 0.905 ms at m = 2); your dequant arm reproduced mine to within 0.7%, which is the tell — at two rows the run is ~1 ms, so fixed pack and fork/join cost is a large fraction of it and any co-tenant moves the ratio. The doc now says to read small m as “between 14x and 18x depending on host load” and treat m ≥ 64 as the stable rows. The disposition is unaffected: the dequant arm cannot go below ~15 ms at this shape because that is what materializing 51 MB costs.

Gates re-run after the changes: cargo fmt --all, 1442 ep-cpu lib tests (19 int8_), clippy --all-targets -D warnings on x86_64 and aarch64 (the new test_kernel_8bit needed #[cfg(target_arch = "x86_64")] — the usual dead-code trap), verify_documented_env_vars.py, workspace_test_packages.py verify.

…ep comments (follow-up to #1404) (#1410) ## What this does Comment-only follow-up to #1404. Corrects two landed doc comments in `matmul_nbits.rs` (the plain `fp16_gate_up_swiglu_two_op_ulp_bound_sweep`) that mischaracterized the default plain-M>1 path. ### What was wrong The comments said plain M>1 "dispatches Marlin **split-K** (default-on), which is documented non-byte-identical and **possibly non-deterministic**." Both halves are wrong. Verified against `crates/onnx-runtime-ep-cuda/src/kernels/marlin_gemm.rs`: - **Split-K is default OFF.** It is gated by `ONNX_GENAI_MARLIN_SPLITK` (`marlin_splitk_enabled()`), separate from `ONNX_GENAI_MARLIN_M_GT_1`. `maybe_launch_marlin_splitk` elects it only when that flag is set *and* `choose_split_k > 1`; otherwise it calls the direct kernel. So split-K does not run under the default config. - **Split-K is deterministic** where it does run: "a fixed-order fp32 partial reduction that is NOT byte-identical to the single-block kernel (it stays within the f64-oracle tolerance and is **deterministic** — greedy/argmax tokens remain byte-identical, validated e2e on glm-4-9b and qwen2.5-14b)." - The actual default plain-M>1 path is the **direct** Marlin int4 tensor-core GEMM (`launch_marlin_gemm`), documented non-byte-identical to the tiled two-op reference by a different accumulation order (parity tests assert a `2e-2 * max_out` tolerance, not equality). ### What the corrected comments say They now name the correct path (direct Marlin, not split-K), drop "possibly non-deterministic", and record that `main` does **not** hold "M>1 is bit-exact to the two-op reference" under the default config — `fp16_gate_up_swiglu_is_bit_exact_to_two_op_path` reds at M=5 with Marlin default-on and passes only under `ONNX_GENAI_MARLIN_M_GT_1=0`. This aligns the code comments with the corrected framing already applied to PR #1404's body, issue #1405 (now titled "direct tensor-core GEMM, not split-K"), and the #1305 correction comment. ## Why it's safe Pure comment change — no code, no behavior change, no test change. The `///` / `//` edits do not touch any executable line. Follows-up #1404. Marlin write-up: #1405. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> Co-authored-by: justinchuby <223556219+Copilot@users.noreply.github.com>

) ## Problem `minijinja`'s built-in `tojson` filter is HTML-safe, so a chat template that renders tool schemas produced output that no HF/llama.cpp deployment ever produces: | difference | ours (before) | Transformers / llama.cpp | |---|---|---| | HTML escaping | `\u003cdir\u003e`, `\u0026\u0026` | `<dir>`, `&&` | | key order | sorted (`type` after `properties`) | insertion order | | separators | `{"a":1,"b":2}` | `{"a": 1, "b": 2}` | Agent clients (OpenCode, and anything else that ships shell/glob tools) put `<dir>` and `&&` all over their tool descriptions, so the model was handed a tool block full of escape sequences. ## Fix * Custom `tojson` filter matching Python's `json.dumps` semantics: no escaping, `", "` / `": "` separators, plus the `indent=` kwarg the built-in supported. * `preserve_order` on `serde_json` and `minijinja` so object keys keep their input order. * Tokenizer: the chat template already emits `bos_token`, and the tokenizer post-processor prepended a second one. The duplicate is now dropped (templates that do not emit BOS still get the post-processor's copy). ## Verification * A real agent request now renders **byte-for-byte identically** (23521 bytes) to llama.cpp's `POST /apply-template` for the same model and body. * Prompt token count against the served model matches the reference tokenization exactly (was +1 from the duplicated BOS). * `cargo test -p onnx-genai-ort --lib --tests` green; one existing assertion updated for the `": "` separator. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> Copilot-Session: 0190e2eb-abe4-451f-b36d-44a035a99b7e

) ## Why `effective_sliding_window` downgrades a metadata `sliding_window` to global attention when the decoder graph carries no GQA `local_window_size`, and the comments call such a window "vestigial", naming Muse-Glimmer-30B as the example. That diagnosis is wrong. Muse-Glimmer-30B trains with a 2048-token window on 39 of its 52 layers (the other 13 are NoPE full-attention). Its ONNX export simply omits `local_window_size`, and running those 39 layers globally corrupts generation once the prompt outgrows 2048 tokens. Measured on a real agent request (5.3k prompt tokens, tools): | decoder graph | result | |---|---| | as exported (no `local_window_size`) | 8189 reasoning tokens, 0 output tokens, `finish_reason=length`, ~390 s | | same graph, `local_window_size=2048` patched into the 39 sliding layers | 731 / 844 reasoning tokens, correct answer, 29-77 s; a tool call now lands in 5.3 s | llama.cpp on the same prompt and weights converges in 536-688 tokens, matching the patched behaviour. ## What changes Behaviour is unchanged: the runtime cannot re-apply a mask the export left out, so the global-attention fallback stays. But the mismatch is now `warn!` instead of `debug!`, the message explains that generation will diverge past the window if the architecture really is windowed, and the comments no longer teach that this model's window is decorative. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> Copilot-Session: 0190e2eb-abe4-451f-b36d-44a035a99b7e

A flattened `channels_first` patch carries both a channel axis and a temporal axis, and models disagree about which one is outer. Qwen2-VL repeats each frame inside its channel block, giving `[C, T, H, W]`, and the packer hardcoded that. Muse Glimmer keeps whole frames contiguous instead, giving `[T, C, H, W]`. The two layouts have identical shapes, so nothing downstream complains; the patch projection just reads colour through the wrong weights. Spatial structure survives, because every frame still holds a complete image, which makes the failure look like anything but a layout bug. ## Evidence Muse Glimmer 30B INT4, native CUDA, before: | image | model says | |---|---| | solid red | `yellow` | | solid green | `purple` | | solid blue | `teal` | | blue circle above red square | "a teal oval above an olive-green rectangle" | llama.cpp on the same checkpoint answered `red`, `green`, `blue`, and "a blue circle and a red square". Dequantising the model's patch embedding and projecting a known pixel both ways isolated it: `[T, C, H, W]` scored **0.994** cosine against the reference weights, `[C, T, H, W]` scored **-0.918** — near perfect anticorrelation, exactly the inversion the model was describing. After adding `temporal_order: temporal_major` to the package metadata, the same server answers `red`, `green`, `blue`, and "A blue oval/ellipse sits above a red rectangle." ## Compatibility The field defaults to `channel_major`, so existing packages keep the previous behaviour. A test asserts the default explicitly. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> Copilot-Session: 0190e2eb-abe4-451f-b36d-44a035a99b7e

## The problem `serve` answers the same four questions as `generate`, `run` and `transcribe` -- which backend, which device, how much VRAM, how much host RAM -- in a different vocabulary. Before this change: | flag | generate | run | transcribe | serve | |---|---|---|---|---| | `--backend` | yes | yes | yes | **no** | | `--device` | yes | yes | yes | **no** (spelled `--native-device`) | | `--host-ram-limit` | yes | yes | yes | **no** | | `--cpu-cores` | yes | yes | yes | **no** | | model positionally | yes | yes | n/a | **no** (`--model` only) | So a command line learned from one subcommand is rejected by another. `onnx-genai serve MODEL --backend native --device cuda:6` -- the obvious extrapolation from `run` -- fails twice in a row, once on the positional model and once on `--backend`, and clap's usage text lists neither. The flag that does work, `--native-device`, appears in no other subcommand's help, so there is nowhere to discover it from. ## The change The flags now live once, in `onnx-genai-server::runtime_args`, and are flattened into all four subcommands. That crate is the lowest layer both the unified CLI and the standalone server binary already depend on, and it already carries clap with the `env` feature. A flag added there appears everywhere with the same name, grammar and help text, so the surfaces cannot drift apart again. Duplicated definitions removed along the way: `EngineArgs`, `CpuArgs`, the device enum, and the `--backend`, `--device` and memory-limit parsers each existed twice, and `parse_native_device` had two implementations that disagreed about whether `auto` and `gpu` were valid. ## Behaviour changes Each has a test: - `serve` accepts the model positionally. Both spellings share the model-source arg group, so giving it twice is refused while parsing rather than after the process has started. - Naming a `--device` implies the native backend on every subcommand. `serve` already worked this way through `--native-device`; previously `generate --device cuda:0` parsed, ran on ORT and ignored the device. - `--vram-limit` and friends gain the env fallbacks `serve` already had. - `--model-id` alongside `--models-dir`/`--models-config` is an error instead of being silently ignored, which hid typo'd invocations. `--native-device` keeps working and is hidden from help, so existing service units are unaffected. `--device` wins if both are given. ## Verification - `cargo test -p onnx-genai-server --lib --features native-backend` — 16 tests across `cli::` and `runtime_args::`, all passing; full lib 237 passed. - `cargo test -p onnx-genai-cli --lib --features native-cuda` — 118 passed. - `cargo fmt --all --check` and `cargo clippy` clean. - End to end on a real native CUDA model: `onnx-genai serve ~/models/... --backend native --device cuda:6 --addr 127.0.0.1:18082` loads and serves a chat completion, which is the invocation that failed before. --------- Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> Copilot-Session: 0190e2eb-abe4-451f-b36d-44a035a99b7e

…#1413) ## Symptom Sending a prompt that references an image without attaching one answered with the raw engine error: ``` missing required pipeline input 'encoder.pixel_values' ``` instead of the advice the REPL is supposed to give ("this model declares image input; attach one with `/image <path>`"). Two `repl_e2e` tests were failing on `main` because of it. ## Cause The front end recognized the condition by matching `"input not found: "` — the wording a *single decoder graph* raises. A *pipeline* raises its own wording from `routing.rs`, so once a multimodal package started failing through the pipeline router the match stopped firing. The unit test guarding this hand-wrote the same string it tested against, so the test and the matcher agreed with each other while neither agreed with the engine. That is why the drift was invisible. ## Fix Name the message once, next to where the pipeline writes it (`MISSING_REQUIRED_INPUT`), and export the recognizer (`is_missing_required_input`) from the engine. A front end now asks the engine whether an error is a missing input rather than guessing at its prose, and the message and its recognizer cannot drift apart because they share the constant. ## Tests The test moves to the engine, next to the constant, and additionally pins the case that must **not** be advised as a forgotten attachment: an optional input that a presence key declared present is a package-authoring fault, not a missing file. `cargo test -p onnx-genai-cli -p onnx-genai-engine` — the two previously failing `repl_e2e` tests pass (47/47). clippy and fmt clean. Note: `decode_position_and_state_e2e::pipeline_load_rejects_kv_page_pool_before_fixed_state_when_host_budget_cannot_fit_either` also fails on `main`, from an unrelated change to how the KV budget resolves under a tiny host limit. It is not touched here. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> Copilot-Session: 0190e2eb-abe4-451f-b36d-44a035a99b7e

…1414) ## What `decode_position_and_state_e2e::pipeline_load_rejects_kv_page_pool_before_fixed_state_when_host_budget_cannot_fit_either` was failing on `main`. The invariant it guards still holds — the wording it matched no longer exists. ## Why it failed A 15-byte host budget is now caught while the **shared KV budget is being resolved**, rather than later when the pool is allocated. The refusal that comes out is the better of the two: ``` cannot satisfy lowered resource limit: requested 15 B, but at least 256 B is required; the remaining 15 B cannot hold one 256 B KV page; raise the limit to at least 256 B ``` versus the old `cannot allocate the pipeline KV page pool ... 1 page(s) across 2 layer(s)`. It names the shortfall, the model's own floor, and the limit that would clear it. So this is a diagnostic improvement whose test was left behind, not a regression. ## Change Assert against the current refusal, read from the **whole error chain** (`{error:?}`) rather than the top-level sentence — the reason a load was refused is the cause, not the context wrapped around it. That is also what kept the failure opaque: `error.to_string()` showed only "failed to resolve the shared pipeline KV memory budget". The negative assertion is unchanged: fixed-state admission must still not be reached, which is the ordering the test exists for. `cargo test -p onnx-genai-engine --test decode_position_and_state_e2e` — 3/3. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> Copilot-Session: 0190e2eb-abe4-451f-b36d-44a035a99b7e

Reasoning models pick their own effort when the chat template receives no `reasoning_effort`, and that self-chosen default is often the maximum. Agent clients are the problem case: OpenCode and pi never send the field on their main request, so every turn runs at maximum effort. On a 5.3k-token agent prompt this model averages hundreds of thinking tokens per turn and, often enough to matter, spends the entire 8192-token budget thinking and returns `finish_reason: "length"` with no content at all -- the operator sees a client that hangs for minutes and prints nothing. `--default-reasoning-effort` (env `ONNX_GENAI_DEFAULT_REASONING_EFFORT`) gives the operator a floor for clients that stay silent. A request that does send `reasoning_effort` always wins, so this can never override a client that has an opinion, and leaving the flag unset keeps today's behaviour exactly: the model's own default stays in place. Measured against a captured OpenCode request (5273 prompt tokens, 6 tools) on Muse Glimmer 30B int4: at the template's default effort a "say hi" turn emits ~900 reasoning chunks and no answer; at `low` the same turn answers in a few seconds with ~20-80 reasoning chunks. Sampling variance means `low` is not a guarantee -- the runaway is a sampler behaviour, not an effort bug -- but it moves the common case from unusable to usable. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> Copilot-Session: 0190e2eb-abe4-451f-b36d-44a035a99b7e Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> Copilot-Session: 0190e2eb-abe4-451f-b36d-44a035a99b7e

…1475) Follow-up to #1473, which shipped query-axis prefill padding together with a raise of `DEFAULT_VARIANTS_PER_NODE` from 4 to 10 so the bound would cover the eight-step ladder. That raise was a regression. A retained kernel variant owns device scratch, and the per-node bound is the only thing capping how much of it accumulates — widening it defeats its own purpose. Isolated on `muse-glimmer-30b-int4`: | configuration | pass-2 recompiles | 12-prompt memory sweep | | --- | --- | --- | | no padding, bound 4 (pre-#1473) | ~888 per forward | 12/12 ok, 39–53 GB | | padding (8 steps), bound 4 | ~888 per forward | — | | padding (8 steps), bound 10 (#1473) | 0 | **fails at 5.5k tokens, 72.8 GB** | | padding (3 steps), bound 4 (this PR) | **0** | **12/12 ok, flat ~50 GB** | So the ladder is sized to the bound that exists rather than the other way round: three prefill widths (`{171, 342, 512}` for a 512-wide chunk), leaving exactly one slot for the single-token decode shape. The step is rounded up so the top step lands on the chunk width instead of just short of it and leaving a stray fourth width above — there is a test for that invariant across a range of declared chunk widths. With no environment overrides, the cache now holds the same 5651 entries it held before padding existed, evictions freeze at 364, and greedy output is byte-identical with padding on and off. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> Copilot-Session: 0190e2eb-abe4-451f-b36d-44a035a99b7e

Follow-up to #1473 / #1475, which landed the explanation as `docs/performance/CHUNKED_PREFILL.md`. It is explanatory writing for people — why prefill can be split into chunks at all, and why the leftover final chunk had to be padded onto a fixed ladder — rather than a specification or an evidence record, so it belongs in `wiki/`. Rewritten to the wiki's conventions from `wiki/README.md`: - YAML frontmatter with `title`, `aliases`, `tags`, `status`, `created`, `updated` - Opens with a `> [!summary]` stating the question the note answers - Callouts for the invariants (`KV cache 变长,不改变任何 kernel 的 input shape`) and for the two cases that refuse padding - `[[wikilinks]]` to `memory/Virtual Memory for KV Cache`, `memory/Memory Management for Beginners`, `execution/CUDA Execution Provider`, `metadata/Metadata Driven Runtime`, and `performance/Performance Engineering Playbook` instead of restating them - Body in Chinese to match its audience, following the precedent of the KV cache and chat template notes; English filename and title so links survive Listed in both maps of content (`wiki/README.md` and `wiki/index.md`), with a one-line description in the index alongside the other narrative notes. All wikilinks in the new note resolve to existing files. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> Copilot-Session: 0190e2eb-abe4-451f-b36d-44a035a99b7e

…et (#1177) (#1477) ## Summary The CUDA test-honesty gate (`CUDA compile (Linux x86_64)` → *Verify CUDA test inventory and skip honesty*) has been **red on `main`** independent of any PR (reproduced at `84a27653`, observed on a rustfmt-only PR). Root cause per #1177: `crates/onnx-runtime-ep-cuda/tests/matmul_nbits_marlin_numerics.rs` is a **mixed target** — five pure-CPU oracle self-checks pass on the CPU lane alongside three CUDA tests that correctly `ignore` without `gpu-tests`. The honesty checker sees CUDA-target tests passing without a GPU and correctly objects. **Diagnosis verified before acting** (per the issue's request): the five CPU tests (`oracle_matches_independent_reference_symmetric`/`_asymmetric`, `oracle_is_exact_on_a_hand_checkable_case`, `envelope_scales_with_output_magnitude_and_has_a_floor`, `parity_flags_a_perturbed_candidate`) issue **no CUDA calls** — they only touch the device-free oracle/envelope helpers. The three GPU tests (`current_path_matches_f64_oracle_group_size_sweep`, `..._projection_shapes`, `fp16_mixed_gemv_matches_f64_oracle_glm_decode`) exclusively drive `run_matmul_nbits_f16`/`maybe_cuda`. The split is clean. ## Fix — split, don't weaken the checker - New shared **non-target** module `tests/marlin_numerics/mod.rs` holds the device-free machinery (`Int4Problem`, the f64 oracle, `Envelope`/`ParityReport`, `f32_dequant_reference`, `GROUP_SIZES`). - `matmul_nbits_marlin_numerics.rs` is now **purely-CUDA**: the 3 GPU tests + the `run_matmul_nbits_f16`/`maybe_cuda` driver, all ignored without `gpu-tests`. - New `matmul_nbits_marlin_oracle.rs` holds the 5 CPU self-checks. - `matmul_nbits_marlin_oracle` added to the checker's `ALWAYS_RUN` set (documented as a genuine CPU-only probe), and run explicitly on the CPU lane via a new CI step so the oracle math stays exercised. - **`verify_cuda_test_honesty.py` pass/fail logic is untouched.** ## Verification (both directions, CPU/script-only — GPU left to the agent using it) - **Honesty base-config phase passes** on the fixed tree for every target: `matmul_nbits_marlin_numerics` = **0 passed / 0 failed / 3 ignored** (now policed & clean); `matmul_nbits_marlin_oracle` correctly **exempt**. (The checker's GPU-execution phase is by design a no-CUDA-host check; the base phase is the #1177-relevant half and was validated in isolation to avoid competing for the busy GPU.) - **Oracle target runs & passes in its new home**: `cargo test -p onnx-runtime-ep-cuda --features cuda --test matmul_nbits_marlin_oracle` → **5 passed / 0 failed / 0 ignored**. - **CUDA tests stay ignored**: `--test matmul_nbits_marlin_numerics` → **0 passed / 0 failed / 3 ignored**. - Both targets compile under `--features cuda,gpu-tests` (inventory reconciliation). - Checker still has teeth: it rejects a simulated pre-split shape (numerics target passing 5 tests without gpu-tests → *"CUDA tests must be ignored, not pass"*) and accepts the post-split shape. Self-test passes. - My three files are rustfmt-clean and clippy-clean (pre-existing `optimizer.rs`/lib lints from a newer local clippy are unrelated and out of scope). Hardware: i7-13800H / RTX 4060 Laptop, CUDA 13.1. No GPU tests were run. Closes #1177 Co-authored-by: justinchuby <223556219+Copilot@users.noreply.github.com>

## The symptom Every OpenCode session opened with a failed tool call: ``` ✗ Invalid Tool Model tried to call unavailable tool 'read.filePath' ``` The agent recovered on the retry, so the session worked, but every task paid for a wasted round trip. ## The root cause is in the chat template `chat_template.jinja` infers a namespace with `fn.name.split('.')[0]` and renders it unconditionally as a pattern. OpenCode offers bare tool names, so the system prompt became: ``` # Valid recipients: "self", "read.*", "bash.*", "glob.*", "user". ``` The model was told it must address `read.*`. It supplied a second segment by borrowing something at hand -- typically a parameter name. `read.filePath` is the model doing exactly what the prompt asked. That template lives in the model package, and it has been fixed there. This PR is the server-side half. ## The change `align_tool_calls` already rewrote qualified names, but only tried the **last** segment. That handles `functions.read`; it cannot handle `read.filePath`, where the tool is the **first** segment. Now every segment is tried, and the rewrite happens only when the offered tools agree on exactly one target. The candidate set is deduplicated so that: - `glob.glob` resolves to `glob` (both segments name the same tool), - `read.write` is left alone (two different offered tools -- no way to tell which was meant). ## Verification - `cargo test -p onnx-genai-server --lib` -- 239 passed, 0 failed. `tool_name_alignment_tests` is 7/7 with two new cases covering the shapes above. - End to end against muse-glimmer-30b-int4 with OpenCode: a read task and a multi-tool debug task (read, edit, bash, todowrite) both completed with no invalid tool errors. CI runners are down; validated locally. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> Copilot-Session: 0190e2eb-abe4-451f-b36d-44a035a99b7e

codecov · 2026-08-19T16:01:50Z

Codecov Report

❌ Patch coverage is 90.59952% with 196 lines in your changes missing coverage. Please review.
✅ Project coverage is 80.67%. Comparing base (f14e307) to head (f2a982d).

Files with missing lines	Patch %	Lines
...es/onnx-runtime-ep-cpu/src/kernels/matmul_nbits.rs	92.31%	57 Missing and 6 partials ⚠️
...rates/onnx-runtime-ep-cpu/src/kernels/x86_sgemm.rs	79.18%	61 Missing ⚠️
crates/onnx-runtime-ep-cpu/src/dispatch_ledger.rs	92.04%	28 Missing and 7 partials ⚠️
crates/onnx-runtime-ep-cpu/src/kernels/matmul.rs	92.72%	9 Missing and 3 partials ⚠️
crates/onnx-runtime-ep-cpu/src/kernels/softmax.rs	60.00%	7 Missing and 1 partial ⚠️
...nnx-runtime-ep-cpu/src/kernels/simd_activations.rs	50.00%	6 Missing ⚠️
crates/onnx-runtime-ep-cpu/src/backend_ab.rs	97.29%	1 Missing and 2 partials ⚠️
crates/onnx-genai-metadata/src/parser.rs	96.00%	1 Missing and 1 partial ⚠️
crates/onnx-runtime-ep-cpu/src/task_runtime/mod.rs	92.00%	1 Missing and 1 partial ⚠️
crates/onnx-genai-cli/src/generate.rs	93.33%	0 Missing and 1 partial ⚠️
... and 3 more

Additional details and impacted files

@@                       Coverage Diff                       @@
##           squad/roy-int4-prefill-gebp    #1403      +/-   ##
===============================================================
+ Coverage                        80.08%   80.67%   +0.59%     
===============================================================
  Files                              362      377      +15     
  Lines                           158227   165621    +7394     
  Branches                        158227   165621    +7394     
===============================================================
+ Hits                            126709   133621    +6912     
- Misses                           26876    27161     +285     
- Partials                          4642     4839     +197

Flag	Coverage Δ
cli-ort-linux	`82.60% <95.23%> (?)`
cli-ort-windows	`82.19% <95.23%> (?)`
offline	`80.61% <90.45%> (+0.53%)`	⬆️

Flags with carried forward coverage won't be shown. Click here to find out more.

Files with missing lines	Coverage Δ
crates/onnx-genai-cli/src/commands.rs	`70.13% <ø> (ø)`
crates/onnx-genai-cli/src/lib.rs	`93.56% <100.00%> (ø)`
crates/onnx-genai-metadata/src/lib.rs	`95.23% <ø> (ø)`
crates/onnx-genai-metadata/src/schema/pipeline.rs	`71.42% <ø> (ø)`
crates/onnx-genai-preprocess/src/image/packed.rs	`72.80% <100.00%> (+0.23%)`	⬆️
crates/onnx-genai-preprocess/src/image/program.rs	`78.13% <100.00%> (+0.16%)`	⬆️
crates/onnx-runtime-cost-model/src/model.rs	`81.46% <100.00%> (ø)`
crates/onnx-runtime-ep-plugin/src/compute.rs	`80.98% <ø> (+1.05%)`	⬆️
...rates/onnx-runtime-ep-plugin/src/dispatch_probe.rs	`68.71% <ø> (ø)`
crates/onnx-runtime-ep-plugin/src/factory.rs	`69.31% <ø> (+0.24%)`	⬆️
... and 22 more

... and 28 files with indirect coverage changes

🚀 New features to boost your workflow:

❄️ Test Analytics: Detect flaky tests, report on failures, and find test suite problems.
📦 JS Bundle Analysis: Save yourself from yourself by tracking and limiting bundle sizes in JS merges.

## Summary Owner-mandated upgrade of the pinned ONNX Runtime from **1.27.0 → 1.28.0** so that native-vs-ORT performance claims are measured against the current ORT rather than a stale baseline. - `crates/onnx-genai-ort/ort-sys/build.rs`: `ORT_VERSION="1.28.0"`, `ORT_API_VERSION="28"`, all four release-archive names + SHA-256 checksums refreshed to the 1.28.0 artifacts. - `ort-sys/src/lib.rs`: platform dylib names (`libonnxruntime.1.28.0.dylib`, `libonnxruntime.so.1.28.0`) and diagnostic version strings. - `src/error.rs`, `src/value.rs`: version strings in diagnostics/doc examples. - `onnxruntime>=1.28,<2` in the CLI/python `pyproject.toml`s and README. - Adds `scripts/ort_ab/ort_cuda_decode_bench.py`: a direct ORT-CUDA io-binding greedy-decode driver used to verify output parity across the bump. ## Cost of the bump **Zero source breakage.** `cargo build --release -p onnx-genai-bench --bin profile_native --features "bench-native,bench-ort,cuda"` succeeds; ort-sys recompiles and bindgen regenerates against the API-28 headers with no changes required beyond the version constants. ORT's C-ABI back-compat held. ## Verification (RTX 4060, CUDA 13.1, ORT 1.28.0 / API 28) - **Provider genuinely CUDA on 1.28**, verified not assumed (loaded-library log reports `1.28.0 (API 28)`; `ort_available_execution_providers` includes `CUDAExecutionProvider`; device memory in use). - **Loading recipe re-verified from scratch on 1.28:** `ONNX_GENAI_ORT_LIB_DIR`=anaconda `capi`, `PATH` prepend `capi;<nvidia>\cu13\bin\x86_64`, `ONNX_GENAI_EP_FALLBACK=1`. On 1.28 cuDNN/cuFFT are optional at runtime and nvrtc is no longer linked — the CUDA EP now instantiates with `capi`+`cu13` alone (cuDNN still recommended for fused-attention perf). The cu13 runtime bin remains **required** (missing it silently drops the provider to CPU — the `cublasLt64_13.dll` trap). - **granite-1b-a400m-f16-mobius, ORT-CUDA:** greedy decode is **byte-identical** on 1.28 vs the stored 1.27 evidence (33/33 shared tokens). Native decode is ORT-version-independent, so native↔ORT byte-identity is preserved by transitivity. Correctness is preserved across the bump. - **qwen05b-main, native graphs-on, 1.28 binary:** 452 tok/s, CUDA-graph capture succeeds (`captures=3 fallbacks=0`). A same-harness 1.27-vs-1.28 throughput delta and a valid qwen05b ORT arm were **not** obtainable on this box (qwen05b's GroupQueryAttention needs onnxruntime-genai input prep that a raw `InferenceSession` lacks, and oga 0.14.1 is a CUDA-12 build that fails to load against CUDA 13.1). This is reported plainly rather than forcing a number the harness cannot honestly produce; full detail is in the decision log (UPDATE 17). Fused-QMoE note: ORT 1.28 loads the Mobius fused model on CUDA, but native still cannot load it (fc1 QMoE-weight-wiring blocker in our loader, ORT-version-independent) and it would still decline CUDA-graph capture (`ai.onnx::Attention-24` growing KV) — the fused-MoE emission line is unaffected by the ORT version. 🤖 Flagged as **needs review** — a version-pin bump with checksum changes should have a squad member confirm the four SHA-256 values against the upstream 1.28.0 release before merge. Co-authored-by: justinchuby <223556219+Copilot@users.noreply.github.com>

A leaked benchmark process spun a full core on the dev box for **eight hours** — `profile_native` pid 9276, 01:00:52 to 09:05, ~**5.6 CPU-hours** — while casual checks reported the machine free. Every wall-clock figure taken in that window was measured a core short. **Both obvious liveness probes fail, in opposite directions, and the guidance circulating during the session named only the harmless one:** | probe | failure | consequence | |---|---|---| | `nvidia-smi`, WMI | retains **post-mortem** rows → false **positive** | looks busy when idle; wastes time | | `Get-Process -Id <pid>` | returned **absent for that live pid**; `Stop-Process -Id` also failed → false **negative** | looks idle when busy; **invites starting a timing run on a contended box** | The false negative is the dangerous one, and it is the one we were telling each other to rely on. **What actually worked**: a repeated `Get-Process -Name <name>` snapshot, confirmed by a **CPU-time delta over a wall interval** — that process gained **5.55 s of CPU in 6 s** of wall time. A process whose CPU time advances is alive whatever any single probe says. Terminating it required `Get-CimInstance Win32_Process -Filter "ProcessId = <pid>"` piped to `Invoke-CimMethod -MethodName Terminate`. Also recorded: - **`0 %` GPU utilisation is not evidence of an idle box.** A live process holding a CUDA context sits at 0 % between kernels, and a spinning host-side loop consumes a core while showing no GPU activity at all. Two agents independently read `0 %` as "free" during this window. - **Prefer contention-invariant counters where the question allows** — CUDA-graph segment counts, `htod_bytes`, page-in and eviction counts, crash/no-crash ratios, token byte-identity. Several conclusions in this corpus rest on them precisely because a shared box cannot be trusted, and those conclusions are unaffected by this episode. Sibling to §40 (a threshold calibrated in one regime misfiring in another) and §41 (a shared primitive's contract change breaking distant callers): all three are failure modes that produce *confident, wrong* results rather than obvious errors. Docs-only. The measurements quoted are from this box today and the affected throughput figures are being re-taken now that it is genuinely quiet. Co-authored-by: justinchuby <223556219+Copilot@users.noreply.github.com> Copilot-Session: d60eb808-7cc6-4abc-b48d-2a6dd3841624

## What Unbreak the two CI lanes that are **red on `main` right now**. Three independent breakages, all pre-existing and all reproduced on an unmodified `dbade34c1` checkout with the same stable 1.97.1 toolchain CI installs: | # | gate | breakage | fix | |---|------|----------|-----| | 1 | `cargo fmt --all -- --check` | 5 sites / 4 files | rustfmt | | 2 | `Rust quality` clippy | `clippy::unnecessary_map_or` — `dispatch.rs:26` | `map_or(true, f)` → `is_none_or(f)` | | 3 | `Fast (Linux x86_64)` clippy `--all-targets` | `clippy::inconsistent_digit_grouping` — `cost-model/model.rs:314` | `2_000_000_000_000_0` → `20_000_000_000_000` | Both lanes build with `RUSTFLAGS: -D warnings`, so #2 and #3 are hard errors, not warnings. **Every PR that merges `main` inherits all three** — verified on #1434 and #1420. Nothing in the queue can go green until this lands. #3 is worth calling out: it is invisible to a plain `cargo clippy` because the literal lives in a `#[cfg(test)]` module. Only the `--all-targets` invocation in the Fast lane sees it. ## Semantics Both non-fmt changes are provably value-preserving: - `is_none_or(f)` is the rewrite the lint itself suggests, and is definitionally `map_or(true, f)`: `None` → `true`, `Some(v)` → `f(v)`. `gqa_shape_capacity_bound_enabled()` is unchanged — unset stays enabled, the falsey spellings stay disabled. - `20_000_000_000_000 == 2_000_000_000_000_0` (both 2e13), which is what the test's own comment already claims — *"2e13 FLOP / 2e13 = 1 s"*. `op_cost_takes_roofline_max` still asserts the compute term dominates. ## Validation Ran locally per the delayed-Actions directive, on this head merged with `origin/main` @ `dbade34c1`: | gate | result | |------|--------| | `cargo fmt --all -- --check` | **0 diffs** | | `cargo clippy --locked --all-targets $(workspace_test_packages.py cargo-args offline-linux) -- -D warnings` | **exit 0** | | `cargo test --locked $(… offline-linux)` | **3943 passed, 0 failed, exit 0** | | `scripts/check_cross_compile.sh` | **PASS** — x86_64 + aarch64 full offline set | | `benchmark_muse_native_local.py --self-test --require-numpy` | 43 cases passed | | `check_publish_order.py` / `check_profile_table.py` / `check_platform_naming.py` | PASS | | `check_dispatch_reachability.py` / `check_feature_gate_coverage.py` | PASS | | `check_dispatch_manifest.py` (`--self-test` and plain) | PASS | | `workspace_test_packages.py verify` | PASS | | `verify_documented_env_vars.py` | PASS — 113 documented, 13 known-unimplemented | | MLAS cfg: `-p onnx-runtime-ep-cpu --no-default-features --features mlas` | `moe::` 19 passed · `qlinear_matmul::` 30 passed · `optimization_registry_excludes_nchwc_without_cnn_ops` 1 passed | | `cargo clippy -p onnx-genai-engine --features native-backend` | clean | | `cargo build -p onnx-runtime-ep-cpu-plugin --features mlas` | clean | ### Windows ARM64: not validated locally — stated as a blocker, then bounded I could not run `Rust (Windows ARM64)` here and I am **not** claiming it as a pass. Two routes were attempted, both fail *identically on unmodified `main`*, so neither can discriminate this PR from baseline: - **`cargo-xwin` / clang-cl** — installed, MSVC CRT + SDK downloaded, correctly targeting `aarch64-pc-windows-msvc`. Fails in vendored `mlasi.h` on NEON intrinsics (`veorq_s32`, `vdupq_n_f32`, …) that MSVC supplies but clang-cl in MSVC mode does not. - **`aarch64-unknown-linux-gnu` + GNU cross toolchain** as an ARM64-NEON proxy — gets much further, compiles most of the ARM64 MLAS source set, then fails on `activate_fp16.cpp`. What makes this safe to merge anyway is **dependency-graph disjointness**, not a judgement call. That lane builds only `mlas-sys` and `onnx-runtime-ep-cpu-plugin --features mlas`. This PR touches `onnx-genai-engine` (tests), `onnx-runtime-ep-cuda` and `onnx-runtime-session`: ``` $ cargo tree -p onnx-runtime-ep-cpu-plugin --features mlas -e normal --prefix none \ | sort -u | grep -cE "onnx-runtime-ep-cuda|onnx-runtime-session|onnx-genai-engine" 0 ``` Zero of the crates this PR modifies are in that lane's graph, so it cannot observe this change. --------- Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

…e it (#1407) ## The UB `nested_dispatch_slot_pressure` (which I added in #1377) reconstructs a `&mut [u64]` over the **entire** row in every task and only then narrows: ```rust let row = unsafe { std::slice::from_raw_parts_mut(base as *mut u64, NEST_INNER) }; for slot in &mut row[start..end] { *slot += 1; } ``` The *stores* are disjoint. The *retags* are not. Under Stacked Borrows the violation is the retag: each task pushes a `Unique` covering all of `NEST_INNER`, popping the previous task's tag. Narrowing before the retag fixes it, and is the shape `parallel_output_rows_repeated` already uses in production. **Diagnosis and fix are Pris's, from #1385.** This PR carries them because that one is still a draft with auto-merge off and the UB is on `main` today. If #1385 goes ready first, close this and I will rebase the Miri half onto it — the two halves are independent. ## Why it survived, which is the part worth keeping CI's Miri lane runs `-p onnx-runtime-ep-cpu --lib task_runtime::`. **Integration tests under `tests/` are never Miri-checked at all.** Fixing this one instance would have left that gap open for the next one, so this also puts the shape under the lane as a lib test. That was not sufficient either, and the first attempt is the useful part: **the new test passed under Miri with the bad retag still in it.** Miri reports `available_parallelism() == 1`, so `resolve_width` builds a one-lane pool, every fan-out returns `Backend::Serial`, and no two tasks ever run against each other. Probe under Miri: ``` pool_width=1 backend=Serial tasks=1 ``` The lane has been type-checking the unsafe blocks in this module without exercising the concurrency they exist for. The workflow comment claims it "runs real threads under Stacked Borrows" — it runs one. So I would have shipped a canary that passes for a reason unrelated to its claim, which is precisely the class Pris named in #1385. Two fixes: 1. **The test asserts it actually fanned out** (`Backend::Native`), so it fails loudly if it ever degenerates to one task instead of passing silently. 2. **A dedicated lane step with `-Zmiri-num-cpus=4`**, scoped to that single test rather than all of `task_runtime::` — multi-CPU Miri multiplies the runtime of what is already the slowest step in the lane. The targeted step costs **18s**. ## Verified in both directions | retag shape | `-Zmiri-num-cpus=4` | result | | --- | --- | --- | | whole row (**#1377 as merged**) | yes | `error: Undefined Behavior: Data race detected between (1) retag write on thread task_runtime::t and (2) retag write of type [u64] on thread nxrt-task-0` | | whole row | **no** (lane default) | **passes** | | own range (**this PR**) | yes | passes | The middle row is the finding: without the flag the falsifier does not falsify. ## Scope Test-only plus one workflow step. No production change — `for_each_chunk_mut` was always correct. 30/30 `task_runtime::` lib tests pass natively and under Miri; the repaired benchmark still runs (`1 of 30 dispatches declined`); fmt and clippy clean. Includes the one-line rustfmt repair of `onnx-runtime-ep-cuda/src/runtime.rs` that `main` is currently failing on (same as #1393/#1395/#1398 — whichever lands first makes the rest a no-op). 🤖 Generated with [GitHub Copilot CLI](https://github.com/features/copilot/cli) --------- Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

…ce (#1173) The production/default CPU EP is pure native Rust: no proactive, default or runtime MLAS, and no ONNX Runtime CPU EP fallback. MLAS stays in the tree as an explicit, non-default research/test/benchmark reference that native kernels are measured against and progressively absorb. - Capability ledger (`dispatch_ledger.rs`), off by default; the three production dispatch sites pay one relaxed atomic load via lazy `record_with`. - A/B surface, differential test (10 tests, 2 MLAS-only) and interleaved A/B benchmark, all non-default. - Seven falsifiers in `default_artifacts_are_mlas_free.rs` prove the shipped plugin, wheel and workspace neither enable, link nor export MLAS: a workspace-wide recursive default-feature sweep, a resolved-graph check, an `nm` symbol check, a wheel-policy check and a manifest-prose check, each with its own vacuity guard. - Wheel MLAS polarity inverted to opt-in (`NXRT_EP_CPU_RESEARCH_MLAS=1`). - `mlas-sys` assembles the GAS aarch64 kernels so the research build can link on non-MSVC ARM64 at all. Validated locally at merged head 50b0d02 (Linux x86-64): fmt clean; clippy -D warnings clean in default and `--features mlas`; ep-cpu 1469 (default) and 1493 (mlas) lib tests pass; mlas-sys 41+2+1+5 pass; plugin suites green with falsifiers 7/7; full-scope `scripts/check_cross_compile.sh` green for x86_64-unknown-linux-gnu and aarch64-unknown-linux-gnu (full offline set, no reduced-scope fallback). Symbol falsifier proven non-vacuous: default cdylib 0/59,736 MLAS symbols, `--features mlas` cdylib 609/63,452. Reviewed-by: Gaff (code review / quality) Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

把整个 wiki 翻译成简体中文,并给每一页加上 `lang` frontmatter 字段。 ## 规则 `wiki/README.md` 新增「语言」章节,把规则写死: **译为中文**:正文、章节标题(H2 及以下)、表格内容、callout 标题文字、列表项。 **保持英文**:文件名、frontmatter `title`、页面 H1、`tags`、**`[[wikilink]]` 的目标**、代码块与标识符、crate 名、路径、环境变量、Obsidian callout 类型关键字。标题保持英文是为了让链接和跨语言导航保持稳定 —— 需要中文显示时用 `[[目标|显示文本]]`,竖线左边不动。 ## `lang` 是真的会生效的不是装饰性字段。Quartz 在 `renderPage.tsx` 里读 `frontmatter?.lang` 直接填进 `<html lang>`;缺失时回落到 `quartz.config.yaml` 的 `locale`。因此本 PR 同时把 `locale` 从 `en-US` 改为 `zh-CN`,让 Quartz 生成的页面(标签页、目录页、404)和它自身的界面文案(搜索、目录、最后更新)也跟随内容语言,而不是回落到英文。 ## 顺带修正的三处不一致 1. **三个早先就用中文写的页面,H1 是中文而 `title` 是英文**,两者不一致。H1 已改回英文,中文名保留在 `aliases` 里,检索入口不丢。 2. **`meta/Using this Wiki.md` 里的笔记模板正文仍是英文**。模板是作者真正复制的东西 —— 只改规则散文而不改模板,新笔记会一路默认回英文。已一并翻译。 3. **`index.md` 里的 `(中文)` 标注已移除**。它原本用来标记「这页是例外」,而中文现在是常规。 ## 验证 - 仓库自带的 `validate_wikilinks.py`:**145 条链接 / 25 篇,与翻译前完全一致** —— 零链接漂移。翻译前先建立了基线数字再对比,而不是只看翻译后能不能通过。 - 完整 `npm run wiki:build` 成功,`validate_site.py` 校验 **4159 条内部链接 / 146 个 HTML 页**。 - 逐页机械比对 HEAD:wikilink 目标、frontmatter、H1、**代码块字节相同**、**数字集合相同**、callout/标题/表格行/列表项计数相同。 - 构建产物里实际抓 `<html lang>` 确认:**25 篇笔记全部是 `zh-CN`**,而不是只读源码推断。 ## 已知限制 51 个 alias 重定向存根仍然是 `lang="en-us"`。这个值硬编码在 `alias-redirects` 插件受 lockfile 校验保护的构建产物里,改它会破坏 `quartz_plugins.py verify`,且下次 `plugin install` 就被覆盖。这些页面带 `robots: noindex`、`refresh` 立即跳转、不渲染任何内容,不值得为此破坏插件锁定校验。如实记录,而不是宣称覆盖率 100%。 Co-authored-by: Copilot App <223556219+Copilot@users.noreply.github.com> Copilot-Session: c80f8522-983c-47f7-8241-2155a823aabe

#1487) ## What Adds `docs/benchmarks/windows-cuda-runbook.md` — the Windows counterpart to the existing Linux/H200 `docs/benchmarks/H200-CUDA-runbook.md`. It documents how to bring up the **native CUDA execution provider** on a Windows dev box that has **no CUDA toolkit installed**, where every CUDA library comes from pip wheels under `site-packages\nvidia`. Cross-linked from: - `docs/benchmarks/README.md` (next to the H200 runbook link) - `.agents/skills/profiling/SKILL.md` (whose "source the CUDA env script" step is Linux-oriented) ## Contents - A copy-pasteable PowerShell env block (cuBLASLt + cuDNN on `PATH`, `CUDA_PATH` at the NVRTC headers) plus a general wheel-discovery snippet for other machines. - A known-good, coherent-output smoke command, validated against `models\qwen2.5-0.5b`. - The three "missing DLL/header" failures, in the order they appear, each with its verbatim native-EP error and one-line fix — with the contrast that the native EP fails **loudly** while ORT silently falls back to CPU. - The `profile_native` `--tokens` > `--decode-skip` guard. - The ORT-CUDA path (`ONNX_GENAI_ORT_LIB_DIR` + provider DLLs + `ONNX_GENAI_EP_FALLBACK=1`; `ort-sys` loads `onnxruntime.dll` by absolute path; the ort-prebuilt ORT is CPU-only). - Shared-box process hygiene and a measured-vs-inferred table. ## Verification **Every command was executed on the box**; results are recorded as passed/failed/ignored. Two of the author's field notes were corrected against what actually happened: - **Failure 2** surfaced as a raw NVRTC error (`could not open source file "cuda_fp16.h"` at `rmsnorm_bf16_v4`), not the friendly `cuda_ep <Op>:` pre-check — the Qwen model's first half kernel is an RMSNorm that skips that pre-check. Same root cause and fix. - **Failure 3 (cuDNN) did not fire** for the `qwen2.5-0.5b` native GQA path even with cuDNN off `PATH`; the run still produced coherent output. The cuDNN error is real code that fires only when a model reaches a cuDNN-backed op. **Not verified:** an end-to-end ORT-CUDA run (no CUDA-enabled ORT is installed on this box); those claims are code-/artifact-verified only. No performance numbers are included (cold NVRTC cache makes first-run timings meaningless). No CI changes. ⚠️ This is a docs-only PR; validated locally, not via CI. --------- Co-authored-by: justinchuby <223556219+Copilot@users.noreply.github.com>

@pris

> **Process note, stated up front.** This defect entered `main` via #1363, which I merged with an admin bypass while every required check was still `queued`. That was wrong, I am not repeating it, and this PR goes through the normal gates. Full disclosure of what I bypassed is in the comment below. `parallel_output_rows_dispatches_to_the_task_runtime` **fails on every stock CI runner** and is live on `main` today. ## The defect The test asserts the flat fan-out reaches the task runtime. But routing reads `rayon::current_num_threads()`, and `flat_fan_out`'s *first* gate is deliberately "stay on Rayon below `MIN_ROUTED_FAN_OUT_WIDTH` (16)". Below that width the test asserts something policy never promised. Measured on unrepaired `main`: | `RAYON_NUM_THREADS` | 4 | 8 | 15 | 16 | 32 | |---|---|---|---|---|---| | result | **FAILED** | **FAILED** | **FAILED** | ok | ok | `ubuntu-latest` is 4 vCPU. The whole `onnx-runtime-ep-cpu` lib suite on unrepaired main at that width: ``` test result: FAILED. 1447 passed; 1 failed; 17 ignored kernels::matmul_nbits::tests::parallel_output_rows_dispatches_to_the_task_runtime ``` It passed for me only because this development host is 16C/32T — the defect needs a *narrower* machine to appear, which is exactly the kind of thing the CI I bypassed exists to find. The existing `task_runtime::width() <= 1` guard does not cover it: task-runtime width and Rayon width are different numbers, and on a 4-vCPU box the first is `> 1` while the second is `< 16`. ## The fix, and the trap in it Install a Rayon pool of exactly the routing width so the decision under test is host-independent. My first attempt only wrapped the fan-out — and **still failed at `rayon=1`**, because `output_chunk_len` reads the same Rayon width and the test's *precondition* carried the identical defect. Moving the precondition inside the pool too is what actually removes the host dependency rather than relocating it. Skipping below the threshold would have been the weaker fix: the test would silently no-op on every real runner and guard nothing. - passes at rayon = **1, 2, 4, 8, 15, 16, 32** - **still falsifies** — forcing `PrefillFanOut::Wide` makes it fail, so it is not vacuous - adds the coverage assertion it should always have had (every output row written exactly once) ## Scope Test-only. No production behaviour changes. Deliberately **not** included: - the **aarch64 dead-code break** #1363 also shipped (`WIDE_PREFILL_MACS`, `prefill_fan_out`, `prefill_column_grain` are dead on a non-mlas ARM64 build, failing `-D warnings`) → **#1382** by @pris was open first and is already armed. I had written the same three `allow(dead_code)` restorations, verified they clear `cargo clippy --target aarch64-unknown-linux-gnu -- -D warnings`, then dropped them from this branch rather than ship a conflicting duplicate. - the **Stacked-Borrows UB** in #1377's test → **#1385** by @pris, and **#1407** (mine, which additionally closes the Miri lane gap that let it through: `miri.yml` runs `--lib task_runtime::` only, so integration tests under `tests/` are never Miri-checked). ## Validation | check | result | |---|---| | `cargo fmt --all -- --check` | clean | | `cargo test -p onnx-runtime-ep-cpu --lib` (host width) | 1447 passed, 0 failed | | `cargo test -p onnx-runtime-ep-cpu --lib` at `RAYON_NUM_THREADS=4` | **1447 passed, 0 failed** (main: 1 failed) | | target test at rayon 1/2/4/8/15/16/32 | all pass | | falsification probe (force `Wide`) | fails as required | Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

## Why CI's Rust quality lane (`cargo fmt --all -- --check`) runs on Linux and is fine. The **local** gate that is supposed to catch rustfmt drift *before* it lands is non-functional on Windows — which is where this repo's agent workflow runs, from git worktrees. The quality lane on `main` has been repaired for rustfmt drift four times (#1260, #1320, #1393, #1400). *That the broken local gate is the cause of those four repairs is a plausible inference, not something I measured — I only verified that the local gate does not work.* This PR makes the local gate runnable on Windows. It does **not** claim to prevent future drift. ## Defects (each verified on this Windows box) Environment: `rustfmt 1.9.0-stable`, `cargo 1.97.1`, Git-for-Windows bash 5.3, `git config core.autocrlf = true`, workspace = **54 members / 972 tracked `.rs` files**, mixed-edition (**52 on edition 2024, 2 on 2021**). 1. **Shell scripts check out as CRLF.** `.gitattributes` only pinned `schema/inference_metadata.schema.json`; `*.sh` and the extension-less `scripts/hooks/*` were unprotected, so all 14 tracked `.sh` files + the hook showed `w/crlf`. Running one under bash printed `scripts/install-hooks.sh: line 10: $'\r': command not found` and `set: pipefail: invalid option name` — unrunnable. 2. **`install-hooks.sh` cannot work in a worktree.** It used `HOOKS_DST="$REPO_ROOT/.git/hooks"` and bailed if that dir was missing. In a worktree `.git` is a **file**, so it always errored "are you in a git repo?". 3. **`cargo fmt --all` fails on Windows regardless.** `cargo fmt --all -- --check` exits **1** with `The filename or extension is too long. (os error 206)`: cargo-fmt passes every path to one `rustfmt`, overflowing the Windows ~32 KB command-line limit. Linux CI is unaffected (`ARG_MAX` ~2 MB). The old `pre-commit` ran this with **`2>/dev/null`** and then told the user `Fix: cargo fmt --all` — a command that also fails with os error 206. *(This fails loudly with exit 1 — there is no false-green here.)* 4. **No hook was installed** in this checkout — a consequence of 1+2. **Mixed-edition trap (why the fix uses `cargo fmt -p`, not raw `rustfmt`):** with the wrong edition, `rustfmt` mis-parses 2024-only syntax (e.g. `let` chains: `error: let chains are only allowed in Rust 2024 or later`) **and fails**. Only cargo knows each package's declared edition, so driving the check per package is the only correct approach. *(An earlier claim of a silent exit-0 false green was traced to a measurement artifact — `rustfmt … | Select-Object -First N` truncates the pipeline and drops the native exit code — and has been withdrawn; the failure is exit 1.)* ## Changes - **`.gitattributes`** — pin `*.sh` and `scripts/hooks/*` to `eol=lf` (comment explains a CRLF bash script is unexecutable) and renormalize. All 15 files now report `i/lf w/lf attr/text eol=lf`. - **`scripts/install-hooks.sh`** — resolve the hooks dir via `git rev-parse --git-common-dir` (the shared gitdir used by the main checkout and every linked worktree), resolving a relative result to absolute. Keeps `--dry` and the "does not clobber foreign hooks" property. - **`scripts/hooks/pre-commit`** — map the staged `.rs` files to their owning workspace packages and run `cargo fmt -p <pkg> -- --check` only for those. Now **mirrors CI's scope exactly**: - Files whose crate is **not a workspace member** (e.g. the root-level `bench-*` crates) are **skipped with a warning**, because `cargo fmt --all` does not cover them either. Blocking on a non-member would recreate the os-error-206 failure shape (`cargo fmt -p <non-member>` → "not a member of the workspace") and wall people off behind drift they never introduced. Membership is taken from `cargo metadata --no-deps` (matched on `manifest_path`, which is unambiguous — bare `"name"` keys also appear on every dependency). - If `cargo metadata` itself fails, the hook **fails open** (warns, lets the commit through) — a format gate must not lock you out of the repo. - Stops suppressing stderr; the printed fix is `cargo fmt -p <pkg>` (works on Windows). - **`wiki/development/Testing and Verification.md`** — state plainly that `cargo fmt --all` does not work on Windows here; give the per-package alternative and `bash scripts/install-hooks.sh`. ## Verification (measured on this box) - **`install-hooks.sh --dry`** succeeds from the **worktree** and (relative-`.git` branch) from a **normal checkout**, both resolving to the same shared `…/onnx-genai/.git/hooks`. Run under **Git-for-Windows bash**, which actually executes hooks. *Note:* WSL bash cannot run git in a Windows-created worktree at all — the `.git` pointer holds a `C:/…` path WSL's git can't resolve; that is a WSL/Windows limitation affecting every git command there, not this script. - **End-to-end, against the committed hook:** - staged a mis-formatted **member** `.rs` → commit **blocked** (exit 1), diff shown, fix `cargo fmt -p onnx-runtime-cpuinfo` printed; ran it → commit **passed**. - staged only a **non-member** (`bench-seqmajor`) `.rs` → commit **passed** with the "not a workspace member … CI's cargo fmt --all does not cover them either" skip warning. - staged a mis-formatted **member** *and* a **non-member** together → commit **blocked**, and the block came **only** from the member; the non-member was skipped and the printed fix command works. - simulated `cargo metadata` failure (stub returning 101) → hook **exited 0** with the fail-open warning. - All test artifacts discarded; nothing committed. - **`main` is clean** by the new check: looping `cargo fmt -p <name> -- --check` over all 54 members → **checked 54, failed 0, ignored 0**. - **Hook wall time** on a realistic single-package staged change: **~1–2 s** (the hook only checks the staged packages, not all 54). - **Full-member confirmation timing** (this is the `main`-clean sweep, not the per-commit hook cost): two consecutive runs **26.1 s** then **25.2 s**, consistent with an independent 23.1 s measurement. An earlier one-off 87.5 s reading was a non-reproducible first-run outlier and is not representative. ## Not touched - `.github/workflows/ci.yml` — CI is not broken; this is a local-gate fix. - Anything under `.squad/`. ## Rebase (onto latest main) Rebased from base `4b1cabb8` onto `origin/main` at `1557a355` (which had advanced through #1482, #1173, #1420, #1487). The **only** conflict was in `wiki/development/Testing and Verification.md`: #1482 translated the whole wiki to Chinese (`lang: zh-CN`), so my originally-English Windows-formatting section collided with the now-Chinese baseline. Resolved by **following the new Chinese baseline** — the added formatting/pre-commit documentation is written in Chinese to match the surrounding prose, and none of #1482's translation was reverted. Per project rules, code, commit messages and this PR title/body stay in English; only that wiki body follows its file's language. Checked that #1487's `docs/benchmarks/windows-cuda-runbook.md` neither overlaps nor conflicts with the wiki formatting note (the runbook covers CUDA benchmarking and contains no formatting/hook content), so no cross-link was needed. After the rebase, re-ran the three end-to-end scenarios (member-block→fix→pass, non-member-only→pass+skip-warning, mixed→blocked-only-by-member) and the `main`-clean sweep (**checked 54, failed 0**) — all still correct. Test artifacts cleaned; working tree clean. Co-authored-by: justinchuby <223556219+Copilot@users.noreply.github.com>

…nt4 (#1434) Building the decode pool is unconditional today: every borrowed-int4 call installs a 16-worker Rayon pool before dispatching, including the decode-shaped calls that then route all of their work to the task runtime and never touch it. Those workers are constructed, parked, and torn down once per call for nothing. This defers the construction for exactly one path — decode-shaped (`m == 1`) borrowed int4 — and hands the routing logic the width it *would* have installed, so the executor choice and the partition grain are both unchanged. ## What is actually covered The scope is deliberately narrow, and the narrowness is what makes it provable rather than measured: - Three call sites are converted: `packed_nbits_gemv`, `gemv_nk`, and the borrowed-int4 site gated on `m == 1`. - Those paths reach Rayon only through `parallel_output_rows_repeated` (`parallel_output_rows` delegates to it), which is the hookable dispatcher. - The `m == 1` gate is what makes this safe by construction, not by inspection: `borrowed_affine_int4_matmul_prefill` drives Rayon directly, and it is unreachable at `m == 1`. The other seven `with_decode_pool` sites are untouched, so the kernels that use Rayon directly — `packed_nbits_gemm`, `int8_matmul`, `int8_row`, `parallel_n16_output_rows`, `parallel_kai_output_rows` — keep eager installation. Routing is preserved rather than assumed. `flat_fan_out` branches on `rayon::current_num_threads()`, which equalled the decode width only because we were installed; `effective_fan_out_width()` reproduces exactly that width while deferred. The second commit extends the same helper to `output_chunk_len`, so the *grain* cannot drift from the installed grain either — without it a deferred call partitioned into 4096-row chunks where the installed one used 64. Five tests, two of them verified as falsifiers by deliberately breaking the thing they check: - relaxing the gate to `m >= 1` makes `only_the_decode_shaped_borrowed_int4_path_defers_the_pool` fail; - dropping the grain fix makes `a_deferred_fan_out_partitions_for_the_pool_it_would_have_installed` fail (chunk 64 vs 4096). ## Measured 16-core budget, interleaved with alternating arm order in a single session. | | before | after | |---|---|---| | process threads | 48 | **32** | | `onnx-genai-decode-*` workers | 16, 340-490 ms CPU | **none** | | voluntary ctxsw / iter | 24.61 | **16.12** | | total CPU | 3.660 cpu-s | 3.580 cpu-s | | dispatches / iter | 1.30 | 1.08 | **Latency is neutral, and that is the claim — not an improvement.** Six A/B reps gave 0.938, 1.180, 1.130, 1.115, 1.031, 1.026 (median 1.073), and an A/A null control measured in the same window gave a 0.83-1.21 band. Every ratio is inside the band, so this PR does not demonstrate a latency change in either direction. The win is structural: 16 fewer threads and a third fewer voluntary context switches. One number deserves an explicit caveat: total CPU is flat, not lower. The decode pool's CPU does not disappear, it reappears on the caller thread. That is attribution changing, not work being removed — the work was always the caller's, it was just being done by borrowed workers. `onnx-runtime-ep-cpu` lib suite: 1453 passed, 0 failed on merged latest main. fmt and clippy clean for this crate. (`cargo fmt --check` currently reports five sites repo-wide, all inherited from main and none in the one file this PR touches; #1393 repairs them.) --------- Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

…verdict (#1489) Coordinator campaign-brain update (doc-only, .squad/identity/now.md). - Bank the honest fp16 fair-test reversal: qwen3.5 fp32 'moats' were export artifacts; fp16io same-graph is ~2.1x ORT-ahead (magnitude flagged unverified pending a real ORT captured harness). - Record PR #1486 native decode fusions (+8.9% qwen3.5-0.8b-hybrid, byte-identical, capture preserved). - Record #1474 memory-stack verification verdict (6 deterministic async-deferred-release GPU test failures; no decode regression). Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

…trip on native decode (#1486) ## Summary Route **f16 `ReduceSum`/`ReduceMean`** to the existing single-kernel NVRTC block reduction (fp16 IO, f32 register accumulation) instead of cuDNN. This eliminates the self-inflicted fp32 round-trip on the hybrid decode hot path. ### The problem (profiling finding) A faithful captured-graph node trace of qwen3.5-0.8b-hybrid (fp16io, keep_io_types=False) decode — agent Deckard, `.squad/decisions/inbox/deckard-fair-hybrid-gap.md` — localized native's #1 fixable inefficiency: the SSM / linear-attention path does its reduction as a **cuDNN fp32 `ReduceSum` bracketed by fp16↔fp32 `op_tensor` cast wrappers**, ~**20.7%** of decode GPU-kernel time (LinearAttention 9.9% + cuDNN ReduceSum 8.1% + op_tensor casts 2.7%), firing 18×/step. `cudnnReduceTensor` rejects a half `reduceTensorCompType` and forces `CUDNN_DATA_FLOAT`, so an f16 reduce becomes a **three-kernel** round-trip (fp16→fp32 cast, fp32 ReduceSum, fp32→fp16 cast) through a full-size fp32 temporary. ### The fix The typed NVRTC block reduction already implements the exact ONNX *accumulate-in-f32-then-cast* rule in a **single capture-safe kernel** — no fp32 temporary, half the memory traffic. It already serves bf16 and every extended f16 reduction, so this is a **general** routing change (all f16 sum/mean, any shape/axis/dim), not a qwen-specific special case. f32 stays on cuDNN (native f32 comp type, byte-identical). One-line behavioral change: the cuDNN reduce gate goes from `Float32 | Float16` → `Float32`. ## Before / after — native decode (H200, `--test-threads=1`, matched A/B on the same build) | | ms/tok | tok/s | captures | fallbacks | |---|---|---|---|---| | before (cuDNN f16 reduce) | 4.664 | 214.41 | 4 | 0 | | **after (fused NVRTC f16 reduce)** | **4.318** | **231.61** | 4 | 0 | **Reclaim: 0.346 ms/tok (+17.2 tok/s, ~7.4%)** — within Deckard's predicted 0.3–0.4 ms. ## Parity - `generated_token_ids` **byte-identical** before vs after (greedy decode). - GPU tests pass, tolerances unchanged: `reduce_comptype_fp16_gpu`, `reduce_capture_gpu`, `nvrtc_reduce_capture_gpu`, `linear_attention_gpu`. ## Capture preserved `cuda_graph: enabled=true captures=4 replays=504 fallbacks=0` — no eager fallback; the fused f16 reduce folds into a single captured segment (verified by `reduce_capture_gpu`'s single-segment assertion). ## Generality The NVRTC block reduction is shape/axis/dim/head agnostic (base/delta offset walk, one block per output element). No dims hardcoded; no model-shape assumptions. bf16 already used this path; f32 remains on cuDNN. Refs profiling finding: `.squad/decisions/inbox/deckard-fair-hybrid-gap.md`. --------- Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

The `Benchmarks` workflow posted a **`🔴 Benchmark Regression Detected`** comment on PRs that had changed nothing relevant. The cause is not the comparison logic — it is the machine. **A GitHub-hosted runner is not a quiet machine.** It is shared, its neighbours are invisible to us, and we cannot verify it was idle for the window we timed. **Why this is worse than publishing nothing.** A red flag that fires on noise trains readers to ignore it. Once the header stops carrying information, a *real* regression posted the same way gets dismissed too — so the failure mode is not "a useless comment", it is "a gate nobody believes". **Why the base-vs-PR design did not rescue it.** Timing base first and PR second in the same job cancels some runner-to-runner variation, but it cannot cancel *within-job* contention: the two halves still run at different moments against different neighbours. The thresholds were calibrated against ~27% worst-case measured noise, and regressions worth catching are often smaller than that — so the signal sat **below** the noise floor, not above it. **Nothing is lost.** This workflow never gated anything; its own header says `Does NOT block CI`. The real regression gates are the throughput floors and dispatch-reachability tests in `crates/onnx-genai-bench/tests/profile_native.rs`, which run on real hardware and are untouched here. ## What changed - Removed the `pull_request` trigger; the job is now `workflow_dispatch`-only, taking a PR number as input. - A manual run checks out the default branch rather than the PR, so it now resolves the head sha and base ref via `gh pr view` and fetches the head explicitly. - Updated the remaining `github.event.pull_request.*` references (concurrency group, comment lookup, comment post) to use the input. The workflow was **also disabled in the Actions UI** (`gh workflow disable`) so the noise stopped immediately rather than waiting on this merge. ## Verified - `python -c "import yaml; yaml.safe_load(...)"` → parses; triggers are exactly `['workflow_dispatch']`. - No remaining `github.event.pull_request` references in the file. - `benchmark.yml` is the only workflow that runs `cargo bench` / `criterion` — grep across `.github/workflows/*.yml` confirms no other PR-triggered benchmark job exists. **Not verified:** the manual `workflow_dispatch` path has not been executed end-to-end (it needs a real dispatch on a macOS runner). The PR-metadata resolution is a straightforward substitution, but it is untested — worth a trial dispatch before anyone relies on it. Co-authored-by: justinchuby <223556219+Copilot@users.noreply.github.com> Copilot-Session: d60eb808-7cc6-4abc-b48d-2a6dd3841624

`cargo fmt --all -- --check` is a hard gate in both `Fast (Linux x86_64)` and `Rust quality`, and it fails on `main` @ `ba52d1702` with four diffs, all in `crates/onnx-runtime-ep-cuda/src/optimizer.rs` — one method chain at line 308, three in the `CudaRsqrtFusion` tests. Every PR that merges main inherits it, so nothing in the queue can go green until it is repaired. Same lane #1393 unbroke four days of breakages ago. Pure `cargo fmt --all` output. No semantic change. ## Local validation | gate | result | |------|--------| | `cargo fmt --all -- --check` | **0 diffs** | | `cargo test --locked $(… offline-linux)` | **3989 passed, 0 failed, exit 0** | | `cargo clippy --locked --all-targets $(… offline-linux) -- -D warnings` | exit 0 | | `cargo check -p onnx-runtime-ep-cuda` | clean | | `scripts/check_cross_compile.sh` | PASS (x86_64 + aarch64) | | 9 guard scripts | PASS | Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

… retired) (#1493) Deckard's trustworthy fair fp16io baseline: ORT eager ~290 tok/s flat (496 not reproducible, retired); ORT runs the hybrid on CUDA (1019/50, no structural moat); native captured 233.6 ÷ ORT 290 ≈ 0.80× → ORT ~1.24× ahead. Real lever is native decode efficiency, not ORT deficiency. Doc-only. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

…1382) ## What this fixes Validating the merged scheduler/fmt wave, I found `#1363` had dropped the `allow(dead_code)` guards on three prefill fan-out symbols, breaking the aarch64 cross-compile gate. **#1443 has since fixed that** — but by `#[cfg]`-ing the three items to `target_arch = "x86_64"`, which removes them outright. That trades one break for two others. ### 1. 🔴 `aarch64 + feature = "mlas"` no longer compiles The three symbols have two callers, gated on **different** things: | call site | enclosing fn | its gate | | --- | --- | --- | | `matmul_nbits.rs:2457` | `run_mlas_shards` | `#[cfg(feature = "mlas")]` | | `matmul_nbits.rs:6914-6915` | `borrowed_affine_int4_matmul_prefill` | `#[cfg(target_arch = "x86_64")]` | Off x86 the second caller disappears, but the **first does not** — it is arch-independent. So on `aarch64 + mlas`, which is Apple Silicon, the caller is compiled and its callee is not: ``` error[E0425]: cannot find function `prefill_fan_out` in this scope --> crates/onnx-runtime-ep-cpu/src/kernels/matmul_nbits.rs:2457:27 error[E0425]: cannot find function `prefill_fan_out` in this scope --> crates/onnx-runtime-ep-cpu/src/kernels/matmul_nbits.rs:6914:8 error[E0425]: cannot find function `prefill_column_grain` in this scope --> crates/onnx-runtime-ep-cpu/src/kernels/matmul_nbits.rs:6915:48 ``` **How that was produced.** MLAS's vendored sources do not cross-build to aarch64 in this container (`arm_neon.h: inlining failed in call to always_inline vaddq_f16 — target specific option mismatch`, an mlas-sys/toolchain issue unrelated to this PR), so instead I reproduced the *exact cfg resolution* on the host: on `main`, retarget the three item gates from `x86_64` to a third arch so they are absent, leave every caller alone, and build the lib with MLAS on — ``` sed -i '395s/x86_64/s390x/; 403s/x86_64/s390x/; 543s/x86_64/s390x/' \ crates/onnx-runtime-ep-cpu/src/kernels/matmul_nbits.rs cargo check --locked -p onnx-runtime-ep-cpu --features mlas --lib # exit 101 ``` That is precisely the configuration `aarch64 + mlas` produces. The `:2457` error is the load-bearing one — that call site is `feature`-gated only, so it is present on aarch64 for real. **Why no gate caught it.** `check_cross_compile.sh`'s aarch64 pass builds **default features**, and `mlas` is off by default — the same blind spot #1443 was written under. The configuration *is* built elsewhere: the `rust-coverage` job's macOS-arm64 leg runs with `RUSTFLAGS: -D warnings` (`ci.yml` L434, L439) and builds `cargo build -p onnx-runtime-ep-cpu-plugin --features mlas` (L501-503), whose `mlas` feature forwards to `onnx-runtime-ep-cpu/mlas`. That is the shipped-wheel build path, so `main` as it stands also breaks the macOS-arm64 MLAS wheel at release time. ### 2. 🔴 The #1363 fan-out policy stopped being tested on aarch64 Gating the items forced gating their tests, so #1443 also put `#[cfg(target_arch = "x86_64")]` on six unit tests. All six are pure functions of explicit literal arguments — `prefill_fan_out(WIDE_PREFILL_MACS - 1, 16, 32)`, `prefill_column_grain(8, 1024, 3072)` — with nothing architecture-specific in them. They are now simply not compiled off x86, so the policy that #1363 rewrote has no aarch64 coverage. ## The fix `cfg_attr(.., allow(dead_code))` instead of `cfg`. The item always exists, so whichever caller survives can reach it; the lint is silenced only where **no** caller exists. The six tests are ungated and run everywhere again. `prefill_tile_grain` in this same file already uses this idiom (`not(feature = "mlas")`) — that is the shape #1363 deleted. The predicate is **per-symbol**, because the caller sets differ: | symbol | callers | predicate | | --- | --- | --- | | `WIDE_PREFILL_MACS`, `prefill_fan_out` | `run_mlas_shards` **and** `borrowed_affine_int4_matmul_prefill` | `not(any(feature = "mlas", target_arch = "x86_64"))` | | `prefill_column_grain` | `borrowed_affine_int4_matmul_prefill` **only** — `run_mlas_shards` takes `prefill_tile_grain` instead | `not(target_arch = "x86_64")` | Giving `prefill_column_grain` the union predicate would leave the lint live on `aarch64 + mlas`, where it has no caller — converting #1443's `E0425` into a `never used` error in the same configuration. Review caught exactly that in the first draft of this branch; the four-way probe below is the regression check for it. ### Four-config probe of the predicates `.validation-worktrees/cfgprobe/probe.rs` reproduces the two items, the two callers and their gates with `mlas`/`x86` standing in for the real cfgs, compiled under `-D warnings`: ``` === union predicate on prefill_column_grain (wrong) === PASS [aarch64 default] FAIL [--cfg mlas] <- error: function `prefill_column_grain` is never used PASS [--cfg x86] PASS [--cfg mlas --cfg x86] === per-symbol predicates (this PR) === PASS [aarch64 default] PASS [--cfg mlas] PASS [--cfg x86] PASS [--cfg mlas --cfg x86] ``` The probe is sharp, not vacuous: it fails on exactly the configuration that is wrong, and only that one. ## Second commit: unbreaking `main`'s required lane `main` currently fails **both** required checks, from merges landed past queued checks: | defect | source | breaks | | --- | --- | --- | | `map_or(true, ..)` in `executor/dispatch.rs` — clippy `this map_or can be simplified` under `-D warnings` | #1427 | `Rust quality` — and it aborts the cross-compile gate *before* its aarch64 pass, which is why the gate never reported defect 1 | | `dispatch.rs`, `gather_block_quantized.rs`, `gpt_oss_20b_decode_lock.rs` unformatted | #1427, #1418 | `Fast (Linux x86_64)` **and** `Rust quality` | `cargo fmt --all -- --check` runs in *both* required jobs (`ci.yml` L162, L276) while `check_cross_compile.sh` runs only in `Rust quality` (L401), and PR checks run against `merge(base, head)`. So while `main` is broken this way a fmt-only PR still fails the cross-compile step and this PR alone still fails fmt — only a branch carrying both can go green. It is mechanical (`cargo fmt --all`, plus `map_or(true, f)` → `is_none_or(f)`, identical on `Option`) and `git rebase` drops it once fixed upstream. ## Verification at `9fb04f5b5` (base `main` `81f99ff42`) | check | step | `main` | this branch | | --- | --- | --- | --- | | `Fast` + `Rust quality` | `cargo fmt --all -- --check` | **FAIL** (4 diffs / 3 files) | **pass** | | `Rust quality` | `bash scripts/check_cross_compile.sh` | **FAIL** exit 1 | **pass** exit 0, `scope: full offline set (aarch64 cross toolchain present)` | | `Rust quality` | 30-crate `cargo clippy --locked --all-targets … -- -D warnings` | **FAIL** exit 1 | **pass** exit 0 | | `Rust quality` | 9 guard scripts | pass | **9/9 pass** | | aarch64 | `cargo clippy --target aarch64-unknown-linux-gnu --all-targets -p onnx-runtime-ep-cpu -- -D warnings` | pass | **pass** (now *with* the 6 tests compiled) | | `aarch64 + mlas` cfg resolution | `cargo check -p onnx-runtime-ep-cpu --features mlas --lib`, items absent | **FAIL** exit 101, 3 × E0425 | **pass** — `cfg_attr` never removes the item, so E0425 cannot occur | | all 4 `(mlas on/off) x (x86 / non-x86)` | `rustc -D warnings` cfg probe | — | **4/4 pass** | | tests | `cargo test -p onnx-runtime-ep-cpu --lib` | — | **1447 passed / 0 failed**; the 8 prefill policy tests pass | ## A note on the gate that found this `scripts/check_cross_compile.sh` **false-passes locally** without an aarch64 cross toolchain: at L191-194 it silently swaps `CRATES_FULL` → `CRATES_NO_FFI`, dropping `onnx-runtime-ep-cpu` — the crate the gate exists for — and still exits 0 with a ✓. The "REDUCED SCOPE" note prints *below* the checkmark. **Read the scope note, not the exit code**; only `scope: full offline set (aarch64 cross toolchain present)` means anything. On Actions it `exit 2`s instead (L178-190), and `ci.yml` L396-399 installs `gcc-aarch64-linux-gnu` + `libc6-dev-arm64-cross` before invoking it, so the fail-loud coverage is intact — this is a local-only trap. All results above were produced with the toolchain installed, at full scope. Two of the three defects in this PR would have been caught by the required checks had they been allowed to run. ## Process No admin bypass, no ruleset bypass, no merge with checks queued or failing. Auto-merge has been armed since 2026-08-19T04:55:37Z and merges only once `Fast (Linux x86_64)` and `Rust quality` are green. Every CI run in this repo is currently `queued` with zero in progress, so the required contexts have not been created yet. Waiting. --------- Co-authored-by: Pris <pris@squad.local> Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

…d (2.31 -> 1.71 at depth 100) (#1387) > **Scope note.** This PR now carries the whole #1077 dispatch-overhead stack. Enabling auto-merge across the stacked PRs caused each to merge into its *parent feature branch* — those branches are unprotected; only `main` requires checks — and GitHub retargeted the children, collapsing the stack into this one. **No `main` protection was bypassed and no work was lost.** This PR still waits on required CI before it can reach `main`. **What is in here, in dependency order — each was independently reviewed as its own PR before it merged:** | was | what | evidence | |---|---|---| | #1387 | deterministic phase/counter instrumentation for the dispatch path | Opus 4.8 REQUEST CHANGES → all findings fixed | | #1392 | stop building debug strings production throws away | 5 → 4 allocations per input per `Run` | | #1394 | the #1077 benchmark grid (depth 1/10/100, static + dynamic) | produced the first real decomposition | | #1397 | stop asking ORT three times where the inputs live | **7 → 4** FFI calls per `Run`, Opus 4.8 APPROVE WITH NITS → fixed | | #1401 | routed-loop instrumentation; O(depth²) retirement scan; lazy absent strides; 8-element dispatch-isolating cases | **2.310 → 1.713** at depth 100 (−24%) | **Headline result** — 8-wide Relu chain, depth 100, 6000 interleaved iterations, pinned cores, ORT arm as drift control (moved 1.7%): | | baseline | after | |---|---|---| | ours p50 | 54.2 µs | **41.0 µs** | | ratio vs ORT | 2.310 | **1.713** | | per-node | 0.52 µs | **0.39 µs** | Two of my own earlier conclusions were falsified along the way and are documented in #1077: a 10-node chain is **one** fused dispatch (so the per-node gap contains no FFI at all), and the original grid was measuring **memory bandwidth**, not dispatch. --- ## Why Issue #1077 says our per-`Run` dispatch overhead is higher than ORT's. Every time I have had to answer *where* that overhead is, the method has been the same: hand-edit `Instant` probes into the hot path, run the A/B, read the numbers, tear the probes out. That is slow, it is unreviewable, and it leaves nothing behind — the next person starts from zero, and nothing catches a regression that quietly adds a round trip back. This makes the measurement a permanent, reviewable part of the crate. ## What `dispatch_probe`, gated behind a **non-default** Cargo feature, wired through the CPU dispatch path: | Phase | Where | |---|---| | `CallbackEntry` | `compute_execute` entry, up to the first real work | | `MetadataQuery` | `read_inputs` | | `TensorBind` | single-node input view construction | | `Allocate` | `allocate_output` | | `DispatchLookup` | `infer_shapes`, `prepare_workspace` | | `KernelInvoke` | `execute_with_workspace` | | `StatusCrossing` | `status_with_code` | Plus event tallies for the three things actually worth removing: `OrtFfiCall`, `DispatchAlloc`, `StatusCreated` (and `ComputeExecute` / `NodeExecuted` as denominators). ## Design notes **The counters are kept twice, deliberately.** The thread-local copy is the precise one — a dispatch is a single-threaded story, and only per-thread accumulation lets one measurement be read without a concurrently-running sibling test bleeding into it. That isolation is what makes the exact-count assertions below possible at all; I tried global-only first and the counts were nondeterministic under the default parallel test runner. The global mirror exists because ORT chooses the thread that runs `Compute`, and it need not be the thread that called `Run`. The e2e harness loads this EP as a cdylib through `RegisterExecutionProviderLibrary`, so it reads the counters back through an exported `nxrt_dispatch_probe_snapshot` symbol — and it has no way to be on ORT's worker thread. A thread-local-only probe would report zero there, which reads as "we made no FFI calls" rather than "you are looking at the wrong thread". Writing both costs one `Cell` store and one relaxed `fetch_add`. **Production pays nothing.** Without the feature, `PhaseGuard` is a zero-sized type with no `Drop` impl and every entry point is an `#[inline(always)]` empty function. Call sites use `guard.end()` rather than `drop(guard)` so the production build does not read as dropping a `Drop`-less ZST. **Timing is gated a second time**, on `ONNX_GENAI_PROFILE_DISPATCH=1`. A test asserting FFI call counts wants the counters but emphatically does not want two `Instant::now()` calls added to every phase it is measuring. ## The counts are pinned, not documented The point of the module is the regression guard, so the numbers are asserted rather than written in a comment. `kernel_ctx` gains a hand-built `OrtApi` — zeroed, then filled in with only the entry points `read_inputs` actually uses, so any call it makes that the test did not anticipate faults loudly instead of silently succeeding — which lets `read_inputs` run with no live ONNX Runtime at all. Today, per `Run`: - **8 FFI round trips** for one input (within `read_inputs`) — one shared `GetInputCount`, then 7 per input - **5 heap allocations** — the `Vec<OwnedInput>` once, then `dims`, the `format!("input {i}")` label, `shape` and `strides` per input - three inputs cost `1 + 3×7` and `1 + 3×4`, so this is a linear model rather than a single data point - an absent optional input short-circuits at 2 calls, which pins the optional-slot path Any change that adds a round trip or an allocation to the per-`Run` path now has to come here and edit a number, in a diff a reviewer can see. ## Verification Both feature configurations: `cargo fmt`, `cargo clippy --all-targets` (zero warnings), and the full `--lib` suite — 251 tests without the feature, 257 with it. The real ORT-driven suite passes unchanged: **55/55 `plugin_ort_e2e`** with `NXRT_REQUIRE_ORT_TESTS=1`, including `every_assigned_node_is_also_executed_by_this_ep`. **Every test here was verified by mutation**, because a test that passes against the known-broken version is not a test: - Injecting one extra `get_type_shape` call and one extra `Vec` allocation into `read_inputs` failed exactly the three count assertions and correctly left the other two green. - The zero-cost assertion **was vacuous when I first wrote it** — it compared `snapshot()` after counting against `snapshot()` before, which are two calls to the same constant function in a disabled build. It passed unchanged against a mutant whose "disabled" `count_n` recorded into a static and whose `snapshot` returned non-zero. It now compares against `Counters::default()` — an absolute claim rather than a self-consistency one — and fails against that mutant. It also asserts `!needs_drop::<PhaseGuard>()`, closing the one way the production guard could regain a cost while staying zero-sized. ## Review round 1 — Opus 4.8, `REQUEST CHANGES`, all addressed The review found a real defect, and it was the module's central claim. **Blocker: `OrtFfiCall` was documented as a per-`Run` total but only `read_inputs` and `allocate_output` were instrumented.** On the happy path of a single-node `Run`, `device_mem_info`, `ort_input_mem_info`, `mem_info_is_device`, `alloc_scratch`, `host_pool::install` and `CreateStatus` all reach into ORT and none were counted. A probe that under-counts while calling itself exact is worse than no probe — it invites exactly the wrong conclusion about where the overhead is. Fixed by instrumenting every leaf `OrtApi` call in the crate, and then by making that property **self-enforcing** rather than a promise. The new `ffi_coverage` tests scan the source of each file that touches `OrtApi` and extract the API members it names — ORT's C API is `CamelCase` and our own fields are `snake_case`, so `.CamelCase` isolates them cleanly — then fail if that set or the instrumentation count moves. Verified by mutation: adding a stray `GetTensorShapeElementCount` reference fails with > `compute.rs now names 10 ORT API members ([… "GetTensorShapeElementCount" …]), not 9. If you added an FFI call, add an ort_call() beside it and update this table.` A second test pins that the scan finds real members, so it cannot rot into a heuristic that matches nothing and passes forever. **`DispatchAlloc` got the opposite treatment: it is now documented as a lower bound.** The reviewer was right that hand-placed allocation counts cannot be exhaustive — `Vec::new()` does not allocate, `Vec::with_capacity(0)` does not allocate, and whether a `collect` allocates once or twice depends on `size_hint`, not on anything visible at the call site. Claiming exactness there would repeat the same mistake in a place where it cannot be fixed. It is exact for `read_inputs`, where the sites are few and each is pinned; callers needing the true whole-`Run` figure get `CountingAllocator`, a `GlobalAlloc` wrapper to install in a test or bench. This crate deliberately does not install it, since a library defining a `#[global_allocator]` takes that choice away from every dependent. Also fixed: - A node with **no** inputs was charged one allocation for a `Vec::with_capacity(0)` that never happens. Now conditional, pinned by `a_node_with_no_inputs_allocates_nothing`. - `CreateStatus` was counted as a `StatusCreated` but not as an FFI call, undercounting error paths. - **Phases are documented as not being a partition.** On success they are non-exhaustive; on error they *nest*, because a guard closes at scope exit rather than at the early `return`, so `StatusCrossing` opens inside whichever phase was live. `DispatchLookup` is entered twice per node and reports the total. Summing `phase_ns` gives neither wall time nor a pie chart, and a reader who assumed otherwise would be misled. - `probe_is_compiled_out_in_production` no longer implies it has proven `count` is side-effect-free — a disabled build has no storage to observe one. It proves what it can: the guard is a ZST that does not need dropping. Confirmed by the reviewer independently: `mem::zeroed::<OrtApi>()` is sound (all 422 fields are `Option<unsafe extern "C" fn>`, so NPO makes all-zero a valid `None` everywhere); the `without_provenance` fake pointers are never dereferenced; and all four `read_inputs` counts are correct, checked against a counting global allocator. Re-validated after the fixes: clippy clean and `--lib` green in both configurations (253 / 260), and **55/55 `plugin_ort_e2e`** against real ORT. ## What this sets up Two allocations are already visibly waste and are removed in a follow-up that this PR exists to measure: 1. `compute_execute` calls `staging_log(&format!(…))` **unconditionally** — `staging_log` early-returns when disabled, but the `String` is built and three values formatted on every dispatch regardless. 2. `read_inputs` builds `format!("input {i}")` eagerly per input, on the success path, purely to label an error that is usually not raised. Relates to #1077. Independent of #1244/#1246 — additive and feature-gated, so it can land in any order. --------- Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

…d MatMulNBits prefill (#959, #1091) (#1176) ## Summary On the **default (`mlas` OFF) build** the `dequant-kn` prefill phase now goes to **zero** for every `MatMulNBits` path that reaches the dense fallback: the native CPU EP gets a transposed-B ("NT") SGEMM and non-MLAS prefill routes through it, reusing the already-cached contiguous `Nk` weight instead of materializing a second, transposed `Kn` copy. Refs #959, #1091. **Does not close #959** — the superlinear per-token *decode* cost (~102 s/token at 14B) is a separate open question this does not touch. This closes the *prefill* `Kn`-materialization term. ## The problem (#959) Int4/int8 `MatMulNBits` prefill (`m > 1`) on the default build dequantized the weight to f32 a **second** time in the transposed `Kn` (`[k, n]`) layout — a strided-scatter transpose (each K step written at stride N), **uncached**, then a dense NN GEMM. #959 measured it at **~2.9x** the contiguous `Nk` pass and *degrading with N* (22 s/GB at 0.5B → 38 s/GB at 14B; **266 s of ~357 s** time-to-first-token on qwen2.5-14b). MLAS hosts already avoided this (`try_prefill_mlas_nt`) by feeding the cached `Nk` weight to MLAS's cache-tiled `sgemm` with `trans_b`. The default build had **no** transposed-B GEMM, so its `#[cfg(not(feature = "mlas"))]` sibling returned `false` and fell back to the slow `Kn` path. ## What changed MLAS's advantage here is not magic — it is a packer that reads the B operand row-wise from `[n, k]`. Ported natively: - **`x86_sgemm.rs`** — `sgemm_simd_nt(a, b_nk, c, m, k, n)` computing `C[m,n] = A[m,k] · B_nk[n,k]^T`. It reuses the entire packed GEBP path (`pack_a`, KC/strip blocking, `micro_6x16`) and adds **`pack_b_nt`**, which gathers each output column from its *contiguous* `b_nk` row (unit stride over `k`) into the L1-resident pack tile at stride `NR`. The transpose becomes a **pack-time reshape of an in-cache tile**, not a full-array stride-`N` scatter — exactly why MLAS wins. - **`matmul.rs`** — `nt_gemm_supported(backend)` is the **single** legality predicate (no two `cfg` arms to drift); `gemm_nt_with_backend` dispatches `Mlas → trans_b`, `SimdX86 → sgemm_simd_nt`. - **`matmul_nbits.rs`** — the two cfg-split `try_prefill_mlas_nt` arms collapse into one **`try_prefill_nk_nt`** gated on `nt_gemm_supported`. It dequantizes once into the pre-existing `weight_nk` `OnceLock` (the same slot decode caches into, so a constant weight pays **one** dequant, not two) and calls `gemm_nt_with_backend`. When the NT route runs, the `dequant-kn` profile phase is skipped by construction. - **Host pool (#1143)** — column strips dispatch onto the installed ORT intra-op pool when present, rather than forking a rayon pool beside it; rayon otherwise. Strip decomposition is numerically transparent. ## Correctness — bit-identical **Byte-for-byte identical** to the existing `Kn` dense route, not "within 1e-5". `pack_b_nt` produces the *same* packed panels `pack_b` would from the `[k, n]` transpose (`b_kn[p·n+j] == b_nk[j·k+p]`); with identical A-pack, K-panel order, and microkernel, every output element's f32 accumulation sequence is unchanged. Strip count and pool choice do **not** affect per-element reduction order, so bit-identity holds regardless. Verified (`assert_eq!` on `f32::to_bits`): - Odd/tail shapes: `m ∈ {1, 2, 7, 33}`, `n ∈ {1, 3, 63, 64, 65}`, plus tile-exact and multi-KC. - 64 randomized shapes (differential NN-vs-NT). - End-to-end by the existing 8-bit prefill oracle test (`matmulnbits_8bit_prefill_batched_matches_dequant_f32_oracle`), which reaches the dense fallback and thus the NT route. This matches the standard the MLAS NT route already held itself to (bit-identity to the no-transpose dense GEMM, recorded at `matmul_nbits.rs` ~2048). ## Measurement **`dequant-kn` → 0 is structural.** The `mm_profile::time_prepack("dequant-kn", …)` call lives in the `if !used_fast_nt {…}` branch; whenever the NT route returns `true` that branch is skipped, so the phase is eliminated by construction. `dequant-nk` is unchanged (still one pass, now the only one). **Finding vs the #959 premise:** on current local builds, 4-bit `acc0` prefill takes the borrowed-int4 in-place path (#979/#1117) and 4-bit `acc4` uses the SDOT prepack — *neither reaches the dense fallback*, so no local q4 model emits a `[mm_prepack] phase=dequant-kn` line to drive to zero end-to-end (confirmed empirically with `ONNX_GENAI_PROFILE_MM=1` on qwen05b q4 / q4-acc4 / symzp). The beneficiaries of this change are therefore: **8-bit** weights (`m>1`), **grouped** quantization, **weight_prepacked**, and **4-bit with `accuracy_level != 0`** that falls to the dense fallback. No local 8-bit/grouped model was available for an end-to-end token-to-first-token arm; per the profiling skill I do not report a contended wall-clock figure I cannot defend. **Per-phase microbench** (`nt_prefill_bench`, `#[ignore]`; best-of-7, `--test-threads=1`, release; **contended box — other agents building concurrently**). Bit-identity asserted in the same harness. The `dequant-kn*` arm is a *plain f32 transpose* standing in for the strided `Kn` materialization the NT route removes (the real int4 dequant is ~2.9x heavier per #959): | shape (m=16) | dequant-kn* (transpose) | NN gemm | NT gemm | old (kn*+NN) → new (NT) | |---|---|---|---|---| | k=5120, n=5120 (100 MiB) | 337 ms | 5.7 ms | 5.0 ms | 343 ms → **5.0 ms** | | k=5120, n=13824 (270 MiB) | 641 ms | 15.0 ms | 11.8 ms | 656 ms → **11.8 ms** | | k=13824, n=5120 (270 MiB) | 1307 ms | 15.3 ms | 11.7 ms | 1323 ms → **11.7 ms** | The eliminated transpose term dominates and grows with size (as #959 predicted); the NT GEMM itself is even **slightly faster than NN** here (contiguous per-column B reads pack better). Per #1132, a native-faster-than-MLAS result on some shape is a graduation event for `benches/native_vs_mlas.rs` — noting it, but **not** changing default routing without that gate's measurement. **RSS.** No new long-lived allocation (reuses `weight_nk`; `apack`/`bpack` are per-call local scratch, freed at return). Removing the second full f32 weight materialization cuts the transient f32 footprint of a prefill that hits the dense fallback by one full `[k, n]` copy (e.g. **270 MiB** at k=13824/n=5120). Not measured end-to-end (no local model reaches that path); the reduction is structural. ## Memory rules No new field that outlives a call or scales with weight size — the NT route reuses the pre-existing `weight_nk` `OnceLock`. `apack`/`bpack` are per-call `vec![]` scratch. The added lines are not matched by `weight-cache-guard.yml` (no `OnceLock<…Vec>` / `(RefCell|Cell)<Vec|Box|Arc>` introduced); its regex and path filter are untouched. ## Gates (exact counts) - `cargo test -p onnx-runtime-ep-cpu --lib` → **1324 passed, 0 failed, 17 ignored** - `cargo test -p onnx-runtime-ep-cpu --lib --features mlas` → **1354 passed, 0 failed, 28 ignored** (MLAS route still works and still wins where enabled) - `cargo clippy -p onnx-runtime-ep-cpu --all-targets -- -D warnings` → **clean** - `cargo fmt --check` → the three changed files are clean (verified with `rustfmt --edition 2024 --check`). Pre-existing diffs remain in three *unrelated* files (`governed_accumulator_budget.rs`, `qlinear_matmul.rs`, `simd_activations.rs`) from a local rustfmt version skew vs CI — left untouched to keep this PR surgical. --- 🤖 Generated with Squad. Flagged **needs review** — please have a squad member review the kernel packing/tail handling before merge. --- ## Update (2026-08-18, Roy) — merged current `main`, plus a production-path A/B ### Merge Merged `main` (`c55a3fab3`), which had since made the #1091 M=1 GEMV the unconditional `SimdX86` route and dropped the `ONNX_GENAI_CPU_MM_SIMD_M1_GEMV` toggle this branch still carried. Conflict resolved by keeping **both**: main's default-route test (`the_default_entry_point_routes_m1_to_the_gemv`) and this branch's NT kernel + bit-identity tests. Two follow-on fixes: - **aarch64 cross-arch lane.** `gemm_nt_with_backend` compiled with neither the `mlas` nor the x86 arm has no reader for any parameter, so the `-D warnings` cross-arch pass rejected all six. Bound them in the unsupported arm. (This is what the old `Rust quality` red was: `Cross-target compile check`, nothing else.) - **`mm_profile` gemv phase.** The MLAS route used to time this GEMM on the `gemv` phase; after the two call sites collapsed into one, that timer was lost. Restored — the default build gets a pass-through (`tick()`, the reporter, is MLAS-only), so the shared call site stays `cfg`-free and MLAS profiling is unchanged. ### Production-path A/B (new harness, `benches/matmul_nbits_prefill_ab.rs`) The original body was right that no local model reaches this route, and honest about not reporting a number it could not defend. That gap is now closed the way #1013 closed its own: drive the **real kernel through the EP's own `get_kernel`/`execute`**, at the shapes and inputs that *do* reach the dense fallback — 8-bit prefill, and 4-bit with `g_idx`. The harness uses no symbol this branch introduces, so the identical file runs on `main` and here; both arms were built and run **interleaved**, 3 repetitions each, on the same box. Host: 32-core x86_64, AVX2, default build (**`mlas` off**), release. **Contended** (other agents building; load ~12), so medians of per-arm medians are reported and the ratios — not the absolute ms — are the claim. **Steady state** (weight already resident, per-call prefill cost): | case | k | n | m | main (ms) | this PR (ms) | speedup | |---|---:|---:|---:|---:|---:|---:| | int8 dense fallback | 2048 | 2048 | 8 | 12.114 | **0.548** | **22.1x** | | int8 dense fallback | 2048 | 2048 | 64 | 12.276 | **1.375** | **8.9x** | | int8 dense fallback | 4096 | 11008 | 8 | 53.858 | **4.170** | **12.9x** | | int8 dense fallback | 4096 | 11008 | 64 | 56.560 | **7.595** | **7.4x** | | int4 + `g_idx` fallback | 2048 | 2048 | 8 | 31.583 | **0.563** | **56.1x** | | int4 + `g_idx` fallback | 2048 | 2048 | 64 | 32.405 | **1.419** | **22.8x** | **Cold** (fresh kernel per repetition, so the one-time weight dequant is inside the measurement — the TTFT term #959 attacked): | case | k | n | m | main (ms) | this PR (ms) | speedup | |---|---:|---:|---:|---:|---:|---:| | int8 dense fallback | 2048 | 2048 | 8 | 6.470 | 6.008 | 1.08x | | int8 dense fallback | 2048 | 2048 | 64 | 12.519 | 7.974 | 1.57x | | int8 dense fallback | 4096 | 11008 | 8 | 53.761 | 37.260 | 1.44x | | int8 dense fallback | 4096 | 11008 | 64 | 57.725 | 38.503 | 1.50x | | int4 + `g_idx` fallback | 2048 | 2048 | 8 | 31.636 | 23.200 | 1.36x | | int4 + `g_idx` fallback | 2048 | 2048 | 64 | 32.083 | 25.883 | 1.24x | **Why steady moves 7–56x and cold only ~1.1–1.6x — and why that is the real finding.** The `Kn` route has **no cache**: `dequantize_weight(WeightLayout::Kn)` is called *inside* `execute`, so every prefill call re-materializes the whole transposed f32 weight. The NT route dequantizes into the pre-existing `weight_nk` `OnceLock`, which a constant weight fills **once**. So this change removes not one transpose but *every repeat of it*. Cold (first call) improves by the layout alone — a contiguous `Nk` write instead of the stride-`N` scatter, 1.2–1.6x at these sizes; steady improves by the caching the `Nk` layout makes possible. Confirmed structurally with `ONNX_GENAI_PROFILE_MM=1` over the same harness run: **main emits 72 `phase=dequant-kn` lines, this PR emits 0 — and 24 `phase=dequant-nk`** (one per kernel instance, i.e. the cold arms only; every steady call pays none). **Bit-identity, across builds.** The harness prints an FNV-style digest of the raw output bits. All six rows have the **identical digest on both arms** (`db6ff07f991d431`, `b4c1df9fd8883789`, `7bd7418eacdcb870`, `2eed012609649617`, `d5834afbecf42494`, `cd8fa77262302969`) — bit-identity of the production `execute` result, not just of the kernel driver, verified across two separately compiled builds. ### Scope, restated honestly On the default build, 4-bit `accuracy_level=0` with contiguous (borrowable) inputs takes the zero-copy borrowed int4 path (#979/#1117/#1126) for both decode *and* prefill and returns before the dense fallback. This PR therefore changes: **8-bit prefill**, **4-bit with `g_idx`**, **`weight_prepacked`/non-borrowable** inputs, and **4-bit `accuracy_level != 0`** that falls through. Those are exactly the cases measured above. MLAS builds already had the NT route; their behaviour is unchanged. ### Gates (re-run after the merge) - `cargo test -p onnx-runtime-ep-cpu --lib` → **1420 passed, 0 failed, 18 ignored** - `cargo clippy -p onnx-runtime-ep-cpu --all-targets -- -D warnings` → clean - `cargo clippy --locked --target aarch64-unknown-linux-gnu --all-targets -p onnx-runtime-ep-cpu -- -D warnings` → clean (the lane that was red) - `cargo fmt --all -- --check` → the files this PR touches are clean; one *inherited* diff remains in `onnx-runtime-ep-cuda/standard_attention.rs` from `main`, fixed separately in #1347. --------- Co-authored-by: justinchuby <223556219+Copilot@users.noreply.github.com> Co-authored-by: Roy <roy@squad.local>

# Conflicts: # crates/onnx-runtime-ep-cpu/src/kernels/matmul_nbits.rs

#1176 landed the native transposed-B SGEMM, so `try_prefill_nk_nt` now succeeds on the default x86-64 build for 8-bit prefill. That made this PR's fused GEBP unreachable: it was consulted only after the NT route declined, and the NT route no longer declines. The two are not redundant, though. NT wins whenever it keeps its dequantized `Nk` weight resident -- one dequant per session, then a cache-tiled NT GEMM. When the #971 memory governor declines that cache, NT has nothing to amortize into and rebuilds the whole `k * n` f32 panel on every call (measured: `dequant-nk calls=5` for five prefills), which is exactly the cost this fused pack removes, in exactly the configuration where memory was already too tight to hold it. So order on `nt_keeps_weight_resident` rather than on NT declining. Measured, 8-bit cells, native-alone, min of 5: governor declined: 2.97x - 21.53x faster than main cache admitted: 0.97x - 1.05x (neutral; NT keeps priority) Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

`nt_keeps_weight_resident` is consumed only by the x86-64 fused-GEBP arm, so aarch64 saw it as an unused variable and `cargo clippy --target aarch64-unknown-linux-gnu -- -D warnings` failed. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

justinchuby · 2026-08-19T19:58:49Z

Reconciled with #1176 (now merged) — and it changed what this PR is for

#1176 merged as 9a452d97d, which landed the native transposed-B SGEMM. That has a direct consequence for this PR that I had to resolve before merging.

#1176 made this PR's kernel unreachable

This PR's fused GEBP was consulted only if !used_fast_nt — i.e. only after the NT route declined. Before #1176, try_prefill_nk_nt returned Ok(false) immediately on any non-MLAS build, so the fused GEBP was the only route for bits == 8, m > 1. After #1176 the NT route succeeds on the default x86-64 build, so used_fast_nt is always true and this PR's kernel never ran. I confirmed that directly with ONNX_GENAI_PROFILE_MM=1: on the merged tree an 8-bit prefill printed [mm_prepack] phase=dequant-nk calls=1, the #1176 route, and never entered the GEBP.

Merging it in that state would have been dead code.

But the two are not redundant — they split on weight residency

The premise in this PR's own doc comment ("the alternative is materializing the whole k * n weight as f32 on every call") is still true, just no longer unconditional. It depends on whether the NT route keeps its dequantized Nk weight resident:

Resident cache admitted (constant weight + int4 weights are materialised to f32: real footprint is ~8x the file, and the memory plan accounts for the file size #971 governor admits): NT pays one dequant for the whole session and then streams each weight row once. NT is the right route and this PR should stay out of the way.
Governor declines the cache (set_resident_dequant_f32_cache_enabled(false), called from engine/load.rs and pipeline/mod.rs when the expanded f32 footprint does not fit the budget): NT has nothing to amortize into and rebuilds the entire k * n f32 panel on every call. Measured on the merged tree with the cache declined: [mm_prepack] phase=dequant-nk calls=5 for five prefills — five full panel materializations. This PR's fused pack allocates only per-strip scratch and materializes nothing.

That second case is exactly where materializing ~51 MB per prefill hurts most: the governor declined the cache because memory was already tight.

So I reordered on nt_keeps_weight_resident = can_prepack && resident_dequant_f32_cache_enabled() instead of on "NT declined", and updated the stale doc comment. This is reconciliation #2 from my cross-PR conflict analysis, now applied as code rather than left as a note.

Measured, 8-bit cells, native-alone, min of 5 (`--native-only`, ORT pool not running)

cell (8-bit)	governor declined main	#1403	speedup	cache admitted main	#1403	ratio
qwen3_0p6b_qkv_t8	4.479	0.208	21.53x	0.203	0.194	1.05x
qwen3_0p6b_qkv_t32	4.533	0.530	8.55x	0.491	0.505	0.97x
qwen3_0p6b_qkv_t128	5.110	1.151	4.44x	1.118	1.123	1.00x
qwen3_0p6b_mlp_t8	5.829	0.387	15.06x	0.249	0.248	1.00x
qwen3_0p6b_mlp_t32	5.936	0.631	9.41x	0.589	0.569	1.04x
qwen3_0p6b_mlp_t128	5.514	1.402	3.93x	1.329	1.355	0.98x
llama3_8b_qkv_t8	17.269	1.747	9.88x	1.927	1.900	1.01x
llama3_8b_qkv_t32	18.529	2.840	6.52x	2.984	3.060	0.98x
llama3_8b_qkv_t128	22.758	7.662	2.97x	7.821	7.655	1.02x
llama3_8b_mlp_t8	39.194	3.664	10.70x	4.003	3.917	1.02x
llama3_8b_mlp_t32	40.526	6.076	6.67x	5.891	5.841	1.01x
llama3_8b_mlp_t128	50.635	15.813	3.20x	15.473	15.588	0.99x

2.97x–21.53x where it now claims ownership, and 0.97x–1.05x (neutral) in the default configuration where #1176 keeps priority. Numerics unchanged: parity=PASS on 8/8 8-bit cells and on 4-bit spot checks.

A cross-compile defect the local gate caught

The first version of the reconciliation bound nt_keeps_weight_resident only in the x86-64 arm, so cargo clippy --target aarch64-unknown-linux-gnu -- -D warnings failed with unused variable. Fixed in f2a982d62. Worth noting because scripts/check_cross_compile.sh would not have caught it — on Linux its TARGET is x86_64-unknown-linux-gnu, the native target, so it exercises neither aarch64 nor Windows.

Local gate matrix on `f2a982d62` (this branch with `origin/main` merged)

fmt --all · offline-crate build · -p onnx-runtime-ep-cpu 1524 passed / 0 failed · clippy offline -D warnings · --features mlas tests · native-backend clippy · aarch64-linux clippy -D warnings · --no-default-features · --all-features · default_artifacts_are_mlas_free · check_cross_compile.sh · 7 guard scripts · workspace_test_packages verify. 19/20 green.

Miri: strided 8, provider 16, dtype 12, task_runtime 29 = 65 tests, 0 failures.

The one non-green step is cargo check --target aarch64-pc-windows-msvc, which cannot run on this Linux host — onnx-genai-ort-sys bindgen needs the Windows SDK (fatal error: 'stdlib.h' file not found). I reproduced that failure on plain origin/main in a clean worktree, so it is a host limitation rather than anything from this PR. Recording the exact scope instead of claiming coverage I do not have.

justinchuby and others added 21 commits August 18, 2026 19:25

justinchuby and others added 8 commits August 18, 2026 22:34

justinchuby and others added 4 commits August 19, 2026 08:19

justinchuby and others added 21 commits August 19, 2026 09:03

Merge remote-tracking branch 'origin/main' into HEAD

1bc95e4

# Conflicts: # crates/onnx-runtime-ep-cpu/src/kernels/matmul_nbits.rs

justinchuby merged commit 9d8a9c4 into squad/roy-int4-prefill-gebp Aug 19, 2026
5 of 6 checks passed

justinchuby deleted the squad/roy-int8-prefill-gebp branch August 19, 2026 19:59

justinchuby mentioned this pull request Aug 19, 2026

perf(cpu-ep): fuse the int4 dequant into the prefill GEMM pack step (23.7x) #1356

Merged

Navigation Menu

Repository navigation

perf(cpu-ep): fuse the 8-bit dequant into the prefill GEMM pack step (17x kernel; every ORT prefill cell loss -> win) - #1403

perf(cpu-ep): fuse the 8-bit dequant into the prefill GEMM pack step (17x kernel; every ORT prefill cell loss -> win)#1403
justinchuby merged 90 commits into
squad/roy-int4-prefill-gebpfrom
squad/roy-int8-prefill-gebp

justinchuby commented Aug 19, 2026

Uh oh!

justinchuby commented Aug 19, 2026

Uh oh!

codecov Bot commented Aug 19, 2026 •

edited

Loading

Uh oh!

justinchuby commented Aug 19, 2026

Uh oh!

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

1 participant

Conversation

justinchuby commented Aug 19, 2026

The loss

The fix

The numbers are bit-identical, not merely close

Result — kernel level

Result — against ONNX Runtime

Controls

Row gate is derived, not inherited

Scope

Tests and gates

Also

One honest note on the matrix

Uh oh!

justinchuby commented Aug 19, 2026

Uh oh!

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

Uh oh!

Codecov Report

Uh oh!

justinchuby commented Aug 19, 2026

Reconciled with #1176 (now merged) — and it changed what this PR is for

#1176 made this PR's kernel unreachable

But the two are not redundant — they split on weight residency

Measured, 8-bit cells, native-alone, min of 5 (--native-only, ORT pool not running)

A cross-compile defect the local gate caught

Local gate matrix on f2a982d62 (this branch with origin/main merged)

Uh oh!

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

1 participant

codecov Bot commented Aug 19, 2026 •

edited

Loading

Measured, 8-bit cells, native-alone, min of 5 (`--native-only`, ORT pool not running)

Local gate matrix on `f2a982d62` (this branch with `origin/main` merged)