Repository navigation
perf(cpu-ep): cheaper softmax exp8 via degree-6 pure-Horner minimax (1 ULP held) - #1245
Merged
Merged
Conversation
The softmax family is the most-lost CPU-EP kernel after #1234 vectorised the exp (1.29-1.63x ORT, thread-matched). A prior experiment falsified the bandwidth hypothesis (a 2-pass streaming softmax was 2.4x worse), leaving exp *throughput* as the lever. This reworks `exp8`'s polynomial while keeping the accuracy-critical Cody-Waite two-part `ln2` range reduction untouched. Change: replace the Cephes `1 + r + r^2·P5(r)` evaluation (5 poly FMAs + an `r^2` multiply + a trailing `+1` add) with a degree-6 Remez minimax on `r` in pure Horner form (6 FMAs, no `r^2` mul, no trailing add). The r^0 and r^1 minimax coefficients round to exactly 1.0f32, so the exact `1 + r` structure is preserved. Net: one fewer FMA-unit op and one fewer FADD per 8-lane exp, the softmax hot path's bottleneck on Zen 4. Rejected by measurement (kept only what wins): - magic-number rounding: ~5% slower (lengthens the t->k dependency chain with add+sub vs a single VROUNDPS). - ILP unrolling (2x/4x independent accumulators): monotonically slower; the loop is FMA-throughput-bound, not latency-bound, so the compiler already hides the FMA latency and manual unrolling only adds overhead. - exp2 reformulation: ~1.15x but 64 ULP (>> the 2 ULP bar). The speed came from computing `f = x*log2e - k` in single precision, which discards the low bits the Cody-Waite split exists to keep. An accuracy-preserving exp2 (two-part log2e) is back to 1 ULP but only ~1.04x, so not worth it. Accuracy: exhaustive f64-reference sweep over all 1.1e9 f32 in [-87.336, 0] gives max 1 ULP for both the old and new exp8 (unchanged). Adds a direct `exp8`-vs-f64 unit test asserting <= 2 ULP; the existing row-level tests cannot police this because softmax normalisation cancels a shared relative error (a 3-ULP degree-5 mutation fails the new test while all 12 row tests still pass). Measured end-to-end (bench_generic, native vs ORT, 1 thread each, interleaved A/B, 12-14 reps, 7 softmax models): native median time improves 1.03-1.13x (median ~1.06x); the native/ORT ratio drops in both min and median on every model (e.g. bert 1.49->1.36, kv8192 1.58->1.45). parity=PASS on all 7. Both binaries are MLAS-free (nm mlas count = 0). Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
## What `main` at 4d231ea fails two blocking CI steps on current stable (rustc/rustfmt 1.9.0, 1.97.1): ``` $ cargo fmt --all -- --check Diff in crates/onnx-runtime-ep-cuda/src/provider.rs Diff in crates/onnx-runtime-memory-api/src/allocator.rs Diff in crates/onnx-runtime-memory-api/src/capability.rs $ cargo clippy --locked --all-targets -p onnx-runtime-session -- -D warnings error: field `0` is never read --> crates/onnx-runtime-session/src/executor/mod.rs:175:45 ``` Both landed while the runner queue was saturated (71 queued / 1 in progress at the time of writing, nothing completed on `main` since 08:18Z), so no PR has seen a red check yet. Every open PR in the repo currently inherits both failures. ## Why these fixes **rustfmt** — mechanical normalisation, no semantic change. **`ActivationPlanForTest`** (from #1226) is a tuple struct whose single field is the `globals_lock()` `MutexGuard`. `dead_code` does not model "this field's value is its `Drop`", so it fires. The guard must stay: releasing it early is exactly the leaked-planner-gate race the struct was added to prevent. So the lint is silenced with a comment explaining the RAII intent, rather than the field removed. ## Verification | check | result | | --- | --- | | `cargo fmt --all -- --check` | clean | | `cargo clippy --locked --all-targets -p onnx-runtime-session -- -D warnings` | clean | | `cargo test --locked -p onnx-runtime-session --lib` | 181 passed, 0 failed | Found while trying to land the CPU task-runtime stack (#1201 → #1202 → #1207 → #1232 → #1238); this is unrelated to that work and is deliberately kept out of it so it can go in on its own. Working as sebastian (Performance Engineer) --------- Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
…ap_bias_ps (#1227) ## What this is `map_ps` is the loop driver behind every unary activation on AVX2. A merge (the `#1037` / `#1074` layering) inserted `map_bias_ps` *between* `map_ps`'s doc comment and `map_ps` itself, so the comment and **both of its attributes** ended up on the wrong function: ```rust /// Apply an 8-lane kernel across a slice. ... <- map_ps's doc #[inline] #[target_feature(enable = "avx2,fma")] <- map_ps's attributes /// Like [`map_ps`], but adds a bias row ... pub(super) unsafe fn map_bias_ps(...) <- ...on map_bias_ps pub(super) unsafe fn map_ps(...) <- bare ``` `map_bias_ps` needs those attributes too, so nothing was broken outright — but `map_ps` was left with no `#[inline]` and no `avx2,fma`. ## Why the missing `target_feature` is a hazard LLVM will not inline a callee that requires a feature its caller lacks (a *strict* superset; equal feature sets inline fine). Every kernel closure `map_ps` takes (`erf_ps`, `tanh_ps`, `sigmoid_ps`, ...) is `avx2,fma`. With `map_ps` compiled at baseline features, those closures are not inlinable into its loop. It is harmless **today** only by luck of ordering: every caller (`tanh_avx2`, `erf_avx2`, ...) is itself `avx2,fma`, so `map_ps` gets inlined *upwards* into the caller first, and once there the closure folds in. That is an inlining decision, not a guarantee. Grow `map_ps` past the inline threshold — which is exactly what an unroll experiment does — and the closure becomes a real call per 8 elements, and a kernel like `erf_ps` has to re-materialise all of its constants on every one of those calls. I hit this while unrolling `map_ps`, which is how it surfaced. ## Evidence that this changes nothing today Disassembled `tanh_avx2` out of the built rlib before and after: ``` diff <(objdump -d ... base) <(objdump -d ... fixed) -> identical ``` Byte-identical, and `erf_avx2` / `sigmoid_avx2` keep their instruction counts (216 / 117). No `call` to any kernel remains in the lib. So this is a latent-hazard and documentation fix with **zero** codegen delta — not a perf change, and it needs no A/B. ## Also recorded here: a negative result on the divide While looking at these loops I measured one algorithmic change and it lost, so it is written down rather than left for someone to retry. `tanh_ps` and `sigmoid_ps` each end in `_mm256_div_ps`. `vdivps ymm` is the only instruction in either loop that is not fully pipelined (Zen 4 retires one about every 4.5 cycles), so replacing it with `vrcpps` + one Newton-Raphson step — `y1 = y0 + y0*(1 - q*y0)`, four pipelined uops for one blocking one — looked like the obvious win. It is not: | case | ORT p50 (control) | divide | rcp + NR | |---|---|---|---| | `bench_tanh_f32_4k` | 0.0031 ms, all 6 runs | 1.518 / 1.525 / 1.526 | 1.550 / 1.570 / 1.570 | Interleaved rebuild-and-alternate, 3 rounds, 1 thread, `taskset -c 8-15`. The 4k case is the only trustworthy one here — it is L1-resident and ORT's own p50 was identical to four decimal places across all six runs, whereas on the 1M cases ORT's p50 moved 4.6x mid-session and neither arm is quotable. Consistent ~3% **regression**. The `rcp` + 2 FMA dependency chain is about as long as the divide it replaces, so a latency-bound kernel gains nothing and just pays two extra uops. It also breaks four exactness proofs that #1121 landed: `tanh(±Inf)` comes back as `0.99999994` instead of exactly `±1.0`, because the input clamp makes `p` and `q` bit-equal at `|v| = 9` and only an exact divide turns that into exactly `1.0`. Buying that back would mean restoring the saturation blend #1121 proved redundant over 1.05 billion inputs — to fund a change that is already slower. Not pursued. No fallback involved; the kernel keeps its own divide. ## Validation - `cargo test --release -p onnx-runtime-ep-cpu --lib` — 1337 passed, 0 failed - `cargo clippy --release -p onnx-runtime-ep-cpu --all-targets` — clean - `cargo fmt --all --check` — clean - `tanh_avx2` disassembly identical to `main` ## Review Opus review: no blockers, no should-fixes. It independently confirmed the three things the PR rests on: every one of the 16 `map_ps` and 2 `map_bias_ps` call sites is already inside a `#[target_feature(enable = "avx2,fma")]` function (so nothing becomes unsound or fails to compile), `map_bias_ps` keeps exactly the attributes it had, and there are **no bare/function-pointer references** to either. It also did the scan I most wanted: all ~130 `fn` definitions in the file were checked for the same orphaned-doc-comment-after-attribute signature that caused this bug. **Zero other anomalies** — every `*_ps` helper and every `*_avx2` dispatcher carries the attributes it should. So this is the only instance. - **NIT (applied)** — the doc comment said LLVM won't inline a callee whose features are "a superset" of the caller's. Read non-strictly that would forbid inlining at *equal* feature sets, which is the opposite of what the fix relies on. Reworded to "requires a feature its caller lacks — a *strict* superset; equal feature sets inline fine". --------- Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
🔴 Benchmark Regression DetectedComparison of criterion micro-benchmarks: PR head vs merge-base, measured on the same runner in the same job (base first → PR second).
Visual flags: Host infoWhat this cannot catch
|
## Summary - Reverts the exact merged squash commit `5c181b49ed9ee073639f1ee3df11ef7c597eec46` from #1192. - Withdraws the bundled #1186 Phase 1 + Phase 2 memory capability split rather than attempting another fix-forward patch. - Preserves unrelated mainline work merged after the squash. Part of #1186. ## Why revert first Subsequent high-effort review found the merged release/identity design structurally unsafe: - shared-mapping release can race free-list address reuse; - `0` conflates a legitimate no-op/refund result with failure; - partial CUDA unmap and handle-release failures can leave software and driver state divergent; - pointer-only identity is vulnerable to ABA/address reuse; - the contract makes mechanism-identity guarantees it cannot actually prove. These are lifecycle and ownership hazards, not isolated implementation bugs. Leaving the contract on `main` while iterating would expose callers to unsafe semantics, so this PR withdraws it completely. ## Related work - #1192 merged the Phase 1 branch into `main` as squash `5c181b49`. - #1206 merged Phase 2 into the Phase 1 branch and was therefore bundled into that squash. - Draft #1233 attempted to structurally pair capability allocation/release. - Draft #1237 attempted a narrower corrective rollback of the identity/atomic-release apparatus. Those drafts do not eliminate the need to withdraw the already-merged contract. A safe Phase 1 and a deliberately narrowed Phase 2 will be re-landed as a new reviewable stack after redesign. ## Validation - `cargo metadata --locked --no-deps --format-version 1` - `cargo check --locked --all-targets -p onnx-runtime-memory-governor -p onnx-runtime-cuda-memory -p onnx-runtime-ep-api -p onnx-runtime-ep-cuda -p onnx-genai-ort` - `cargo test --locked -p onnx-runtime-memory-governor --lib` — 64 passed - `cargo test --locked -p onnx-runtime-cuda-memory --lib` — 8 passed - `cargo test --locked -p onnx-runtime-ep-api --lib provider::tests::` — 6 passed - `cargo test --locked -p onnx-runtime-ep-cuda --lib provider::tests::` — 8 passed, 6 GPU-gated tests ignored - `cargo test --locked -p onnx-genai-ort --lib governed_allocator::tests::` — 20 passed - `git diff --check` Co-authored-by: Copilot App <223556219+Copilot@users.noreply.github.com>
Running [OpenCode](https://opencode.ai) against the OpenAI-compatible server, with Muse Glimmer 30B INT4 on the native CUDA pipeline, surfaced a set of gaps between what a real agent client sends and what the server accepted. Each was found by an actual failing agent turn, not by inspection. ## Fixes **Session header.** OpenCode attaches `X-Session-Id` to every request; the server answered `400 X-Session-Id is not supported by pipeline models`, so no agent turn ran at all. The header is a transport hint — the client resends the full history each turn, and a pipeline retains and rewinds its own context — so it is now ignored on pipeline models rather than rejected. Validation of the header itself is unchanged, and non-pipeline session semantics are untouched. **Tool call arguments.** OpenAI carries `function.arguments` as a JSON *string*, while chat templates index it as a mapping (`args.items()`), and the jinja sandbox cannot parse a string back into one. Sending a tool result back failed the whole turn: ``` chat template render failed: ... requires tool_call.function.arguments to be a dict (mapping) ``` Arguments are now decoded where the wire format is known. Anything that is not a JSON object is left as the original string rather than guessing at a shape the caller did not send. **Tool result identity.** A `tool` message's `name` and `tool_call_id` never reached the template, so a tool-calling template could not name the function a result answers. Both are now carried through, and `name` is accepted on the request. **Reasoning effort.** A reasoning model's template picks how much thinking to request from a caller-supplied keyword, defaulting to its own (often maximal) value. There was no way to set it, so an agent turn could spend its entire budget thinking and return nothing. OpenAI's `reasoning_effort` is now accepted and exposed to the template under both spellings in the wild, `reasoning_effort` and `reasoning_strength`. Measured against the real model, same prompt and same correct answer: | `reasoning_effort` | completion tokens | | --- | --- | | `low` | 64 | | `high` | 258 | **Sampling defaults.** Pipeline handles hard-coded `generation_defaults: None`, so every request decoded greedily at the OpenAI schema defaults and ignored the package's declared `generation` block. A model shipping `do_sample: true` degenerated into exact repetition — the agent's first turn looped one sentence until it hit the token budget. The pipeline already exposes `generation_defaults()`; the server now uses it. Relatedly, `top_k` is absent from the OpenAI schema and `0` is its "disabled" sentinel, so every silent client was overriding a model's declared `top_k` with "no truncation at all". An absent `top_k` is now treated as unspecified; an explicit one still wins. **ATEM protocol.** Muse emits tool calls as `<atem:invoke>` blocks and splits a turn into channels addressed to `self` (private reasoning) and `user` (the answer). Tool calls in that format are now parsed into standard OpenAI `tool_calls`, and only the user channel is returned as content. A turn truncated before reaching the user channel yields empty content rather than leaking the model's private reasoning as if it were an answer. ## Testing - `cargo test -p onnx-genai-server` — 191 pass; `cargo test -p onnx-genai-ort --lib chat_template` — 9 pass. - New regression tests for each fix: session header on a pipeline model, arguments decoded to a mapping, tool identity reaching the template, reasoning effort reaching the template, declared `top_k` preserved, and reasoning withheld without a user channel. - End to end with OpenCode 1.18.18 against Muse Glimmer 30B INT4 on native CUDA: the agent completes a full tool loop — model emits a `bash` tool call, OpenCode executes it, the result is sent back, and the next turn renders. Image attachments are accepted as `image_url` data URIs and reach the vision encoder. > [!NOTE] > Stacked on #1162, which exposes the pipeline `generation_defaults()` this consumes. > [!NOTE] > Agent turns against this model are still not reliably *useful* — under OpenCode's ~10k-token system prompt it often exhausts its budget mid-reasoning. That is model behavior, not a protocol gap, and is out of scope here. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> Copilot-Session: 0190e2eb-abe4-451f-b36d-44a035a99b7e
…6.85x -> 1.15x vs ORT) (#1183) ## What this changes #1091 added a native `M == 1` f32 SGEMV to the SIMD SGEMM but gated it behind `ONNX_GENAI_CPU_MM_SIMD_M1_GEMV`, defaulting **off**. The shipped build therefore never executed it: `MatMul` f32 at `M = 1` — every decode step — kept running the packed prefill kernel, which packs a one-row panel, walks it once and discards it. This makes the `m == 1` GEMV route unconditional in `sgemm_simd`, deletes the env probe and the `m1_gemv_enabled()` helper. `use_m1_gemv` survives only as a parameter of `sgemm_simd_variant`, which is how the `bench_f32_gemm_ab` A/B harness drives both arms in one process. Production no longer pays a locked `getenv` per GEMM call. ## Measurement (reviewer re-measurement) Measured by the squad reviewer on **Intel Core i7-13800H (14C/20T), Windows 11**, single-threaded (`RAYON_NUM_THREADS=1`), via the in-repo `bench_f32_gemm_ab` same-binary A/B harness (`simd_gemv` vs `simd_packed`, min-of-N timing). This harness drives exactly the two routes the default flip chooses between, so the `packed/gemv` ratio is the speedup this PR unlocks in production. It is **ORT-independent**; see the caveat below on the "vs ORT" figure. `packed/gemv` (higher = the newly-default GEMV is faster), M=1 decode shapes, across 4 runs: | shape (M=1) | packed | gemv | speedup | | --- | ---: | ---: | ---: | | 1×5120×5120 | 19.8–27.5 ms | 3.8–5.6 ms | **4.7–5.2×** | | 1×13824×5120 | 63.3–75.9 ms | 12.2–14.1 ms | **4.7–5.4×** | | 1×5120×13824 | 97.7–115 ms | 11.6–12.6 ms | **7.8–9.9×** | | 1×5120×7168 | 27.6 ms | 5.4 ms | **5.2×** | | 1×5120×152064 (lm_head) | 745 ms | 141 ms | **5.3×** | **Control (M=128 prefill, byte-identical code both arms):** `gemv/packed` ranged 0.97–1.05 across all runs — i.e. ≈1.0 as required, worst drift ~4.6% on this shared box. The unchanged-code rows not moving is the honesty check that the M=1 movement is real and not machine noise. The win is a robust **~5×** (shape-dependent, 4.7–9.9×) at M=1, reproduced over four interleaved runs. The route is reached only for `m == 1`; every `m >= 2` shape generates byte-identical code before and after (the control). ### On the "vs ORT" number The original body quoted "6.85x → 1.15x vs ORT". I did **not** reproduce that specific ratio: I have no pinned ORT f32 reference build on this Windows box, and a ratio against an unnamed ORT build is exactly the kind of figure this team has withdrawn before. What I can stand behind is the packed→gemv delta above (~5×), which is consistent with the packed path being several× slower than ORT and the GEMV landing near parity. The headline is stated as the measured internal speedup, not an unverified ORT multiplier. ## Numerics f32 addition is not associative, so the GEMV changes M=1 results by reassociation versus the packed path. This is bounded by the test tolerance (`1e-3 * (1 + max_ref)`) and is deterministic: the GEMV does no horizontal reduction (each output column is a lane reduced across `k` in a fixed unroll-by-4 order) and every strip start is a multiple of 8, so the same input yields the same bits at any thread count. `m >= 2` is unaffected (byte-identical). ## Tests - `the_default_entry_point_routes_m1_to_the_gemv` — asserts `sgemm_simd(m=1)` equals `sgemm_simd_m1` **bit-for-bit**, with `m=2` equal to the packed path as a control. Reverting the default to a probe fails it. - `m1_route_matches_packed_within_tolerance` — now compares against the genuine packed path (`sgemm_simd_variant(.., false)`) with an `assert_ne!` proving the two arms are different code, so the tolerance check is non-vacuous. - The `env::set_var` test from the original was correctly dropped (data race against parallel `env::var` readers; also couldn't fail for its own regression). ## Validation (reviewer) - `cargo test -p onnx-runtime-ep-cpu --lib` — **1346 passed, 0 failed, 16 ignored** (base 1345 + 1 new route test) on base `8a90878a`. - `cargo clippy -p onnx-runtime-ep-cpu --all-targets -- -D warnings` — clean. - No MLAS feature enabled for the shipped path; the `mlas` arm of the harness is a reference arm only and was not used for the numbers above. --------- Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
…e (replaces #1224) (#1249) Replacement for #1224, rebased onto current `main` after #1197 landed as a squash (the stacked base diverged, so the original went `CONFLICTING`). Same two commits, re-validated on top of the merged #1197. A real agent client driving this server through the OpenAI API exposed two defects that share a cause: the server answered questions the client never asked. ## A reasoning turn was cut off mid-thought OpenAI accepts only `max_completion_tokens` for a reasoning model and rejects `max_tokens`, and the server implemented neither — serde fell back to a fixed 256-token default, which a reasoning model can spend entirely on private thinking, ending the turn before it addresses the caller. `max_tokens` and `max_completion_tokens` are now both optional, the newer field wins, and a request naming neither falls back to the server cap rather than to 256 — what OpenAI means by an unspecified limit. The cap is a ceiling on generosity, not a reservation, so an unnamed budget yields to the prompt (`admit_output_budget`) instead of rejecting every budget-less request on a short-context model; a budget the caller *did* name still reports honestly that it does not fit. ## Private reasoning must not reach the client The model spells a turn as channels addressed to `self` (private), `user` (the answer), or a tool (a call). The buffered path already filtered via `atem_visible_content`; streaming sent tokens as they arrived, and a streamed token has its special tokens stripped, so the markers deciding visibility were invisible token-by-token and private reasoning went out live as content. `PrivateChannelGate` keeps the turn's special-token spelling and asks the same `atem_visible_content` what is visible so far, emitting only the growth — streaming agrees with the buffered path by construction. Two properties: - **Armed by declaration, not guesswork.** The gate is armed from the package's own `response_template` (`declares_private_channels`); a model that declares no private channel is never gated and streams exactly what it generated. - **Fails closed.** A token the gate cannot place is withheld. Also applied to per-token logprob chunks, constrained-JSON completions, and the stop-sequence drain, which previously dropped whatever the gate still held. ## Verification (mine, on this branch atop merged #1197) - **208 lib + 31 http + 11 bundle tests pass**; `cargo clippy -p onnx-genai-server --all-targets -- -D warnings` clean. - **Falsified both claims** (revert hunk → RED → restore → GREEN): - Privacy: forcing the streaming gate to pass tokens through made `reasoning_never_streams` leak `" to=selfsecret plan"` and `only_the_user_channel_streams` leak the `secret` channel; restored, both assert the reasoning is absent. - Budget: reverting the unspecified-budget fallback to 256 made `an_unspecified_budget_is_the_server_cap` return 256 instead of the 4096 cap; restored, GREEN. - **End to end** (`qwen05b-q4`, native): a budget-less request now completes with `finish_reason: stop` (no 256 truncation); streaming for a non-ATEM model streams unchanged (un-gated), confirming no regression for ordinary OpenAI clients; over-cap `max_completion_tokens` is rejected naming the exact field. The ATEM streaming-leak path itself is covered by the unit gate tests above rather than e2e, since that model declares no private channel. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> --------- Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> Copilot-Session: 0190e2eb-abe4-451f-b36d-44a035a99b7e
…is silent (replaces #1236) (#1250) Replacement for #1236, rebased onto current `main` after #1197/#1224 landed as squashes (the stacked base diverged). Same fix, plus one added guard test, re-validated on top of the merged stack. ## The defect `temperature` and `top_p` carried the OpenAI schema's documented `1.0` defaults via `#[serde(default = ...)]`, so the server could not tell a caller who *chose* 1.0 from a caller who said nothing. Both were then handed to `resolve_sampling_defaults` as explicit overrides, so a model that declares its own sampling regime lost it on every request. An absent field is "no opinion", not the documented default. ## The fix Both fields become `Option<f32>`; an absent one is passed as `None`, so the package's declared value wins — the same treatment `top_k` already had. An explicit value (including `temperature: 0`, greedy) still wins. Validation (`validate_sampling_range`) now judges only what the caller actually sent. The same fix applies to legacy `/v1/completions`. ## Verification (mine, on this branch atop merged #1197 + #1224) - **211 lib + 31 http + 11 bundle tests pass**; `cargo clippy -p onnx-genai-server --all-targets -- -D warnings` clean. - **Falsified the core claim** (revert hunk → RED → restore → GREEN): reverting `chat_sampling_overrides` to forward `Some(1.0)` for an absent field made `a_silent_request_keeps_the_declared_temperature_and_top_p` see `Some(1.0)` instead of `None`, and `silent_chat_request_honors_model_do_sample` override the declared `0.6` with `1.0`; restored, the declared value survives. - **Absent-vs-default trap, added test:** `an_explicit_default_valued_temperature_still_overrides_the_package` asserts a caller that explicitly sends `temperature: 1.0` (equal to the old schema default) still overrides a package declaring `0.6`. This documents that absence (`None` → package wins) and an explicitly-sent default (`Some(1.0)` → caller wins) are now distinct — the exact distinction this PR exists to make. - **End to end** (`qwen05b-q4`, native): a silent request answers normally; legacy `/v1/completions` at `temperature: 0` is bit-identical across two runs, so the greedy path is untouched; an explicit `temperature: 1.0` is accepted and a negative one rejected. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> --------- Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> Copilot-Session: 0190e2eb-abe4-451f-b36d-44a035a99b7e
…to repay it (3.15x on small shapes) (#1149) ## What this changes `half_gemm::gemm_impl` forked **every** half GEMM across the rayon pool with no work guard, so it paid fork/join overhead to multiply tiny operands a single core finishes far faster. Since the f16 prefill gate (#1140) this kernel only receives the *small* shapes that decline widening, so the unguarded fork was mis-sized for essentially every shape it still serves. Two guards, scheduling-only (arithmetic is untouched): 1. `PARALLEL_MIN_WORK = 1_048_576` on `m*k*n` — stay serial below it. 2. Never fork when the split yields a **single block** (`m.div_ceil(split_mc) == 1`); it cannot use more than one thread, so it is pure overhead however large the operand. `m == 1` always lands here. ## Measurement (reviewer re-measurement) Re-measured by the squad reviewer on **Intel Core i7-13800H (14C/20T), Windows 11**, via the in-repo `bench_half_gemm_parallel_threshold` A/B harness (`m = 8`, serial vs pool-split interleaved rep-by-rep, p50 of 9), at `RAYON_NUM_THREADS = 4/8/16`. The last column is `serial/par` — **>1 means splitting wins**. Smallest shapes (the ones this guard sends serial) — splitting is catastrophic, so declining it is the win: | `m*k*n` | T=4 | T=8 | T=16 | |---|---|---|---| | 32,768 (8×64×64) | 0.08 | 0.10 | 0.09 | | 65,536 | 0.13 | 0.16 | 0.15 | | 131,072 | 0.21 | 0.31 | 0.23 | So the worst shape (8×64×64) goes from a forked path to serial and is **~10–12× faster** on this box (par 0.39/0.25/0.42 ms → serial 0.033/0.025/0.037 ms at T=4/8/16). That is even larger than the 3.15× the original body reported on its Linux reference box, because this laptop's relative fork/join overhead is higher. Around the threshold and above: | `m*k*n` | T=4 | T=8 | T=16 | |---|---|---|---| | 1,048,576 (threshold) | 0.98 | 0.68 | 0.76 | | 1,572,864 | 0.95 | 0.79 | 0.96 | | 2,097,152 | 0.91 | 1.19 | 1.01 | | 3,145,728 | 1.29 | 1.49 | 1.17 | **Honest hardware caveat.** On *this* box the crossover where splitting starts to pay sits **higher than `1_048_576` — around 2–3M MACs**, not at ~1M. The original constant was derived on a 16-core Linux box. This does **not** make it unsafe here: the guard only ever moves *below-threshold* work to serial, and leaves everything `>= PARALLEL_MIN_WORK` splitting exactly as `main` does today. So relative to the current unguarded base there is **zero regression at any shape** — every shape either improves (below threshold) or is byte-for-byte the same code path (at/above). The only imperfection on this box is that shapes in ~[1M, 2.5M] still fork and lose 0.68–0.98× as they already do on `main`; the guard leaves that (pre-existing) money on the table rather than creating a new regression. I deliberately did **not** retune the constant to this laptop: it is a defensible, evidence-backed value pinned by a test and an evidence table, and overfitting it to a single dev box would be worse than a conservative shared default. If the CPU-EP owner wants the constant re-derived on the reference hardware, that is a separate, non-blocking follow-up. ### No large-shape regression Explicitly checked, because the risk of a "don't fork below N" heuristic is that N is wrong for some *other* shape: every shape `>= PARALLEL_MIN_WORK` takes the identical split path as base (the guard's `parallel` branch is unchanged), so nothing large got slower. Confirmed above — the 1.5M–3.1M rows behave as they do on `main`. ## Correctness The guards change **scheduling only, never arithmetic**, so both routes must agree **bit-for-bit** (compared on `to_bits()`), for f16 and bf16. All pass on this box: - `both_routes_agree_bit_for_bit` / `bf16_routes_agree_bit_for_bit` — spans `m > MAX_MC` so the serial multi-block indexing is actually exercised. - `half_gemm_declines_to_split_work_below_the_crossover` — proves the route via a thread-local counter, both directions. - `the_crossover_itself_splits_and_one_mac_below_it_does_not` — exact boundary. - `a_single_block_is_never_split_however_large_the_operand`. - `half_gemm_parallel_threshold_matches_the_measured_crossover` — pins the constant. ## Validation (reviewer) - `cargo test -p onnx-runtime-ep-cpu --lib` — **1352 passed, 0 failed, 17 ignored** (stacked on #1183; the 6 new guard tests account for the delta). - `cargo clippy -p onnx-runtime-ep-cpu --all-targets -- -D warnings` — clean. ## Scope This does not move the `K=N=2048` plugin benchmarks (those shapes are ≥4.2M MACs, far above the threshold). The win is confined to half GEMMs below ~1M MACs, which is the range this kernel actually serves post-#1140. --------- Co-authored-by: Roy <roy@squad.local> Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
## What this is This is #1159 (`fix(cuda): run Muse vision encoder natively`) rebased onto current `main` (`2d09da4e`), plus a follow-up commit adding real-hardware GPU numeric coverage. Opened as a replacement branch because the original `split/*` stack carries metadata that makes `gh pr merge` 403 even after the base PR closes. Closes #1159. ## The real bug Running the Muse vision encoder natively on CUDA exercised code paths the native EP did not implement, so the encoder could not run without falling back. The diff fixes four concrete gaps in `onnx-runtime-ep-cuda` + the session executor: 1. **f16/bf16 Attention with an f32 additive mask.** Half-precision attention previously rejected a floating mask whose dtype did not match Q/K/V. Vision encoders feed an f32 additive mask into an f16 attention; the fix accepts an independent f32 additive mask for half attention. 2. **Typed f16/bf16 SkipLayerNormalization.** The kernel was f32-only and errored on `Float16`/`BFloat16` input. Now typed. 3. **i32 integer elementwise / Range / Sum-Max-Min reductions.** Integer Add/Sub/Mul/Div/Min/Max, `Range`, and `ReduceSum/Max/Min` on `Int32` were "not yet implemented". Now implemented. 4. **Compress output-shape inference + Bool shape-input support** in the session executor's dynamic-shape path, so a data-dependent `Compress` resolves its output shape. ## How it was verified (real hardware: RTX 4060, CUDA 13 runtime) - Built `cargo build -p onnx-genai-cli --features native-cuda` clean on current `main`. - Ran the changed kernel modules under `--features gpu-tests -- --test-threads=1` (GPU-gated tests actually execute, not silently ignored): **59 passed** in `elementwise/range/reduce/normalization/standard_attention`. - The pre-existing tests for the new i32/typed-skip-layernorm kernels only asserted NVRTC source strings / dispatch planning — **no GPU numeric coverage**. Added `claim_probes` modules that allocate real device memory and check numerics on the GPU for claims 1–3 (i32 elementwise, i32 Range, i32/i64 reductions, typed f16/bf16 SkipLayerNorm). Claim 1 has a device-level test. All pass on the 4060. - **Falsified every claim.** For each fix I re-introduced the pre-fix bug, rebuilt, and confirmed the matching test goes RED with the exact pre-fix error (e.g. `Add with A dtype Int32 not yet implemented`, `input dtype Float16 (this slice is f32-only)`, `Range for dtype Int32 not yet implemented`, `ReduceSum with input dtype Int32 not yet implemented`, `Float16 attention should accept an f32 additive mask`, and Compress resolving to `Some([])` instead of `Some([[2,3]])`), then restored to GREEN. ## Honest caveats - **No Muse model was available locally**, so I could not run the full vision encoder end to end. The local CUDA VLM fixtures' vision "encoders" are trivial MatMuls that do not route the changed ops. Verification is therefore at the kernel-numeric + shape-inference level, with every claim independently falsified — not a full-model E2E. - The CI **benchmark-regression** comment on #1159 (CPU `matmul` +131%, gather, sampling) is **not caused by this PR**: the diff touches only CUDA kernels + the session executor's Compress/Bool shape inference, with no CPU-EP matmul/gather/sampling path. Most likely runner noise or a base/head build-config difference. - The engine E2E test `native_cuda_pipeline_decoder_parity` fails with a **pre-existing** "full CUDA mask reservation size overflow" in `native_decode/cuda.rs` — that file is byte-identical to `main`, so it is not a regression from this PR. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> --------- Co-authored-by: justinchuby <223556219+Copilot@users.noreply.github.com> Copilot-Session: 0190e2eb-abe4-451f-b36d-44a035a99b7e
…aces #1195) (#1253) ## What this is This is #1195 (`fix(runtime): infer Compress output shapes for large conditions`) rebased onto current `main` (which now includes #1251, the replacement for its stacked base #1159), plus a follow-up commit that adds the falsifying test the original was missing. Opened as a replacement branch because the original `fix/compress-condition-shapes` was stacked on a `split/*` branch, which makes `gh pr merge` 403. Closes #1195. ## The bug A dynamic `Compress` output shape is the count of selected entries, so the executor must read the boolean condition tensor to size the output. The dynamic-shape reader (`resolve_node_outputs`) admitted a shape input only if it passed `bounded_shape_input`, whose cap is `MAX_SHAPE_DATA_ELEMS = 1024` — sized for scalar/vector shape metadata. A vision graph whose `Compress` condition spans a full image is longer than 1024, so it was rejected, `dynamic_output_shapes` received a `None` condition, and the run failed with `UnresolvedShape { op: "Compress" }` (a 500). ## The fix - `geometry.rs`: add `bounded_compress_condition` — a 1-D `Bool` tensor up to `MAX_COMPRESS_CONDITION_ELEMS = 1 << 20`. - `dispatch.rs`: route the `Compress` condition input (default domain, `op_type == "Compress"`, input index 1) through a new `compress_condition_i64` that uses the larger bound, while every other shape input still goes through the ordinary `shape_input_i64`. So the one operand `Compress` genuinely needs is materialized, and an unrelated large tensor still is not. ## Test-discipline gap I closed The PR shipped `compress_condition_allows_image_sized_boolean_vectors`, which only asserts the `bounded_compress_condition` **predicate**. Reverting the `dispatch.rs` routing (the behavioural half of the fix) leaves that test green — it cannot catch a regression of the fix. I added: 1. `compress_runs_with_image_sized_condition_including_empty_and_full` — builds a `Compress` graph whose `Bool` condition initializer is `MAX_SHAPE_DATA_ELEMS + 500` (1524) elements and runs it end to end on the CPU EP, asserting the output shape resolves and exactly the selected rows are gathered. Covers the all-false (empty output) and all-true (full output) boundaries through the real resolve + kernel path. 2. `dynamic_output_shapes_compress_boundary_counts` — the selected-count sizer at its boundaries: condition longer than the axis (clamped), shorter than the axis, all-false, all-true. ## Falsification (revert -> RED -> restore -> GREEN) Reverting only the `dispatch.rs` routing back to `shape_input_i64` (keeping the helper + method so it stays a clean behavioural RED, not a compile error) makes `compress_runs_with_image_sized_condition_including_empty_and_full` fail with exactly: ``` UnresolvedShape { value: "y", op: "Compress" } ``` while the shipped predicate-only test and the counting test stay green (confirming they do not cover the routing). Restoring the routing returns all 4 Compress tests to green. ## Validation - `cargo test -p onnx-runtime-session --lib`: **185 passed, 0 failed** on current `main`. - Compress subset: 4 passed (shipped predicate test + shipped counting test + my 2 new tests). Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> --------- Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> Copilot-Session: 0190e2eb-abe4-451f-b36d-44a035a99b7e
…aces #1248) (#1254) Replacement for #1248, rebased onto current `main` after #1197/#1224/#1236 landed as squashes (the stacked base diverged). The original functional commit plus two added tests, re-validated on top of the merged stack. ## Symptom A model driven with flat tool names (`bash`, `edit`, `glob`, `grep`, `read`, `write`) emits a namespaced call such as `glob.glob`, which the client rejects: it can only dispatch a name it offered. ## Cause (as claimed by the PR, and what I could and could not confirm) The caller offers flat tool *definitions* over the API. The instruction text the model reads is rendered by the **package's own chat template**, which the PR says assumes namespaced names, takes `name.split('.')[0]`, and renders `# Valid recipients: "self", "glob.*", ..., "user".` — so the model follows a grammar that is wrong for a flat offer and writes `glob.glob`. **I could not inspect that template with my own eyes.** It ships inside the model package (Muse Glimmer / harmony-style), which is not in this repository or in my local model set; searching the repo and local packages for the `split('.')` / "Valid recipients" grammar found nothing. What I *can* confirm from the server side corroborates the addressing scheme: the ATEM parser recognises `to=functions.<name>` channels and `<atem:invoke name="...">` blocks (`extract_atem_tool_calls`, `ATEM_USER_CHANNEL = "to=user<|message|>"`), so a namespaced recipient like `functions.glob` producing a call named `glob.glob` is consistent with the diagnosis. **Template fix vs dispatch fix — my conclusion:** resolving at dispatch is the right layer, and this is a considered choice, not merely the first available. The template belongs to the model author and is rendered in the jinja sandbox; the server cannot rewrite an arbitrary third-party package's template at runtime, and the namespaced grammar is not universally wrong (it is correct for a caller that offers namespaced tools — only a flat offer breaks it). The server is the **only** component that sees both the caller's offer and the model's emitted call, so it is the only place that can reconcile them without editing model files. A template correction for flat offers would still be worthwhile in the package, but it cannot be the server's runtime responsibility. ## Fix When a parsed call names a tool that is not on offer but whose **final `.`-segment exactly equals an offered tool's whole name**, the call is resolved to that offered tool; anything else is left exactly as the model spelled it. Applied at the three call sites that already have the request in scope (buffered, streaming, and the logprob path), via `ParsedAssistantOutput::aligned_to` / `align_tool_calls`. **Ambiguity is impossible by construction, not merely avoided:** the resolution target is always the single final segment string, and it resolves only if that exact string is an offered (unique) name — there is never a choice between two candidates. When two offered tools merely share a final segment (`fs.read`, `net.read`) and the model writes a third namespace (`svc.read`), the bare segment `read` is offered by neither, so the call is left alone for the caller to reject. ## Verification (mine, on this branch atop the merged stack) - **216 lib + 31 http + 11 bundle tests pass**; `cargo clippy -p onnx-genai-server --all-targets -- -D warnings` clean. - **Falsified the positive claim** (revert → RED → restore → GREEN): neutering `align_tool_calls` to a no-op made `a_namespaced_call_resolves_to_the_offered_tool` and the new full-parse-path test return `glob.glob` instead of `glob`. - **Confirmed the negative tests have teeth, and characterised their limits honestly:** a no-op revert leaves the negatives green (they guard against *over*-resolution, not absence). Dropping the exact-offer guard makes `an_offered_name_is_left_alone` go RED (`fs.read` → `read`), which is exactly what that test protects. - **Added tests:** `a_shared_final_segment_is_not_resolved_arbitrarily` (the ambiguity case above) and `a_namespaced_atem_call_dispatches_through_the_full_parse_path`, which drives a raw `<atem:invoke name="glob.glob">` turn through the real `parse_assistant_output` → `align_tool_calls` path and asserts the dispatched name is `glob`. - **End-to-end note:** I could not reproduce the OpenCode/Muse path — no local model emits ATEM namespaced calls — so the end-to-end evidence is the full-parse-path unit test above (the server's actual code path), clearly labelled as such rather than an HTTP round-trip. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> --------- Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> Copilot-Session: 0190e2eb-abe4-451f-b36d-44a035a99b7e
… (1.66x -> 1.36x vs ORT) (#1218) ## What this changes `erf_ps`'s big-branch `R` polynomial and the `exp_ps` polynomial it feeds are switched from Horner's rule to **Estrin's scheme**. Horner is a six-deep dependent FMA chain; Estrin regroups it as `a^3·(c0 a^3 + c1 a^2 + c2 a + c3) + (c4 a^2 + c5 a + c6)`, three deep for two extra multiplies, whose halves the out-of-order engine runs in parallel. Only the two chains actually on `erf`'s critical path are converted; the small branch stays Horner (it runs in parallel and is off-path). Rebased onto current `main` so it **preserves #1227** (`map_ps`'s `#[target_feature(enable = "avx2,fma")]`). The original branch predated #1227 and its raw diff would have moved that attribute back onto `map_bias_ps`; that hunk is dropped. Only `simd_activations.rs`'s erf/exp polynomials and one accuracy test change here. ## Numerics — measured, not assumed (reviewer) Estrin reassociates f32 FMAs, so it is **not** bit-identical to Horner by construction. The squad reviewer measured the actual error with a two-build dump-and-diff (base Horner build vs this Estrin build, identical input grid of 404,045 points spanning `[-6,6]` densely, the `0.921875` split boundary, large `|x|`, denormals, and NaN/±inf), on **Intel Core i7-13800H, Windows 11**: - **Estrin vs Horner: max difference = exactly 1 ULP.** 4,401 of 404,043 finite points (1.09%) differ, every one of them by a single ULP; max absolute diff 5.96e-8. Nothing moves by more than one ULP anywhere in the domain. - **vs f64 `erf` truth:** max |Horner − erf| = 5.77e-8, max |Estrin − erf| = 6.55e-8 (worst at x≈0.933, just past the split). Both within ~1 ULP of the correctly rounded result; Estrin is at most ~0.8e-8 looser than Horner and never worse than 1 ULP from it. - **Special values are all correct and identical between the two:** large `|x|` (7,8,10,20,50,100,1000,1e30,f32::MAX) saturate to ±1.0; denormals produce bit-identical tiny outputs; NaN → NaN; +inf → +1.0; −inf → −1.0. This is a conscious, documented 1-ULP change, not a hand-wave. The accuracy gate `erf_reassociation_costs_no_accuracy` (worst scaled error pinned at 5.97e-8 = 2^-24 over `[-6,6]`) **passes** on this box, so a future regrouping that widens the envelope fails instead of being absorbed by the looser `ERF_BOUND`. ## Performance — measured on this box, ORT-independent The original body quoted "1.66× → 1.36× vs ORT" (≈ −17.7% kernel time). I could not reproduce that specific figure: no pinned ORT reference on this Windows box, and it is different hardware. Instead I timed `erf_avx2` directly, same machine, two builds, single-thread, min-of-200 × 9 rounds over a 1,048,576-element mixed-range buffer: | build | per-round min (ms) | best | | --- | --- | ---: | | Horner (base) | 0.6346–0.6348 (8/9 rounds) | 0.6346 | | **Estrin (this PR)** | **0.6042–0.6043 (8/9 rounds)** | **0.6042** | The ranges do not overlap, so the direction is solid: **Estrin is ~5.0% faster (1.05×)** on the erf kernel here. That is a real win but smaller than the −17.7% originally claimed — reported honestly as the measured kernel-internal delta rather than an unverified ORT ratio. The mechanism (a latency-bound chain freed by shorter dependency depth) is consistent with a positive but hardware-dependent magnitude. ## Validation (reviewer) - `cargo test -p onnx-runtime-ep-cpu --lib` — **1353 passed, 0 failed, 17 ignored** (base 1345 + #1183/#1149 + this test), including the accuracy gate. - `cargo clippy -p onnx-runtime-ep-cpu --all-targets -- -D warnings` — clean. - No MLAS feature involved; `map_ps`'s target_feature (#1227) verified intact in the merged result. Co-authored-by: justinchuby <223556219+Copilot@users.noreply.github.com> Co-authored-by: resch <resch@users.noreply.github.com>
…ionStore (#1255) ## What & why The engine kept **two parallel session implementations** selected by `if self.decode_backend == EngineDecodeBackend::Native`, one over `self.native_sessions` and one over `self.sessions`. The *policy* around a session — the "session {id} not found" lookup, the rewind bound-checks and their exact error strings, the checkpoint arithmetic, and `truncate`/`clear` — was **backend-independent but duplicated verbatim** across the two arms (the `checked_sub` bound check and its message in `rewind_session_by`, the `position > current` check in `rewind_session_to`, etc.). This is the same class of drift risk that #1170 closed for the KV *length* layer with `KvPrefixStore`. This PR extracts that policy into one place and routes both backends through it, following the `KvPrefixStore` precedent rather than inventing a new pattern. ## The seam New module `crates/onnx-genai-engine/src/engine/session_state.rs`: - **`SessionStore` trait** — mirrors `KvPrefixStore`. Backends implement `logical_len` (via `SessionLen`), `validate_rewind`, `rewind`, `reset`, `close`, with **identical method names/signatures** on both sides, exactly as the `NativeKvStore`/`OrtKvStore` adapters already do. - **Free policy functions** `require_session`, `checkpoint`, `token_count`, `rewind_by`, `rewind_to`, `reset`, `close` own everything that is backend-independent: the lookup + not-found text, both rewind bound-check strings, `SessionPosition::new(len)` checkpoint arithmetic, and the dispatch into the single backend truncation method. - **Adapters in `runtime.rs`**: `OrtSessions`/`OrtSessionsRef` (always compiled) and `NativeSessions`/`NativeSessionsRef` (`#[cfg(feature = "native-backend")]`) implement `SessionStore`. The six public methods (`checkpoint_session`, `rewind_session_by`, `rewind_session_to`, `reset_session`, `close_session`, `session_token_count`) now select an adapter and call `session_state::*`. The `require_ort_backend(...)` guard stays in the ORT arm. Error text is now authored **once**. Rendered strings are byte-identical to the originals (only internal variable names differ, e.g. `{id}` vs `{session_id}`), so behaviour is unchanged. No test asserted these strings before; the ones that assert *backend* rewind errors (`failed_rewind_of_*`) are untouched. ## Asymmetries I deliberately kept (and why) Forcing these into false symmetry would be worse than the small duplication that remains, so — mirroring how `KvPrefixStore`'s `reset` is legitimately different per backend — they stay backend-specific behind the trait: - **`validate_rewind`**: ORT validates the target *and* draft/speculative state and paged-KV page ownership (`validate_target_state_rewind_to_len`) so a rejected rewind leaves the session untouched; the three `failed_rewind_of_*` tests pin this. Native's `validate_rewind` is a deliberate `Ok(())` — it always admits and clamps against its in-process decoder length. These are genuinely different objects: ORT owns paged KV + a draft runner; native owns one persistent in-process decoder. - **`rewind`/`reset`/`close` mechanism**: ORT does validate → `scheduler.complete` → remove → `rewind_target_state_to_len` → reinsert; native truncates its token vec, rewinds/resets the single persistent decoder (gated on `native_active_session`) and touches LRU `last_access`. - **`create_session`** is left backend-specific and is **not** routed through the seam: it is object *construction*, not policy, and the two session value types (`EngineSession` vs `NativeSessionState`) are genuinely different. Unifying it would be a forced abstraction with no shared policy to host. - **Token `truncate` for ORT is intentionally *not* lifted** into the shared policy. It lives inside `rewind_target_state_to_len` (`kv_bridge.rs`), which is **shared with the speculative-decode hot path** (`speculative/mod.rs`). Lifting it would fracture that path for no gain. The shared policy therefore owns the logical-length **bound check** (truly common) and delegates the truncation. ## DRY guard (mandatory, per #1170) Two-part, mirroring `ReusedPrefix` + `KV_REWIND_CALLERS`: 1. **Compile-time** `CheckedPosition` newtype with a private field. It is only ever constructed by `session_state::rewind_to` after the bound check, and it is the required argument to `SessionStore::rewind`/`validate_rewind`. A future backend therefore *cannot* rewind KV without going through the one bound check — there is no other way to obtain the token. 2. **Test-time tripwire** `the_rewind_bound_check_lives_only_in_the_shared_policy` scans the crate sources for the `"cannot rewind session"` fragment and fails if it appears anywhere except `engine/session_state.rs`. A third backend that open-codes a copy turns the suite red. (`fork_session` uses the distinct `"cannot fork session"`, so there is no collision.) ## Performance These are **per-turn / per-request** session ops, not the per-token decode hot path (`generate_incremental_with_callback`). The trait dispatch happens once per rewind/checkpoint/reset, off the hot path, so the dynamic call cannot show up in decode throughput. No hot-path code changed; no per-token dispatch was introduced. (Restructured, not measured — because there is nothing on the hot path to measure.) ## Falsification (the important part) I broke the shared policy deliberately — inverted the one bound check in `session_state::rewind_to` (`requested > current` → `requested < current`) — and confirmed **both** backends go red from that single line: - **ORT** (default backend): `failed_rewind_of_windowed_evicted_position…`, `failed_rewind_of_ort_owned_kv…`, `failed_rewind_of_runner_backed_state…` — all 3 FAILED (they now hit the bound check before their backend validation). - **Native** (`--features native-backend`): `native_session_rewind_by_truncates_logical_length` FAILED with `cannot rewind session 1 to token 3; current length is 5`. Reverting the one line restored both to green. A break in the shared path turns *both* backends red, which is the evidence that the extraction actually unified them rather than leaving two look-alikes. I added `native_session_rewind_by_truncates_logical_length` because there was previously **no** integration test exercising a *valid* explicit rewind on the native backend (the existing native "rewind" test drives prefix-reuse, a different path) — this closes that gap and gives the native half of the falsification. ## Verification - `cargo test -p onnx-genai-engine --lib` → **414 passed, 0 failed, 1 ignored**. - `cargo test -p onnx-genai-engine --features native-backend` → all targets pass **except** the pre-existing `pipeline_load_rejects_kv_page_pool_before_fixed_state_when_host_budget_cannot_fit_either` (a host-RAM KV-budget test — confirmed failing identically on a stashed baseline; it does not touch session code). - `cargo clippy -p onnx-genai-engine --all-targets -- -D warnings` and the same with `--features native-backend` → clean. - **CUDA (RTX 4060, driver 591.55):** `cargo test -p onnx-genai-engine --features native-backend,cuda --test native_engine -- --test-threads=1` ran on the GPU (not ignored): **16 passed, 1 failed, 0 ignored**. Passing includes the new `native_session_rewind_by_truncates_logical_length`, `native_sub4_cuda_fallback_generates_coherent_decode`, and `qwen15b_native_decode_locks_accurate_near_tie_token` (real CUDA decode). The 1 failure, `engine_native_cuda_matches_cpu_tokens` ("full CUDA mask reservation size overflow" during native decoder load), is a pre-existing CUDA mask-workspace sizing issue on this 8 GB card — confirmed failing identically on a stashed baseline, unrelated to session state. ## Behaviour note `close_session`/`reset_session` now run the existence check (via `require_session`) before `scheduler.complete`. Previously ORT `close` called `scheduler.complete` before its remove-based existence check. This is not observable through any error (`scheduler.complete` on a missing id is benign) and `reset` already checked existence first. Do not merge — please review. Co-authored-by: justinchuby <223556219+Copilot@users.noreply.github.com>
…-ORT headline) (#1259) ## What Records the correction of record for **#1218** (erf's Estrin-scheme reassociation), appended as **§32** to the CPU-EP-vs-ORT ledger — the same document and style where the #1226 (§30.2) and #1230 withdrawals live. #1218 landed as *"1.66x -> 1.36x vs ORT"* (~17.7% kernel-time reduction). On the review host that ratio is **unreproducible** — there is no pinned ORT reference build here, and the original number came from different hardware — so it is **withdrawn**. Stated plainly: it was **neither reproduced nor refuted**, because it was measured against a reference we do not run. ## What the doc now records - **The withdrawal**, with the reason (no ORT arm on this box), and a note that the overstated figure reached **no benchmark table, README, or source comment** — only the (unrewritable) squash-commit subject. - **The measured number:** ORT-independent same-binary A/B on `erf_avx2`, single-thread, min-of-200 × 9 rounds, 1M elements, on **Intel Core i7-13800H** — Horner 0.6346 ms → Estrin 0.6042 ms, non-overlapping ranges → **~5.0% (1.05x)**. Real, but a fraction of the implied ~17.7%. - **The durable correctness fact, stated in as many words:** this reassociation is **NOT bit-identical**. Estrin differs from Horner by **at most exactly 1 ULP** (1.09% of 404k points, max abs 5.96e-8); vs f64 truth Horner 5.77e-8 / Estrin 6.55e-8 (both within ~1 ULP of correctly-rounded erf); large |x| → ±1.0, denormals bit-identical, NaN → NaN, ±inf → ±1.0. Names the gate test **`erf_reassociation_costs_no_accuracy`** as the tripwire. ## Validation Documentation only — no source, no benchmarks re-run. No markdown linter/build exists in-tree, so there is nothing to build; the change is prose appended to an existing report. Refs #1218. Companion to the separate threshold-tuning issue filed for #1149. Co-authored-by: justinchuby <223556219+Copilot@users.noreply.github.com>
CI's format lane had fallen behind — `cargo fmt --all -- --check` was failing on `main` at `7d669c76`. The drift came from several changes landing in the same window: the session-state unification (#1255), the server stack (#1197/#1249/#1250/#1254), and the native CUDA vision kernels (#1251). **No logic changes.** This is `cargo fmt --all` and nothing else — 7 source files, plus two `.squad/decisions/inbox` notes that were already sitting uncommitted in the tree. Verified locally: `cargo fmt --all -- --check` exits clean on this branch. Co-authored-by: justinchuby <223556219+Copilot@users.noreply.github.com>
…ersedes #1259 framing) (#1262) ## What Revises **§32** of the CPU-EP-vs-ORT ledger to reframe the #1218 erf/Estrin result as a **hardware-dependent measurement**, not a withdrawal. Per owner clarification: #1218 was originally measured on a **many-core AMD server part**; the reviewer's re-measurement was on an **Intel Core i7-13800H** laptop. The difference between the original ~1.36x-vs-ORT and the reviewer's ~5% is a **microarchitecture / core-count / reference-baseline difference, not an overstatement**. The previous version of §32 (from #1259) framed it as a withdrawal in the #1230 style; that was the wrong framing and this corrects it before it misleads a later reader. ## What §32 now says - **32.1** presents both measurements as **peers**, each with hardware named — original ~1.36x vs ORT on a many-core AMD server part; reviewer ~5.0% (0.6346→0.6042 ms, single-thread, min-of-200 × 9 rounds, 1M elements) on i7-13800H via ORT-independent same-binary A/B — and states explicitly they **do not conflict** (different microarch, core count, and reference). - **32.2** draws the general, reusable lesson: CPU perf numbers here are hardware-dependent, and any recorded number must carry its **hardware, thread count, and reference baseline** or it is not interpretable later. Flags (without fixing) that some older entries lack that metadata. - **32.3** keeps the **hardware-independent** 1-ULP numerics record unchanged, and still states plainly that the reassociation is **not bit-identical**. ## Validation Documentation only. No source or benchmarks touched; no markdown linter/build exists in-tree. Supersedes the framing introduced in #1259. Refs #1218. Co-authored-by: justinchuby <223556219+Copilot@users.noreply.github.com>
…red native-CUDA parity test) (#1264) ## What this fixes `main` currently has a **red end-to-end native-CUDA parity test**: `native_cuda_pipeline_decoder_matches_cpu_token_ids` fails during `DecodeCudaState::new()` with `full CUDA mask reservation size overflow`. ### Root cause `DecodeCudaState::new` computed the mask *reservation* extent unconditionally: ```rust let full_mask_bytes = batch .checked_mul(capacity.max_len) // <-- capacity.max_len == usize::MAX .and_then(|e| e.checked_mul(size_of::<i64>())) .context("full CUDA mask reservation size overflow")?; ``` For a model with **no `max_sequence_length` metadata**, `resolve_cuda_kv_capacity` returns `capacity.max_len == usize::MAX` — the deliberate *"unbounded until growth fails"* sentinel (added in #367, asserted by `cuda_kv_capacity_without_metadata_is_unbounded_until_growth_fails`). So the multiply overflows and aborts construction — **even on the non-VMM decode path, whose mask binding is a plain `[batch, max_len]` allocation that never reserves the full logical extent and never reads `full_mask_bytes` at all.** The `tiny-gemma4-vlm-cuda` fixture used by the parity test is exactly such a metadata-less model, which is why the test is red. ### First bad commit Bisected to **#682 (`9fd7ee56`, "Reserve CUDA KV with VMM committed ranges")**, which introduced the unconditional `full_mask_bytes`. The `usize::MAX` sentinel predates it (#367). My earlier suspects **#1221** (two-pass workspace reservation) and **#1228** (KV mask growth) were both **cleared with evidence** — they are much later and touch different code (#1228 only added the `grow_vmm_mask_in_place` shape assertion). ### The fix Compute `full_mask_bytes` (and the `mask_committed_bytes` it feeds) **only inside the `kv_commits_on_demand` (VMM) branch that actually consumes them.** The `usize::MAX` sentinel is intended design and is left untouched. Behaviour is **byte-identical on every previously-working path**: the non-VMM branch always committed `mask_bytes`, and `seq_major_fixed` implies `kv_commits_on_demand`, so `full_mask_bytes` was only ever live on the VMM path. ### Captured operands at the failure (RTX 4060, non-VMM path) ``` batch = 1 capacity.max_len = 18446744073709551615 (= usize::MAX) source = "unbounded (model.max_sequence_length unavailable)" bytes_per_token = 40 kv_commits_on_demand = false (non-VMM path -- full_mask_bytes is dead here) ``` ## Verification (RTX 4060, driver 591.55, CUDA 13 runtime) - **Regression test** `native_cuda_decoder_constructs_with_unbounded_metadata_less_capacity` (CUDA-gated, `ONNX_GENAI_RUN_CUDA_SMOKE=1`): constructs a metadata-less non-VMM decoder and asserts `new()` succeeds + `capacity.max_len == usize::MAX`. - **GREEN** with the fix. - **RED** with the exact `full CUDA mask reservation size overflow` when the computation is hoisted back out of the branch (falsified by revert). - **Parity test** `native_cuda_pipeline_decoder_matches_cpu_token_ids`: **RED on `main` → GREEN with this fix.** - Engine lib suite (`--features "cuda,native-backend" -- --test-threads=1`): **548 passed, 0 failed, 2 ignored** (default). ## Scope note — the VMM half is NOT fixed here This fix makes only the **non-VMM** path correct. On the **VMM** path (`ONNX_GENAI_CUDA_VMM=1`) a metadata-less model would still overflow for real, because `full_mask_bytes` — and `full_vmm_kv_allocation_bytes(..., capacity.max_len)` — become *live* values that genuinely need a concrete bound. Resolving the unbounded sentinel to a real capacity before the VMM path can reserve is a design decision, not an obvious mechanical fix, so it is **filed as a separate issue** rather than papered over here (see linked issue). ## Not covered by this PR Running the lib suite with `ONNX_GENAI_RUN_CUDA_SMOKE=1` surfaces **3 pre-existing failures** (`native_cuda_capture_replay_...`, `native_cuda_symbolic_batch_...`, `native_cuda_symbolic_total_seq_...`) that fail with `cannot resolve model.io.token_input ... 3 ports match`. These are a pre-existing test-harness `model.io` ambiguity unrelated to this change (this PR does not touch graph building or io resolution); they only surface because CI never sets the smoke env. Co-authored-by: justinchuby <223556219+Copilot@users.noreply.github.com>
…nt ORT activation deferrals) (#1235) ## What this changes Adds native CPU-EP kernels for **`Celu`** (opset 12) and **`Mish`** (opset 18), which had no kernel at all and were therefore silently handed to ORT's CPU EP, and vectorises **`Log`** (f32), which was still taking a scalar `libm` call per element. Celu/Mish are wired through all four capability filters (registry, `PHASE1_OPS`, dtype table, shape-inference table, plugin descriptors) and added to the two activation-coverage tests. ## Behavioural claim — verified independently (reviewer) The headline "removes the last silent ORT activation deferrals" is a behavioural claim, so I falsified it rather than trusting it, on **Intel Core i7-13800H, Windows 11**: - **The deferral was real and silent.** `GetCapability` runs three fail-closed filters; an op absent from the shape-inference table (Filter 2) or the dtype descriptor list (Filter 3) is handed to ORT with no diagnostic. Celu/Mish had no kernel, so they never reached `supports_op`. - **Falsification (RED/GREEN).** With the PR applied, `activation_and_norm_ops_clear_every_capability_filter` passes. Removing just the Celu/Mish entries from the shape-inference table (`compute.rs`, Filter 2) turns it **RED** with `::Celu: no shape rule (shape filter declines it) … would be handed to ORT's CPU EP` for both ops; restoring returns it to **GREEN**. The test genuinely guards the deferral. - **"The last" is defended structurally, not by spot-check.** The coverage test's `OWNED` list now enumerates the full activation/normalisation family and holds every member to both the shape and dtype filters; the companion `every_registered_op_has_a_shape_rule_or_is_a_known_gap` enumerates all registered ops, and its remaining shape-table gap (data-dependent, internal-fusion, and inferrable-but-unwritten ops) contains **no activations**. `Log` was never a deferral — it already had a scalar kernel — so this PR is a perf change for it, not a coverage change. ## Correctness The kernels' own tests (dense sweeps vs the scalar reference, subnormals, and the full IEEE special set) pass; the scalar reference, the f64 path, and the AVX2 path are held to agree, including NaN propagation (`f32::max`/`min` drop NaN, so each kernel guards it to match ORT). Full `cargo test -p onnx-runtime-ep-cpu --lib` = **1362 passed, 0 failed, 17 ignored** (base 1353 + 9 new). `cargo clippy --all-targets -D warnings` clean. `map_ps`'s `#[target_feature]` (#1227) verified intact. ## Performance — measured on this box, ORT-independent The gaps doc's ratios were taken on an **AMD EPYC 9V74** against ORT. I have no ORT reference on this host, so I measured the ORT-independent internal win instead: the AVX2 kernel vs the scalar-native fallback (which is what runs on non-AVX2 hosts and small tensors), on **Intel Core i7-13800H (14C/20T), Windows 11**, single thread, 1,048,576 elements, min-of-200 × 9 rounds. | kernel | scalar-native (ms) | AVX2 (ms) | speedup | | --- | ---: | ---: | ---: | | `Log` (positive domain) | 2.89 | 1.14 | **~2.54×** | | `Celu` (alpha=1, mixed) | 2.77 | 0.68 | **~4.07×** | | `Mish` (mixed) | 17.82 | 2.50 | **~7.13×** | These are hardware-dependent (see `docs/performance/CPU_ACTIVATION_GAPS.md` §32 lesson) and are **not** the same quantity as the doc's ORT-relative ratios — mine compare AVX2 against the native scalar path on this machine; the AMD EPYC figures compare native against ORT. Presented as peers, each with its hardware named. Both agree on direction: the native kernels are a clear win. --------- Co-authored-by: Resch <resch@squad.local> Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
…1201) ## What Adds a CPU task runtime (`task_runtime`) with an adaptive-spin native pool, plus an SMT/hybrid-aware `core_topology` module. This is the *foundation* PR: it introduces the machinery but wires no production kernel to it yet (RoPE/Softmax routing is #1202). `matmul_nbits` only gains a `decode_thread_budget()` getter. ## Design - **One pool per machine, not per subsystem.** Inside an ORT plugin-EP compute call, fan-outs route to `KernelContext_ParallelFor` (`Backend::Host`) and the native pool stays parked, so the new pool never spins alongside ORT's intra-op pool. Outside ORT it uses the native adaptive-spin pool (`Backend::Native`); tiny work runs `Backend::Serial`. - **Width is capped to the physical cores the process may run on** (min floor 8), and reuses the decode thread budget so one knob sizes the engine. - **Adaptive spin** grows the spin window while dispatches keep arriving and decays to a park when they stop. ## Review notes (validated on this hardware) **Hardware: Intel i7-13800H, 14 physical / 20 logical (6 P-cores×2 SMT + 8 E-cores), Windows.** Baseline reference = origin/main. - **Correctness / thread count.** In this PR no production kernel dispatches to the pool (rayon remains everywhere), so a running inference gains **zero** steady-state threads from it; the pool is lazily built on first use only. Verified there are no `task_runtime::` callers outside the module + tests. - **Bug found & fixed during review.** The Windows `GroupMask` offset in `core_topology` was computed as the address delta between two unrelated stack copies, so `GetLogicalProcessorInformationEx` detection returned `None` on this box. The SMT/physical-core cap then silently no-op'd and the pool resolved to **20 workers on 14 physical cores** — the exact oversubscription this exercise exists to remove. Fixed with `offset_of!`; detection now reports `core_count=14`, hybrid split `[2×6, 1×8]`, and the resolved width drops **20 → 14**. Added a non-vacuous Windows regression test (the pre-existing self-consistency test returned early on `None` and was vacuous on Windows). - **Dispatch latency reproduced** (release, 14 workers, 400 rounds), `tests/task_runtime_latency.rs`: | gap before dispatch | p50 | p90 | | concurrent sessions | p50 | p90 | |---|---|---|---|---|---|---| | 0 µs | 1.7 µs | 1.9 µs | | 1 | 2.0 µs | 2.2 µs | | 20 µs | 1.9 µs | 2.1 µs | | 2 | 1.9 µs | 2.4 µs | | 100 µs | 1.9 µs | 2.6 µs | | 4 | 2.3 µs | 2.7 µs | | 500 µs | 142.9 µs | 186.5 µs | | 8 | 2.5 µs | 3.2 µs | The point is the flatness: 0 µs and 100 µs cost the same, so a decode step no longer pays for having been idle; past the 500 µs spin ceiling the workers park and the cost returns to a futex wake, which is correct. `slot exhaustion: 0 of 9744 dispatches fell back to serial`. (These are on this Intel laptop; the PR's original 4.8 µs figures were on an AMD EPYC 9V74 server — peers, not a contradiction.) - **Tests:** `cargo test -p onnx-runtime-ep-cpu --lib` = 1395 passed / 0 failed (baseline main = 1353; +42 new). Task-runtime integration tests pass individually. Clippy `--all-targets -D warnings` clean. Note: the idle test skips on Windows (no per-process CPU accounting), and one lib timing test flaked once when the whole suite ran concurrently with the idle-window binary under full CPU load — it passed on every isolated and repeat run; correctness never depends on the pool (serial fallback). Merged onto latest main (resolved a benchmark-doc section-numbering conflict; the task-runtime writeup is §33). --------- Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
) ## Symptom Driving this model from OpenCode, a plain `hello` shows nothing at all for minutes, and no reasoning is ever displayed. Measured against this server: ``` first_visible_text = 280.4s total = 280.8s finish = stop ``` 280 seconds of complete silence on the wire. A client cannot tell that apart from a hung server. ## Cause #1249 stopped this package's private channel being served as an answer, which was right — but it did so by *discarding* it. The model addresses `to=self` while it works and only later opens `to=user`, so for the whole of the thinking there is nothing to send, and the connection simply sits there. Thinking was never the problem. Passing it off as an answer was. ## Fix The private channel is now reported on the channel that exists for it — `reasoning_content`, beside `content` rather than inside it — on the streaming delta and on a buffered message alike. Nothing moves into `content`, so the answer is byte for byte what it was. The gate keeps one transcript and asks the channel definitions what each has revealed since the last emission, so the two channels cannot disagree about where a token belongs. A model that declares no private channel reports no thinking rather than having one invented for it. Reasoning streams even on the tool-detection path, where the answer must be buffered until the turn is known not to be a tool call — that is exactly the agent case, and exactly where the silence was worst. ## Verification Same prompt, same server, after the change: ``` first_reasoning = 6.9s first_content = 7.8s total = 8.3s finish = stop answer: Hello! 👋 How can I help you today? ``` - 220 lib + 31 http + 11 bundle tests pass; clippy `--all-targets` clean. - Four new gate tests: thinking is reported as it is produced while the answer stays gated; a turn that never answers still reports what it was thinking; thinking resumed after a tool call extends rather than replaces; an unarmed gate reports none. - End to end from OpenCode: reasoning is displayed live, the tool calls dispatch, and the answer is unchanged. --------- Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> Copilot-Session: 0190e2eb-abe4-451f-b36d-44a035a99b7e
…#1022) (#1155) Closes #1022 ## What this is — and what it is not (verified by review, not restated) This adds the **engine-side seam, per-row router, capability predicate and byte-accounting** for continuous-batch device sampling. It does **not** change runtime behaviour today, and it makes **no throughput claim**. I verified both of those statements directly: - **No CUDA/native producer emits `BatchStepLogits::Device`.** The only construction of that variant and the only `impl DeviceRowLogits` in the tree are the test mock (`MockDeviceRowLogits`). So with every real backend, `device_available` is `false`, every row plans `Host`, and the demux is byte-identical to the pre-PR whole-batch copy. The per-row win is **latent infrastructure** that a future CUDA producer implements against the `DeviceRowLogits` trait; it is exercised here only by mocks. - Therefore the previous whole-batch host copy is genuinely the slow path this is *designed* to remove, but this PR does not yet remove it in production — it makes the routing possible and proves it correct. That is the honest scope. ## What changed 1. **Lifted the capability predicate** (`decode_loop.rs` → `processors.rs`). The two inline `let greedy_fastpath` / `let sampled_fastpath` bindings are now one `device_sampling_plan(chain, options, has_custom_sampler, greedy_supported, sampled_supported) -> DeviceSamplingPlan {Greedy|Sampled|Host}`. I diffed the lifted logic against the originals: it is condition-for-condition identical (empty chain, `top_logprobs`, custom sampler, portable chain, backend support). 7 no-GPU unit tests pin each disqualifier (`processors::device_sampling_plan_tests`). 2. **Per-row routing off one device-resident buffer** (`BatchStepLogits::Device(Box<dyn DeviceRowLogits>)`). The manager consults `device_sampling_plan` **per row**: device-portable rows call `sample_row` (4 bytes back), host-required rows call `copy_row_to_host` (full `[vocab]`). The seeded categorical path draws the request RNG (`rng.value_for(options)`) at the same point the host path would, so routing does not perturb output. 3. **Bytes-moved accounting** (`LogitsD2hStats` gains `rows_host_copied / rows_device_sampled / token_id_bytes`), reported by the manager when a device buffer is used. ## Verification (this review, RTX 4060 + default-feature host) - **Heterogeneous-batch tests are genuine and falsifiable.** `mixed_batch_moves_only_host_required_rows` (2 greedy device rows + 1 repetition-penalty + 1 `top_logprobs` host rows) asserts `rows_host_copied == 2`, `rows_device_sampled == 2`, `bytes == 2·vocab·4` (**not** `4·vocab·4`). `device_routing_matches_host_token_stream` drives a device-routed manager and an all-host manager to completion over greedy + **seeded** categorical + a host-only penalty row and asserts identical token streams. - **Falsification I ran:** injecting the exact "apply row 0's greedy plan to the whole batch" bug (`let plan = DeviceSamplingPlan::Greedy;` for every row) turns **both** tests RED (`mixed_batch`: `rows_host_copied` 0 vs 2; `device_routing`: token stream diverges / host buffer mis-sampled). Restoring returns green. The tests genuinely depend on per-row routing, not a shared config. - **Default-feature suites:** `onnx-genai-engine --lib` 424 passed / 0 failed / 1 ignored; `onnx-genai-ort --lib` 144 passed / 0 failed. - **GPU suite (`--features "cuda,native-backend" -- --test-threads=1`, RTX 4060):** compiles under CUDA; **558 passed / 0 failed / 2 ignored** (no smoke env). With `ONNX_GENAI_RUN_CUDA_SMOKE=1`: 555 passed / **3 failed** / 2 ignored — the 3 failures are the **pre-existing** `model.io.token_input` harness ambiguity (`native_cuda_capture_replay_...`, `native_cuda_symbolic_batch_...`, `native_cuda_symbolic_total_seq_...`), unrelated to this PR, which touches neither `native_decode` nor graph building. - `cargo clippy -p onnx-genai-engine -p onnx-genai-ort --all-targets -- -D warnings` — clean. ## Not exercised here The device kernel math is out of scope (covered by `device_sampler.rs` host-oracle parity). The end-to-end device win cannot be measured until a CUDA producer emits a `Device` buffer; no wall-clock number is asserted. --- _Reviewed and merged on local evidence (not CI) by the Copilot CLI agent on an RTX 4060. Verified the predicate lift is faithful, the heterogeneous-row tests go RED when per-row routing is removed, and the GPU suite shows no new failures._ Co-authored-by: justinchuby <223556219+Copilot@users.noreply.github.com>
## What Routes `RotaryEmbedding` and `Softmax` off Rayon and onto the CPU task runtime (#1201). Both previously sized their fan-out from `rayon::current_num_threads()` and split with `par_chunks_mut`, paying Rayon's park latency on every region and starting a second pool beside ORT's inside a plugin-EP call. They now go through `task_runtime::chunks_mut`, which lands on ORT's intra-op pool inside a compute call and on the native adaptive-spin pool outside. Grain policy changes from a static per-worker share to a dynamic-claim floor. ## Review notes (validated on Intel i7-13800H, 14 physical / 20 logical, Windows; baseline = origin/main) - **Correctness / bit-identity.** This is a pure-scheduling change (reordering *parallel* work), so output is unchanged. Verified by the existing bitwise kernel tests, which all pass: `rotary_*_bit_identical`, `rotary_*_bitwise`, `rotary_strided_cache_fallback_matches_contiguous_fast_path_bitwise`, and the softmax/log-softmax reference tests. The range-partition hazard is guarded by `task_chunking_covers_every_unit_exactly_once`, which drives the real partition through `task_runtime::chunks_mut` and asserts each unit is written exactly once. - **Reachability.** RoPE and Softmax now call `task_runtime::chunks_mut`/`chunk_runs_mut`; `rayon`/`par_chunks` is fully removed from both files — not an unwired placeholder. - **Cross-cluster conflict (activation kernels).** #1202's original commit `773d6a94` force-inlined `map_ps`/`map_bias_ps` with `#[inline(always)]` (no `target_feature`) to stop the RoPE/Softmax rewiring de-vectorizing `tanh_avx2`/`sigmoid_avx2` via a whole-crate inline-cost-model flip. The activation cluster has since landed the *competing* fix on main (`#[inline] + #[target_feature(avx2,fma)]`, mutually exclusive with `inline(always)`). Per the no-revert directive I kept **main's** version and dropped `773d6a94`, then **verified by disassembly** that the de-vectorization does not recur: with #1202's routing applied, `tanh_avx2` = 124 lines / 0 calls / 68 ymm ops and `sigmoid_avx2` = 127 / 0 / 69 — both fully vectorized, `relu`/`erf` inlined away. So the two clusters are compatible and #1202 no longer touches `simd_activations.rs`. - **Tests:** `cargo test -p onnx-runtime-ep-cpu --lib` = 1405 passed / 0 failed / 17 ignored. Clippy `--all-targets -D warnings` clean. - **Perf.** The park-latency win is inherited from #1201's reproduced dispatch-latency result (p50 1.9 µs flat across 0–100 µs gaps on this box). The PR's per-shape RoPE/Softmax ORT ratios were measured on an AMD EPYC 9V74 and are not re-run here — they need the `scripts/ort_ab` ORT session harness and are hardware-dependent by nature. Rebased onto latest main; benchmark writeup renumbered to fit after the erf section. --------- Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
Replacement for #1190, rebased onto current `main` (the original could not be squash-merged because #1210 is stacked on its head branch, which triggers GitHub's stacked-PR 403). Content is identical to #1190; only the base is refreshed. ## What this adds A 20-file explanatory wiki under `wiki/` (crate architecture, execution backends, EP contracts, memory model, metadata-driven runtime, request lifecycle, testing/verification, performance playbook, repository map). It repeatedly labels itself non-authoritative and points at `docs/` and source for verification. ## Accuracy review (the bar for docs in this repo) I spot-checked the load-bearing claims against `main` as it stands today (after the memory capability-split revert #1247, session-state unification #1255, `reasoning_content` #1257, CUDA mask-reservation fix #1264): - **Crate architecture** — the crate roster and layering match the actual workspace `Cargo.toml` members and `crates/` tree. The map is explicitly non-exhaustive and says so. - **Memory model (the highest-risk area)** — `wiki/memory/Memory Management for Beginners.md` is marked `status: proposed` and frames the capability split as a *recommendation*, while describing the **current combined `DeviceAllocator`** as the present reality. I verified against `crates/onnx-runtime-memory-governor/src/allocator.rs`: the trait indeed still carries `allocate`, `commit_allocation_range`/`decommit_allocation_range`, `allocation_committed_bytes`, `mapped_bytes`, `create_shared_prefix`, and the successful-no-op default it warns about. Because the doc calls the split a proposal and the combined trait the current state, it **survives** the #1192→#1247 revert rather than being invalidated by it. The proposed `onnx-runtime-memory-api` crate correctly does not exist. `with_memory(...)` — cited as the mechanism that already satisfies issue #513's core goal — exists (`onnx-genai-ort`, `ep-cpu`, `ep-cuda`). - **Execution backends** — the ORT-vs-native-nxrt description and EP contract are accurate and high-level; nothing contradicts the #1255 session-state unification. - **Repository Map** — every code path it points at exists (`decode_loop.rs`, `sampling.rs`, `speculative/`, all crate dirs). - **Links** — all 58 repo-relative links resolve. (The one apparent miss, `docs/path.md`, is inside a fenced note-*template* code block, not a real link.) - **Benchmark-number rule (`docs/benchmarks/2026-08-15-...` §32.2)** — no perf figure is copied without metadata; the only "speedup" mentions are the docs *teaching* that a bare speedup is insufficient, and the GB values are illustrative diagram quantities, not measurements. ## Verification - Nothing under a test, build script, or CI job references the new `wiki/` files (checked with `git grep`), so this cannot break the build or suite. - No in-tree markdown lint / mdbook config exists, so there is no docs build to run here — stated rather than invented. No inaccuracies found inside the diff; no corrections were needed. Closes #1190. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> --------- Co-authored-by: Copilot App <223556219+Copilot@users.noreply.github.com> Copilot-Session: c80f8522-983c-47f7-8241-2155a823aabe
…1240) # Zero-copy activation fast path, and an AVX2 Elu kernel Two changes in the same seam of `ActivationKernel::execute`: 1. **Generalise the contiguous-f32 fast path.** `silu_contiguous_f32` becomes `contiguous_f32(input, output, kernel)` — the kernel is now a parameter — so Celu, Mish and Elu write their result straight into the output buffer instead of taking the general path's two allocations and three passes (`to_dense_f32_widen` → `map` → `write_dense_f32_narrow`). 2. **An AVX2 Elu kernel** (`elu_ps` / `elu_f32_slice` / `elu_avx2`) built on the `exp_full_ps` that Celu and Mish already use, replacing a scalar `expf` loop. Base: `origin/main` with #1235 (Celu/Mish/Log) landed first. --- ## Reviewer measurements (Intel Core i7-13800H) All numbers below were taken by the reviewer on an **Intel Core i7-13800H (14C/20T, AVX2+FMA, no AVX-512), Windows 11**. No ORT reference build is available on this box, so the perf figure is an **ORT-independent, same-binary A/B**: the scalar Elu reference (`elu_scalar` loop) versus the AVX2 slice kernel (`elu_f32_slice`), 1M f32 elements, single-thread, min-of-200 iterations × 9 rounds, `black_box`ed on both arms of the same release binary. | Elu, 1M elems, 1 thread | scalar-native | AVX2 slice | speedup | |---|---|---|---| | min-of-200 × 9 rounds | 2.5842 ms | 0.5107 ms | **~5.06×** | This measures the vectorisation win itself (scalar vs AVX2) rather than a ratio against ORT, which is what the same-binary A/B can establish without an ORT build present. ### Peer measurement (original PR, AMD EPYC 9V74) The PR was originally measured on an **AMD EPYC 9V74, 16 physical cores, AVX2/FMA/F16C, no AVX-512**, as a ratio `ours / ORT` in a shared process. There, Elu went from **2.06× slower than ORT to ~4.2× faster** at 1M (4.96 ms → 0.57 ms against ORT's 2.39 ms), with Celu/Mish also improving from the fast path alone (their kernels are unchanged in this PR). These two results are **peers, not corrections**: different microarchitecture, different core count, and different reference (one is `ours/ORT`, the other is an internal scalar-vs-AVX2 A/B). Both show the AVX2 Elu kernel is a substantial win; neither refutes the other. --- ## The fast-path fallback is proven, not assumed `contiguous_f32` forms `&[f32]` and `&mut [f32]` over raw tensor pointers, which is only sound when both tensors are dense f32, agree on shape and strides, and do not overlap in memory. A fast path that silently produced wrong results on an unusual layout would be worse than a slow one, so the reviewer added a direct test — `contiguous_f32_fast_path_declines_every_unsafe_layout` — that pins the guard: - **dtype mismatch** (f64 input) → declines; - **non-dense strides** (a broadcast stride of 0) → declines; - **aliasing** (input and output over the same buffer) → declines; - **control**: two distinct dense f32 buffers of equal shape → the fast path runs *and* is the thing that produced the output. The closure passed on the decline arms panics if invoked, so the test fails if the fast path is ever wrongly taken. Verified by falsification: neutering the overlap check makes the aliasing arm run the kernel and the test goes **RED** (`fast path ran on a layout it must decline`); restoring it returns **GREEN**. The overlap check refuses even exact in-place aliasing, because forming `&[f32]` and `&mut [f32]` over the same bytes is UB in Rust regardless of what the elementwise machine code would compute. --- ## Numerics - `elu_special_values` — `-Inf → -alpha`, `-0 → -0` (bit-checked), `+0 → +0`, `+Inf`, `NaN`, `±MAX`, `±1e30`, and ordinary values against the scalar. - `elu_scalar_path_agrees_with_the_vector_path` — nine specials run once through a sub-`SIMD_MIN_LEN` slice (scalar) and once padded past it (vector), compared **bit for bit**, at alpha ∈ {0.25, 1, 2, 7.5}. - `elu_matches_scalar_across_alphas` — 20 003 points over `[-90, 90]` at four alphas with an alpha-scaled bound. `-0.0`: Elu is a select so the sign survives (`Elu(-0) = -0`), where Celu's trailing addition gives `Celu(-0) = +0`; both are pinned. The scalar reference moves from `alpha * x.exp_m1()` to `alpha * (x.exp() - 1.0)` so the scalar path, the vector path and ORT are the same function (there is no AVX2 `expm1`); the two spellings differ by at most ~6e-8 absolute, inside this module's 4e-7 bound. --- ## Verification (reviewer, Intel Core i7-13800H, Windows 11) - `cargo test -p onnx-runtime-ep-cpu --lib` — **1366 passed / 0 failed / 17 ignored** (+4 over #1235: three Elu tests plus the fallback falsifier). - `cargo clippy -p onnx-runtime-ep-cpu --all-targets -- -D warnings` — clean. - #1227's `#[target_feature(enable = "avx2,fma")]` confirmed intact on **both** `map_ps` and `map_bias_ps` after the merge. - MLAS remains non-default; `elu_f32_slice` uses `dispatch!` like its neighbours and touches no MLAS path. Co-authored-by: justinchuby <223556219+Copilot@users.noreply.github.com>
…ver (#1269) Wallace's different-architecture GPU probe result + coordinator standing-table update. - **qwen3.5-2b (hybrid linear-attn, head_dim=256)** native beats ORT in BOTH framings: eager 1.10x (within noise), capture 1.65x (clean, fallbacks=0). ORT CUDA EP can't place `LinearAttention`/`CausalConvWithState` on GPU (25 Memcpy nodes, graph blocked); native runs+captures the whole hybrid graph. - New lever (in flight, Deckard): **general head-size fused GQA decode kernel** — head_dim=256 falls out of the head<=128 fast path to `gqa_attention_reference_f32` = 31% of decode. Generalized (per user directive) over head sizes {64,80,96,112,128,192,256} + asymmetric dual-head-size audit (Gemma/DeepSeek MLA). - Support gaps: Phi-3.5-mini generic-cpu `If` op; qwen3.5-9b multimodal harness. Squad-state only (`.squad/decisions.md`, `identity/now.md`, `agents/wallace/history.md`). Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
Makes the native CUDA device sampler's kernels loadable on a machine
whose CUDA
toolkit is ahead of its driver. Independent of everything else: the only
file
touched is `crates/onnx-genai-ort/src/device_sampler.rs`.
## The defect
`CudaSampler` asked NVRTC for the **virtual** architecture
`compute_{cc}`. NVRTC
can only emit that as PTX, and a driver will only JIT PTX produced by a
toolkit
no newer than itself. On any machine whose CUDA toolkit is ahead of its
driver —
the normal shape of a box with a system toolkit and a vendor driver —
every
construction failed:
```
CUDA_ERROR_UNSUPPORTED_PTX_VERSION: the provided PTX was compiled with an unsupported toolchain
```
**Fix.** Prefer real-architecture SASS via `nvrtcGetCUBIN` on
`--gpu-architecture=sm_{cc}`; nothing is JIT-compiled, so a newer NVRTC
is
harmless. Fall back to PTX only after checking `nvrtc_version <=
driver_toolkit_version`. If neither is possible, fail with an error
naming the
device, the NVRTC and the driver, and the three ways to resolve it —
never a
silent fallback to a host path.
Side effect (reported by the author, not re-measured here): sampler
construction
dropped from 24.9 s to 4.2 s because SASS needs no JIT.
## Verification on this reviewer's hardware (RTX 4060, driver 591.55)
This box happens to be exactly the failing shape the PR targets, so the
defect
reproduces here rather than being taken on faith. The device smoke test
prints:
```
device sm_89 | NVRTC 13.3 | driver supports CUDA 13.1
NVRTC is newer than the driver: only the SASS path can succeed here
```
NVRTC 13.3 > driver's CUDA 13.1, so `ptx_jit_supported` is false and
only the new
SASS path can load.
- `cargo test -p onnx-genai-ort --features cuda device_sampler --
--include-ignored`
→ **18 passed, 0 failed, 0 ignored**, including the real-device
`device_smoke::device_argmax_runs_in_the_default_environment` (no
`LD_LIBRARY_PATH`/PATH reordering to hide the skew).
- Full crate lib suite `--features cuda --lib -- --include-ignored` →
**166
passed, 0 failed, 0 ignored**.
- `cargo fmt -p onnx-genai-ort --check` clean; `cargo clippy -p
onnx-genai-ort
--features cuda --all-targets -- -D warnings` clean.
### Falsification (revert → RED → restore → GREEN)
To prove the fix is what makes it load — not incidental — I forced the
pre-fix
PTX path (`cubin_supported` → false, `ptx_jit_supported` → true) and
re-ran the
smoke test. It went **RED** with the exact original failure on this
hardware:
```
DriverError(CUDA_ERROR_UNSUPPORTED_PTX_VERSION,
"the provided PTX was compiled with an unsupported toolchain.")
```
Restoring the two functions returned the smoke test to **GREEN**. So on
this box
the SASS path is load-bearing, not cosmetic — the regression guard
actually
discriminates here.
## Tests added
- `device_sampler::device_code_selection` — pure logic for `arch_flag`
(real vs
virtual, two-digit minors), `ptx_jit_supported` (incl. the 13.3-vs-13.0
case),
`cubin_supported`, and `decode_driver_version`.
-
`device_sampler::device_smoke::device_argmax_runs_in_the_default_environment`
— real device, default environment, 202048-wide microcase.
## Performance attribution
This PR moves **no** kernel benchmark number: it changes which device
image is
compiled and loaded, not what the kernels do. The kernel's own
performance is a
separate, stacked concern (#976, not landed here). Its value is that the
native
sampler works at all on a skewed toolkit/driver, in the default
environment,
without `LD_LIBRARY_PATH` help — verified above on real sm_89 hardware.
---
Extracted from the Muse stock-ORT performance work so it can land
independently
of the inference-metadata redesign (#828).
Signed-off-by: Justin Chuby <justinchuby@users.noreply.github.com>
Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
## Symptom
Driving this model from OpenCode, a tool call fails on its arguments
rather than its name:
> ✗ bash failed — The bash tool was called with invalid arguments:
`SchemaError(Missing key at ["command"])`
Captured off the wire, the call this server produced was:
```json
{"name": "bash", "arguments": "{\"bash\":\"\\n<atem:parameter name=\\\"bash\\\">command\"}"}
```
The argument the tool needs is not merely wrong — it is absent, and a
raw `<atem:parameter>` tag is sitting in the value.
## Cause
The model opened the same parameter twice. The parser read the value
from the first open tag to the first close tag, so the stray second open
tag was swallowed into the value and carried the real argument in with
it.
## Fix
A parameter cannot contain a parameter, so an open tag appearing before
the close means the *outer* one is the stray. The inner tag is now taken
as the parameter and the outer discarded, which recovers the argument
the model meant to pass.
This matches the leniency the format already asks for — it states
outright that its output "is not expected to be valid XML and is parsed
with regular expressions" — and it is confined to the parameter that was
opened twice.
## Verification
- 223 lib + 31 http + 11 bundle tests pass; clippy `--all-targets`
clean.
- Three new tests: the parameter opened twice keeps its inner value; a
stray open tag leaves its neighbours untouched; a well-formed call is
read as written, including a value that legitimately carries angle
brackets.
Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
Copilot-Session: 0190e2eb-abe4-451f-b36d-44a035a99b7e
Replacement for #1210, rebased onto current `main` after the wiki base (#1190) landed as #1267. The original is stacked on #1190's head branch and cannot be squash-merged directly (GitHub stacked-PR 403). Content is identical to #1210's own diff: 154 files under `site/`, one Pages workflow, `.gitignore` rules, and two small edits to `wiki/`. ## What this adds The Quartz 5 static-site machinery that publishes the `wiki/` content (added in #1190/#1267) to GitHub Pages: - `site/quartz/` — the Quartz engine, `package.json` + committed `package-lock.json`, `quartz.config.yaml`, plugin lock (`quartz.lock.json`). - `site/scripts/` — Python validators (`validate_wikilinks.py`, `validate_site.py`) and a test suite. - `.github/workflows/wiki-pages.yml` — build + deploy workflow. ## Practical review (the questions that matter for publishing machinery) - **Build dependency / CI job — acceptable and documented?** Yes. The new `wiki-pages.yml` pins every action to a full SHA, runs with top-level `permissions: contents: read` (the deploy job separately requests `pages: write` / `id-token: write`), sets `persist-credentials: false`, installs via `npm ci` from the committed lockfile (reproducible), and **deploys only on push to `main`** (build-only on PRs/branches). It is self-contained and does not touch the Rust build or any existing workflow. - **Commits generated output that will rot?** No. `.gitignore` excludes `site/quartz/node_modules/`, `public/` (the built site), `.quartz/`, `.quartz-cache/`, vendored static assets, and `tsconfig.tsbuildinfo`. No `node_modules`, built HTML, or `public/` is committed. The committed `package-lock.json` pins exact dependency versions. - **Conflicts with the existing docs layout?** No. There is no pre-existing `site/` or Pages/Quartz/mkdocs workflow on `main`; the existing workflows (`ci`, `publish`, `wheels`, `benchmark`, `squad-*`, …) are untouched. The site lives under `site/` and *reads* `wiki/` (`quartz build -d ../../wiki`); it does not move or rewrite `docs/`. - **Config coherence.** `quartz.config.yaml` `baseUrl: justinchuby.github.io/onnx-genai` matches the `/onnx-genai/` base path the post-build validators (`validate_site.py public /onnx-genai/`, `rewrite_repository_links.py`) expect. ## Verification I could run - **Wikilink validation passed** (`validate_wikilinks.py ../../wiki`): *Validated 96 wikilink(s) across 20 note(s).* - Nothing in cargo or the Rust test suite references `site/`, so this cannot affect the crate build or tests. ## Gap I must be honest about — the end-to-end site build I could **not** run the actual Quartz build (`npm run wiki:build`) in this environment. `npm ci` fails because the pinned transitive dependency `@xmldom/xmldom@0.8.14` (pulled via `pixi.js`) is **not available** on the corporate npm proxy (which only serves `0.9.x` for that package), and the public registry is TLS-blocked here. There is no lockfile-preserving workaround from this sandbox. Consequently the Python test suite shows 37/78 failures — but **every** failure is a missing-node-dependency error (`Cannot find package 'acorn' …`) from tests that shell out to the un-installed Node scripts, not a logic defect; the network-free checks pass. The definitive build proof therefore rests on the `wiki-pages.yml` job this PR adds, which runs on `ubuntu-latest` against public npm where `0.8.14` exists. Blast radius is contained: a build failure would surface only in this isolated workflow / Pages deploy and cannot break the Rust build or test suite. Merging on the machinery review + link validation, with the JS build unverified-locally and delegated to CI, clearly flagged. Closes #1210. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> Co-authored-by: justinchuby <223556219+Copilot@users.noreply.github.com>
… decode) (#1323) ## Summary Two **byte-identical** levers on the fused SwiGLU gate/up expert GEMV (`qmoe_gate_up_activate_f32`) — the #1 decode kernel on DeepSeek-V2-Lite int4 QMoE (19.8% of GPU decode after the #1317 TopK win). 1. **Occupancy `_occ` entry → DEFAULT-ON.** The `__launch_bounds__(256, 6)` sibling was already shipped default-OFF behind `ONNX_GENAI_QMOE_OCC`. `__launch_bounds__` only caps register allocation, so the accumulate order and both `block_sum` reductions are unchanged → bit-identical. Opt out with `ONNX_GENAI_QMOE_OCC=0`. 2. **Read-only-cache (`__ldg`) int4 loads, scoped to gate/up only.** A new `ReadOnly` template param on `qmoe_int4_chunk` routes the packed weights / scales / zero-points through the read-only data cache. Same bytes, same decode → bit-identical, but it cuts the Long-Scoreboard weight-load latency the kernel is bound on. The fc2 `qmoe_linear` path **regressed** under `__ldg`, so it keeps the default cached load (`ReadOnly` defaults to `false`). ## Kernel A/B (nsys, H200, eager, GPU pinned) | Kernel | Baseline | This PR | Δ | |---|---|---|---| | `qmoe_gate_up_activate_f32` | 42.68 µs | **34.24 µs** | **−19.8%** | | `qmoe_linear_f32` | 28.79 µs | 28.77 µs | ~0 (no collateral regression) | | `qmoe_route` | 6.46 µs | 6.46 µs | 0 | The selective `ReadOnly` scoping is what keeps `qmoe_linear` flat — a blanket `__ldg` on the shared chunk helper sped up gate/up but regressed linear +14.6%. ## E2E A/B (capture, 128 tok, medians of 5, H200 `CUDA_VISIBLE_DEVICES` pinned) | Arm | Runs (tok/s) | Median | |---|---|---| | Baseline (`ONNX_GENAI_QMOE_OCC=0`, no ldg) | 124.58 / 124.65 / 124.98 / 125.01 / 125.73 | **124.98** | | This PR (occ default-on + gate/up ldg) | 127.52 / 127.55 / 127.77 / 128.13 / 129.22 | **127.77** | **+2.23% V2-Lite decode**, **byte-identical** token stream (128-tok md5 match between arms). ## Numerics gate - `qmoe_gpu` `gpu-tests`: **32/32 pass**, incl. `qmoe_occ_is_bit_identical` (0 differing bits across fp16/bf16/fp32 × rows∈{1,4,6,8}). - V2-Lite native-CUDA **golden decode lock** (`deepseek_v2_lite_native_cuda_matches_golden_greedy_sequence`): **byte-identical** greedy sequence — the strongest proof both levers preserve bits end-to-end. ## Scope / safety - Single file: `crates/onnx-runtime-ep-cuda/src/kernels/qmoe.rs`. QMoE-only kernel; dense models unaffected. - Both levers are byte-identical (no fp32 partial reorder), so no oracle-tolerance argument needed — exact token match. - `ONNX_GENAI_QMOE_OCC=0` remains as an opt-out for A/B / bisection. GO. Please **admin-merge** after review — not self-merging. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
… attribution fix (#1325) ## Issue #1295 — VMM admission cap (cliff → plateau) + `vram_free_ms` attribution fix Follow-up to #1315 (the Phase-1 characterisation). Owner-greenlit Phase 2. ### What this does **1. Admission cap sourced from *usable* VRAM.** The device-authority ledger ceiling already bounds the *sum* of all Device-tier mapped physical (weights + KV + activations). It was mis-**sourced**: `Fraction(0.90)` resolved against **nominal** total (8188 MiB → 7369 MiB), which *exceeds* measured usable free (~7100–7959 MiB on this shared box). A ceiling above usable lets the ledger permit leases the driver can't satisfy → the WDDM driver OOMs first. This clamps the ceiling to `cuMemGetInfo` **free × fraction** (default 0.90), so the ceiling is ≤ usable and the **ledger refuses cleanly before the driver falls off the fault-in cliff**. No parallel mechanism — the refusal routes through the existing `try_commit_span` / `reserve_mapped_growth` G3 path. **This is a cliff→plateau conversion, NOT a raised `N_max`.** The ceiling does not move; the *failure mode* changes from a catastrophic driver crash (and 90 s `vram_free` storms) into a graceful governed refusal. The next reader should not expect more sequences to fit — they should expect over-subscription to be *refused* instead of *thrashed*. **Derivable, hardware-recorded fraction (not magic — cf #1261).** `DEFAULT_VMM_MAPPED_FRACTION = 0.90`, overridable via `ONNX_GENAI_VMM_MAPPED_FRACTION`. The 0.90 protects the ~740 MB margin between usable-free and the *fault-in* cliff observed in #1315: at `fraction = 1.0` the governed refusal lands exactly at free with no margin and the run sits on the edge of the fault-in degradation. What would justify changing it: a box where the fault-in cliff sits at a different offset below usable-free (measure it, don't guess). **2. `vram_free_ms` instrument fix (first-class deliverable).** `weight_paging.rs::release_allocation(true)` synchronised **both** streams *inside* the `GLOBAL_VRAM_FREE_NS` bracket — so "free" time could be waiting on paging-slowed work. That headline (90 s) drove the whole investigation. Split the stream-drain into `GLOBAL_VRAM_FREE_SYNC_NS` (surfaced as `vram_free_sync_ms`); the free bracket now spans only unmap/release/free. **Corrected result (2 release samples, qwen14b-zp, RTX 4060 8GB, single-thread):** `vram_free_ms` = 41.5–45.2 s, `vram_free_sync_ms` = 1.6–1.9 s → sync is a stable **~4 %** of the bracket. **The headline is not substantially retracted** — the corrected free time is still tens of seconds of *genuine* unmap/release churn. An honest instrument was still worth building; it refines the number rather than overturning it. ### Evidence - **Live cliff→plateau:** no-limit 14B VMM load — *before* (nominal-sourced / no headroom): `cuMemSetAccess failed: CUDA_ERROR_OUT_OF_MEMORY` (driver crash). *After* (cap 0.90): clean `MemoryError` at ceiling `6699456921 B` = `floor(free × 0.90)`. At `fraction = 1.0`: governed refusal at `7443841024 B` (= free exactly), proving the sourcing-from-free moves refusal driver→ledger. - 13 governor unit tests pass (0 failed, 0 ignored), incl. a **non-vacuous** G3 refusal test. ### Honest limitation A live multi-request **N=8 throughput plateau could not be captured on this box** — the governed batch-N sweep OOMs at *load* for N≥4 (14B weights ≈ usable VRAM on 8GB). The plateau *mechanism* is proven by the unit test + the live load-time refusal; the specific 0.80→peak tok/s curve is **not measurable here and not claimed**. ### Cross-slice - **MoE #1308:** the cap is a *single shared authority* — MoE per-expert paging reuses the same device-authority ceiling, no duplication. (And #1308 *lowers* peak mapped physical, so it does not worsen this cliff.) - **OS-paging / residency-hint direction:** this is evidence *for* it — if the binding cost is WDDM fault-in/eviction beyond usable VRAM and our own map/unmap is not the driver, our machinery is largely fighting for control it doesn't benefit from. The admission cap is correct under either architecture: refusing to over-subscribe is right whoever does the paging. ### Test / build `--features "cuda,native-backend"`, single-thread. Governor suite: 13 passed / 0 failed / 0 ignored. Pre-existing failures unchanged and unrelated: #1284 (three harness), #1305 (four kernel on this GPU). Closes #1295. Co-authored-by: justinchuby <223556219+Copilot@users.noreply.github.com>
… tok/s (#1326) ## What & why Composes the two prior trace-driven studies (#1321 single-tier residency, #1322 tiering) into **the number a deployer needs**: how much VRAM a large MoE needs to hit a target tok/s. Assumes a **DRAM tier holds the whole bank** (so SSD is out of the critical path — #1322 result), leaving **PCIe DRAM→VRAM as the only per-step expert-transfer cost**, and applies the best practical policy (**static hot-pin**) from #1321. Routing is **measured** on the granite trace (32 experts, top-8, 24 layers, 192 decode steps, real trained router; **CPU-only, no GPU**). Every **bandwidth constant is inferred/assumed**, not measured on this box. Box: **i7-13800H (14C/20T), RTX 4060 8 GB, WDDM**. ## Headline (measured routing, inferred bandwidth; 16 MiB experts) **VRAM 25%, batch 1, PCIe 4.0 x8 (~13 GB/s): 1.63 GB/token → ≈ 8 tok/s from expert transfer ALONE, before any compute.** A bank that overflows VRAM is PCIe-bound in the single-digit-to-low-double-digit tok/s range. **DirectStorage/faster-SSD cannot help — the SSD is already out of the path; the wall is PCIe.** ## The curve & the two levers - Static hot-pin sits **~17–27% above the Belady oracle** across the range (consistent with #1321). - **No sharp knee.** Concentration is milder than the assumed 80/20 Zipf, so VRAM must cover a **large** share of the bank (~50–65%) to make transfer cheap — *later* than a Zipf assumption predicts. Even at 50% VRAM you still stream ~49 experts/token. - **Batching** W=8 roughly halves GB/token (25%: 1.63 → 0.81). ## Rule of thumb — min VRAM % for target tok/s (transfer floor only, inferred BW) | batch | link | 5 | 10 | 20 | 40 tok/s | |---|---|---|---|---|---| | W=1 | x8 (13) | 5% | 35% | 65% | 75% | | W=1 | x16 (26) | 5% | 5% | 35% | 65% | | W=8 | x16 (26) | 5% | 5% | 5% | 50% | | W=8 | PCIe5 (55) | 5% | 5% | 5% | 5% | Batch size and PCIe generation matter as much as VRAM. ## Caveats (recorded in the doc) - **The W=8 byte savings are NOT free in wall-clock.** A sibling measurement finds **M=1→2 costs ~5.4×/step for 2× work** (batch-decode penalty, GEMV excluded) — cross-referenced. The batching win holds in **bytes**, not yet in tokens/second. - Trace-driven sim gives no wall-clock, no paging-mechanism cost, and no proof skew generalises to a 256-expert router (expected *more* concentrated → earlier knee; granite is the conservative flat end). - Mechanism choice (VMM vs OS page cache vs DirectStorage) is a separate slice; this strengthens the cheapest-first (OS page cache) instinct there. ## Files - `scripts/moe_pcie_floor.py` + `scripts/moe_pcie_floor_results.json` - `docs/benchmarks/2026-08-18-moe-offload-pcie-floor.md` (links #1321/#1322 docs) Reproduce: `python scripts/moe_pcie_floor.py`. Analysis/measurement only — ours to merge. No runtime policy, no #82/#87 wiring. Co-authored-by: justinchuby <223556219+Copilot@users.noreply.github.com>
#1327) Docs-only follow-up to #1324. Records, from the already-measured M>=2 capture-segmentation breakdown (no new measurement), the **fixed per-step batch cost** the owner asked for as the input to the expert-aware MoE batching question. ## What this answers Route-aware / expert-aware batching only pays if being a batch is cheap per step. This PR states, as a standalone number, what batch-N costs before any clever scheduling: | arm | fixed per-step batch cost | marginal per row | |-----|--------------------------:|-----------------:| | **before fix** | **~22-24 ms** (flat, ~independent of M) | ~0.4 ms/row | | **after fix** | **~1.4 ms** | ~1.9 ms/row | *(qwen05b-q4, resident, RTX 4060 Laptop 8 GB, i7-13800H, CUDA 13.1, single-threaded decode; fit `total_ms ~= fixed + slope*M` over M=2/4/8, agrees with the `kernel_host_dispatch` mechanistic term.)* **Verdict:** the penalty is **structural (capture segmentation), not diffuse**, and therefore fixable — the byte-identical fix (`squad/batch-swiglu-capture-fix`) removes it, turning the step into a normal `~1.4 ms + ~1.9 ms/row` linear model. So expert-aware batching is viable on this stack **conditional on the fix landing**; unfixed, the ~22 ms flat penalty would likely swamp any expert-load bandwidth saving. **Backend neutrality (DRY):** measured on native decode; the seam and fix live in the shared CUDA EP kernel + capture machinery, so the ORT path is inferred (not separately measured) to inherit both the penalty and its removal. Recommendation to the MoE routing-trace simulation: model per-step batch cost as **~22 ms fixed (today) vs ~1.4 ms fixed (fix landed)** and treat the pay/no-pay verdict as conditional — it may invert between the two. Docs + decisions-inbox only; no code change. Ours (batching slice). Co-authored-by: justinchuby <223556219+Copilot@users.noreply.github.com>
Campaign-brain: #1323 (QMoE gate_up occ + selective __ldg, +2.2% V2-Lite decode, byte-identical) and the ORT upstream-contribution recon (NO-GO — both our TopK/argmax wins already in ORT main, verified). Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
… VRAM (#1329) ## What & why Extends the composed PCIe-floor model (#1326) with an **expert-size sweep**. PCIe bytes/token is the wall, and it scales **linearly with expert size**, so **quantisation attacks it most directly of all levers** — no policy, prediction, or scheduler. Priced on the **same axis** as the VRAM (static-hot-pin) lever, using the same measured granite routing curve. Routing **measured** (CPU, no GPU); expert sizes and bandwidths **inferred**. ## Same-axis result (bytes/token, W=1, static hot-pin) | expert | @25% VRAM | equivalent VRAM lever | |---|---|---| | 16 MiB | 1670 MiB/tok | baseline | | 8 MiB (2×) | 835 | ≈ static-pin @ **50%** VRAM | | 4 MiB (4×) | 418 | ≈ static-pin @ **65%** VRAM | **A 4× expert-size cut at 25% VRAM equals raising VRAM to 65%** — a 40-point VRAM saving unattainable for an over-subscribed bank — and the two **compose** (4 MiB @ 50% = 197 MiB/tok). Confirms the prediction *measured against our own curve*: quantisation beats pinning in the 25–50% VRAM range. **Revised min-VRAM for 20 tok/s (x8):** 16 MiB → 65%, **4 MiB → 5%.** ## Honesty constraints (these decide if the factor is real) 1. **Quantisation trades accuracy** — a product decision, not a free win (unlike lossless pinning/batching). Labelled as such on every row. 2. **Baseline precision changes the answer completely.** 4× holds from **f16** (→int4). If the 16 MiB target is **already int4** (plausible for DeepSeek-class), the 8/4 MiB rows mean int2/int1 — severe, and **no further quant lever on the offload path**. Nail the baseline before quoting a factor. 3. **Mechanism constraint (from the mechanism agent):** only **canonical MatMulNBits int4** is file-backable; Marlin/MLAS-prepacked are excluded from any storage-to-device/mmap path. So the offload quant lever means "get to canonical int4" — which coincides with the accuracy sweet spot. ## Files - `scripts/moe_quant_lever.py` + `scripts/moe_quant_lever_results.json` - `docs/benchmarks/2026-08-18-moe-quant-lever.md` (+ forward-link from the PCIe-floor doc) Reproduce: `python scripts/moe_quant_lever.py`. Analysis/measurement only — ours to merge. No runtime policy, no #82/#87 wiring. Co-authored-by: justinchuby <223556219+Copilot@users.noreply.github.com>
Scribe consolidation of #1317 (router TopK) + #1323 (gate_up_activate) decision notes into decisions.md; compact deckard history. (Session/orchestration logs are gitignored, local-only.) Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
…sibility gate (#1331) ## MoE offload — the prediction-window measurement The feasibility gate for **dynamic** expert prefetch vs the established static hot-pin: is the window between "router output for layer L known" and "layer L expert weights needed" longer than a page-in? If not, predictive prefetch cannot hide the latency regardless of prediction quality. **Model/hardware:** granite-3.0-1b-a400m-instruct (32 experts, top-8, 24 layers), onnxruntime 1.27.0 CPU EP, batch 1 — i7-13800H (14C/20T), RTX 4060 8 GB, WDDM. ### Measured - **Intra-layer window ≈ 0** (topology, EP-independent): each layer's `TopK` router indices feed 32 `Equal` masks whose immediate successors are the expert MatMuls — routing is the direct predecessor of expert compute, no intervening work. Demand-prefetch of a layer's own experts hides nothing on any hardware. - **CPU decode:** 188 ms/token median (20 steps), 7.83 ms/layer. ### Inferred / cross-referenced - Page-ins (mechanism agent, measured this box): PCIe H2D pinned 1.5 ms, NVMe cold 3.7 ms per 16 MiB expert. - The 7.83 ms/layer CPU-dense-f16 window **overstates** the GPU-sparse target (the whole token collapses to a few ms → per-layer window below a single page-in), and using it requires predicting next-layer routing, which is data-dependent + weakly cross-correlated (second-order, measured earlier). ### Verdict **Dynamic demand-prefetch cannot beat static hot-pin.** The only experts prefetchable with certainty *and* lead time are the always-on core (zero-lookahead-predictable) — which static pin already keeps resident. Own-layer experts have a ~0 window; future-layer experts need prediction the routing does not support. Cold-NVMe (3.7 ms) exceeds any realistic decode window → cold-tier fetch is not viable in the decode critical path (residency, not prefetch). Dynamic prefetch is not a lever here; the window forecloses it. Measured vs inferred kept separate. No runtime policy, no #82/#87 wiring. This closes the current investigation line. Co-authored-by: justinchuby <223556219+Copilot@users.noreply.github.com>
…ts (not commit-and-subdivide) (#1332) ## #1295 follow-up — MoE VMM sub-granule packing: the trade curve (analysis, no impl) **Reporting before implementing, as the owner asked.** This answers "can VMM commit more at once and subdivide ourselves?" from the already-measured granite router-skew distribution — no packing code was written. ### The question and the trap The 2 MiB CUDA granule wastes **2.7×** on granite int4 experts (~0.75 MiB). The proposed escape — commit a large region (tens of MiB), pack ~85 experts, subdivide ourselves — is favourable on granule *waste* alone. But committing is also the unit of *residency*, so the real question is whether sub-division keeps selectivity. ### Finding (measured, granite-3.0-1b-a400m, 32e/top-8/24L, real IBM router, CPU EP, 192 decode steps; expert 0.75 MiB, granule 2 MiB) **The granularity that minimises resident mapped physical — and the owner's own `waste × transferred` product — is ONE 2 MiB granule holding TWO experts (E=2), not a large region.** | E | waste | resident MiB/layer/step | transfer MiB | sel% | |---|---|---|---|---| | 1 (per-expert) | 2.67× | 16.00 | 6.00 | 100% | | **2 ← optimum** | 1.33× | **13.34** | 10.01 | 167% | | 8 | 1.00× | 19.07 | 19.07 | 318% | | 32 (whole bank) | 1.00× | 24.00 | 24.00 | 400% | Resident physical is **U-shaped with its minimum at E=2 and rises back to whole-bank levels for every larger chunk** — the "tens of MiB" proposal is at the *worst* end. E=2 beats per-expert (16.00→13.34, **−17%**) because the second expert in a granule you already pay for is **free physical**. ### Why large regions are fatal (inferred) Every decode step activates **all 24 layers**, so any chunk spanning layers is hot in 100% of steps. A 64 MiB region (~85 experts) can't fit one layer's 24 MiB bank, so it **must** span layers → always resident → whole-bank pinning. The "near-zero waste" arithmetic is real but irrelevant; the binding objective is resident physical. ### The other answers - **Hotness-sorted packing rescues selectivity, modestly** (~5–14% vs arbitrary at every E; does *not* rescue large chunks). - **Static packing is stable**: decays only **3.4–7.1%** across prose/code/math vs a per-prompt oracle — so packing once is justified. Layer-1/2 always-on experts stay free pins. - **Interaction with #1295/#1325**: span count is **not** the cliff driver, so E=2 does **not** help the cliff — but lower peak mapped physical raises the admission-cap **plateau height** (not `N_max`). Single shared authority preserved. This is not conflated in the doc. ### Decisive framing (which architecture this assumes) - **If we keep our own VMM paging:** the minimal correct change is **E=2** (two experts/granule, hotness-packed, within-layer) — provably the trade-curve optimum, and small (still commits exactly one granule). **Do not build large-region commit-and-subdivide** — dominated on every metric. - **If OS hinting works, sub-division is unnecessary:** 4 KiB OS paging dissolves the granule problem (zero waste + full per-expert selectivity, strictly better than any VMM packing). - **My #1295 finding leans toward OS paging:** the binding cost is WDDM fault-in beyond usable VRAM, not our bookkeeping — so we may be fighting for control we don't benefit from. The condition to **not** build even E=2 is the OS-hint agent confirming WDDM honours residency hints with an adequate prediction window. I did **not** measure hinted-vs-unhinted fault-in (their job) but recorded that per-page fault-in is exactly what a prefetch hint would target. ### Reproduction `scripts/moe_granularity_analysis.py` reuses the exact `moe_router_skew.py` probe, captures per-step selection sets (the committed JSON has only aggregates), and prints the curve for both packings. Deterministic, bit-for-bit. Box: i7-13800H, RTX 4060 Laptop 8GB, driver 591.55, CUDA 13.1, WDDM. Analysis is CPU-EP (the *which experts* question is dtype-/EP-independent); granule is the CUDA device granule. No GPU decode claimed. Doc-only + reproducible script; no source change, nothing to test beyond the script's deterministic output. Follow-up to #1295 (closed by #1325). Co-authored-by: justinchuby <223556219+Copilot@users.noreply.github.com>
…e page cache, PCIe wall, what to build (#1333) ## MoE expert-offload architecture — synthesis of three converged workstreams Composes the mechanism, policy, and batching findings into one authoritative document, `docs/memory/MOE_OFFLOAD_ARCHITECTURE.md`, in the honesty register of `MEMORY_MANAGEMENT_MODEL_DESIGN.md` (measured vs inferred separated; hardware/model/regime on every figure; per-agent attribution so any claim traces to its evidence). ### The architectural conclusion of the week **Three regimes** decide whether a routing-aware cache is worth building (measured on this box — RTX 4060/WDDM, 63.8 GB DDR5-6400, gemma-3-27b, routing-skewed pattern under swept memory pressure): | regime | condition | OS captures | headroom | |---|---|---|---| | 1 | cache ≥ full bank | ~100% | ≈0 | | 2 | hot-set < cache < bank | ~88% (LRU exploits skew for free) | small | | 3 | cache ≲ hot-set | **33% hot-warm** | **~3× pin win** | - **Host/PCIe crossover:** warm host reads (~10.5 GB/s) exceed every host→device path (zero-copy 5.9, staged 10.9 GB/s), so in regimes 1–2 the host tier is not the bottleneck — **PCIe is**. Only regime-3 eviction-miss traffic drags host throughput below PCIe. - **CUDA UM is empirically closed on WDDM** (~0.6 GB/s, `concurrent=0`, `cuMemAdvise_v2`/`cuMemPrefetchAsync_v2` return `INVALID_DEVICE`) — recorded as a hard negative so nobody re-proposes it. - **Layout constraint gates the whole storage family:** only canonical MatMulNBits int4 / `BlockQuantizedMoE` are file-backed-usable; Marlin and MLAS-prepacked are excluded. - **DirectStorage/GDS** are PCIe-bound Regime-B backstops (bank > host RAM), not breakthroughs; the D3D12→CUDA interop link is flagged inferred-not-measured. ### Composed with the other workstreams (cited, not re-derived) - **Policy** (MoE agent, #1321/#1322/#1326/#1331): concentration milder than 80/20 (25% → 45.6%); static hot-pin within 17–27% of Belady; PCIe floor **1.63 GB/token → ~8 tok/s** (25% VRAM, batch 1, x8); no knee; **dynamic prefetch foreclosed** by a ≈0 intra-layer window. - **Batching**: **−61% bytes/token** (W=2→8) is the largest lever, but **blocked** behind a ~5.4× M=1→2 step overhead. ### What to build / not build 1. Routing-aware **static** hot-pin — **regime 3 only** (near-worthless in 1–2; dynamic prefetch foreclosed). 2. Target (hundreds-of-GB MoE) sits in regime 3 → worth building, but **cannot be validated on this 8 GB laptop**; the doc states what would (Linux datacentre GPU + real MoE). 3. **Quantisation** = best return per complexity (linear PCIe relief, no policy), at a stated accuracy cost. 4. DirectStorage/GDS = Regime-B backstop only. 5. Never re-propose CUDA UM on WDDM. Cross-links the result back into `MEMORY_MANAGEMENT_MODEL_DESIGN.md` (which said MoE "should be measured before any residency policy is designed"); its dense conclusions (`reads_per_step=1.000` #944, ~30× OS win #864/#874) kept intact and reframed as regime 1 with no reuse. Docs-only; nothing under `docs/` is consumed by a build script (verified). Measured on RTX 4060 / driver 591.55 / CUDA 13.1 / WDDM. Closes #1266. Co-authored-by: justinchuby <223556219+Copilot@users.noreply.github.com>
…s.md (#1335) ## Summary Consolidates two evidence-backed NO-GO findings from Deckard into the canonical decisions log. Both findings close QMoE decode optimization levers on DeepSeek-V2-Lite int4 (H200). ### Decision 1: Port ORT int4 decode-GEMV ideas — NO-GO - Byte-identical load-batching regressed gate_up 34.18→40.96µs (+19.8%) via occupancy spill (40 reg/thread cap at 6 blk/SM from #1323) - Layout interleave inapplicable (loads already fully coalesced, 128B/warp) - qmoe_linear already at 8 blk/SM hardware ceiling (85.4% occupancy, zero headroom) - Both kernels: memory-latency bound (~11% DRAM), occupancy-max ### Decision 2: Fuse gate_up_activate→qmoe_linear pipeline — NO-GO - Intermediate round-trip = 20.99 KB = **0.075% of DRAM traffic** (decisive measurement) - Intermediate is served from L2, not DRAM — no real traffic saving - True fusion collapses grid to ~6 blocks/route, ~4.5% SM occupancy vs occupancy-max - Dominant cost: irreducible int4 expert-weight reads (27.85 MB combined) ### Strategic conclusion Both QMoE decode levers exhausted with quantitative ncu/nsys evidence. Kernels at M=1 decode roofline (occupancy-max, latency-bound, weight-traffic irreducible). Further QMoE gains require lower-bit weight format — model/format change, not kernel work. **Next lever: `attention_row` (12.4%, 24.7µs).** ## Files changed - `.squad/decisions.md` — two NO-GO entries appended - Inbox notes deleted (were untracked; not in this PR) ## Archive check All entries in decisions.md are dated 2026-08-18 — nothing >30 days old. No archive needed. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
…don't build it) (#1336) ## #1295 follow-up — re-scope MoE sub-division: the honest "no" The owner narrowed the question after the granularity trade-curve analysis was written: **small experts fit and never page, so granule waste is irrelevant; the real problem is over-subscription with large (≥ granule) experts** (GLM/DeepSeek-class, ~16 MiB = 8× the 2 MiB granule). This PR re-scopes the merged trade-curve doc accordingly and records the conclusion. ### Surviving question and answer *For large ≥-granule experts in the over-subscribed regime, commit large chunks and subdivide, or per-expert commit?* → **Each expert its own commit. Don't build sub-division.** The two arguments that could have survived both fail: - **Call count / churn.** A 16 MiB expert already maps **8 granules**; packing K experts into one chunk cuts the number of `cuMemMap`/`cuMemUnmap` *calls* (1 vs K) but **not granules mapped** (8K either way). The binding cost is **WDDM fault-in** (per-granule), and #1295/#1315 proved **span/commit count is not the cliff driver**. The reducible per-step API overhead is **3.39 ms/step** (16 MiB skewed) vs H2D's **10.9 ms/step** — well under ¼ of per-step paging cost, and **zero** of the fault-in cliff. - **Coherent residency grouping.** At ≥ granule there is no granule-sharing to force, so grouping is a **policy** decision served by per-expert commits + eviction. Packing into a shared commit **removes** independent per-expert eviction (a chunk is resident if *any* member is hot) — strictly worse, for no benefit. Grouping belongs in the offline residency-policy replay, at per-expert granularity. The sub-granule **E=2** result is retained only as arithmetic for the case now set aside. What survives: the **admission cap** (PR #1325, correct under our-VMM *and* OS-paging) and a **per-expert residency policy** (separate agent). If OS 4 KiB hinting is honoured, packing is moot entirely. Doc-only. Box: RTX 4060 Laptop 8GB, driver 591.55, CUDA 13.1, WDDM; model granite-3.0-1b-a400m / GLM-class 16 MiB experts inferred. Follow-up to #1295 (closed by #1325). Co-authored-by: justinchuby <223556219+Copilot@users.noreply.github.com>
…y (+5.5% V2-Lite long-ctx, byte-identical) (#1337) ## Summary `attention_row` (the default-domain standard-MHA kernel) launches **one block per (batch, q_head, query) row**. In the **decode** regime (`q_seq == 1`) the grid collapses to `batch * q_heads` blocks — just **16** for a 16-head model like DeepSeek-V2-Lite — leaving the kernel badly under-occupied and memory-latency bound. This PR widens the **decode** launch to **256 threads/block** so each of the few per-head blocks has enough warps to hide the per-key K/V load latency. It is **byte-identical** (more threads, identical per-thread reduction order) and touches only the decode path; prefill keeps 128 threads. `ONNX_GENAI_ATTN_ROW_THREADS` overrides the width for tuning. ## Why (ncu profile verdict, H200 pinned) `attention_row` decode at ~65-key context: **Grid=16 blocks, Achieved Occupancy 6.23%** (3.99 of 64 warps/SM), **DRAM 0.17%, Compute 0.47%, 62.7% barrier stalls** → pure memory-latency bound with too few warps to hide it (each active SM stalls ~88% of the time). Unlike the occupancy-maxed QMoE GEMVs, this kernel has structural headroom. **It also scales linearly with context and is the #1 decode kernel at realistic prompt lengths:** | context | attention_row share of GPU decode | |---|---| | short "Hello" (~65 keys) | 12.4% | | ~500-tok prompt, 64 tok | 25.3% | | ~500-tok prompt, 256 tok | **40.0%** | ## A/B (H200, PINNED `CUDA_VISIBLE_DEVICES`, median of 5) | workload | baseline (128) | this PR (256) | delta | |---|---|---|---| | V2-Lite, ~500-tok prompt, 128 tok — **attention_row kernel** | 172 µs | 148 µs | **−14%** | | V2-Lite, ~500-tok prompt, 128 tok — **E2E tok/s** | 81.94 | 86.47 | **+5.5%** | | V2-Lite, short "Hello", 128 tok — E2E tok/s | 127.96 | 128.27 | neutral | | dense qwen2.5-0.5b-int4, long ctx — E2E tok/s | 584.41 | 584.67 | neutral | 512 threads was no better than 256 (stage-5 P·V only has 128 channels), so 256 is the decode default. ## Byte-identity / numerics gates (PINNED, single-threaded) - 128-tok token-id **md5 unchanged**: V2-Lite (short + long ctx) and dense qwen2.5-0.5b. - **V2-Lite golden decode lock PASSES** (`deepseek_v2_lite_native_cuda_matches_golden_greedy_sequence`). - `standard_attention_fp16_gpu` (exact-rounded f32 oracle + fp32 accumulation), `standard_attention_bf16_gpu`, `standard_attention_capture_gpu` (decode capture == eager) all pass. - `standard_attention_gpu`: 23/24 pass; the 1 failure (`..._requires_homogeneous_floating_input_dtypes`, a dtype-claim validation test) is **pre-existing on origin/main** and unrelated to this change. - New unit test `row_threads_tests::decode_uses_wider_blocks_prefill_keeps_base` locks the decode(256)/prefill(128) widths. ## Scope / next lever This captures the latency-hiding win available on the **16 active SMs** byte-identically. The remaining headroom is the **116 idle SMs** — using them requires splitting each head's key-reduction across many blocks (split-KV / FlashDecoding), which reorders fp32 partials and needs the #1150 f64-tol path + a cross-block combine. Recommended as the scoped follow-up lever; this PR is the clean, byte-identical first step. Do NOT self-merge — for admin review/merge. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
…1339) Records in the campaign brain: - **#1337** attention_row decode block-widening (byte-identical, −14% kernel / +5.5% E2E at realistic prompt) — merged @ 37bdefe - **QMoE pipeline-fusion NO-GO** (profile-refuted: intermediate round-trip 0.075% of DRAM traffic) - **QMoE int4 ORT-idea port NO-GO** (occupancy-spill +19.8% regression) Both QMoE decode levers now exhausted with evidence; moat kernels at M=1 decode roofline. Next lever: split-KV/FlashDecoding to fill the 116 idle SMs. Docs-only (.squad/identity/now.md). Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
## Problem Every PR runs the full Rust matrix in `ci.yml` (build/test/clippy/coverage/CUDA compile across Linux/Windows/macOS) plus `benchmark.yml` and `audit.yml`, even when the diff is pure `docs/`, `wiki/`, `site/`, or Markdown content. ## Constraint that shaped the design `Fast (Linux x86_64)` and `Rust quality` are **required status checks** (repo ruleset `merge rules`, confirmed via `gh api repos/justinchuby/onnx-genai/rulesets/20017687`). A plain `paths-ignore` on `ci.yml` would leave these two checks permanently "pending" on docs-only PRs and block merge forever — so that approach is unsafe for `ci.yml`. ## What changed **`ci.yml`** — added a `changes` job that diffs the PR (`base.sha`/`head.sha`) or push (`before`/`sha`) and classifies it `docs_only` only if *every* changed path matches `docs/**`, `wiki/**`, `site/**`, `*.md`/`*.mdx`/`*.markdown`, `LICENSE`, or `NOTICE`. Any other file (Cargo.lock, `scripts/**`, `.github/workflows/**`, `crates/**`, etc.) anywhere in the diff flips it back to full CI. All 7 Rust jobs (`fast-linux`, `rust-quality`, `rust-coverage`, `rust-windows-arm64`, `ep-conformance`, `cli-ort`, `cuda-compile-linux`, `cuda-compile-windows`) get `needs: changes` + `if: needs.changes.outputs.docs_only != 'true'`. A job skipped by `if:` reports conclusion `skipped`, which GitHub [treats as passing](https://docs.github.com/en/pull-requests/how-tos/merge-and-close-pull-requests/troubleshooting-required-status-checks) for required status checks — so `Fast (Linux x86_64)` / `Rust quality` still resolve (fast) on every PR instead of hanging pending. The `changes` job only triggers for `pull_request`/`push`; `schedule`/`workflow_dispatch` never run it, so `docs_only` stays unset and every `if:` evaluates true — those triggers are **unchanged**. **`benchmark.yml`** (PR-only) and **`audit.yml`** (push/PR/schedule/dispatch) are *not* required checks, so they just get a `paths-ignore` with the same allowlist on their `push`/`pull_request` triggers (path filters don't apply to `schedule`/`workflow_dispatch`, so those stay unchanged too). **Untouched:** `miri.yml` and `weight-cache-guard.yml` already scope to targeted Rust paths. `wiki-pages.yml` keeps validating site/wiki content on every PR/push. No new third-party actions; permissions and concurrency groups preserved everywhere. ## Validation - `python3 -c "import yaml; yaml.safe_load(...)"` — all 3 edited files parse. - `actionlint` on all 3 files — same 4 pre-existing shellcheck warnings as `main` (confirmed via diff against `main`'s `ci.yml`), no new issues. - Manually verified `needs`/`if` present on all 8 non-`changes` jobs in `ci.yml` via a small parser script. - Extracted the classifier shell logic and unit-tested it against a trigger matrix: | Scenario | Files | `docs_only` | |---|---|---| | Docs-only | `docs/x.md` | `true` | | Docs-only, multiple dirs | `README.md`, `wiki/page.md`, `site/quartz/config.ts` | `true` | | Mixed docs + Rust | `docs/x.md`, `crates/foo/main.rs` | `false` (full CI) | | Cargo.lock only | `Cargo.lock` | `false` (full CI) | | CI workflow change | `.github/workflows/ci.yml` | `false` (full CI) | | Script change | `scripts/check_publish_order.py` | `false` (full CI) | | Empty/unclassifiable diff | (none) | `false` (fail-safe full CI) | | `schedule` / `workflow_dispatch` | n/a | `changes` job doesn't run → unset → full CI (unchanged) | Closes: none (proactive CI optimization, not tied to an issue). Co-authored-by: Copilot App <223556219+Copilot@users.noreply.github.com>
…Lite wide-ctx) (#1340) ## Summary `attention_row` launches **one block per (batch, q_head, query) row**, so at decode it uses only ~16 of 132 SMs and is memory-latency bound. Profiling V2-Lite eager decode at wide context (GPU3, ~390-word prompt, 256 tok) showed `attention_row` is the **#1 decode kernel at 36% of GPU decode** — ncu: **grid=16, waves/SM=0.02 (~116 idle SMs), DRAM 0.79%, achieved occupancy 9.8%** → a definitive under-utilization signal. This PR adds a **two-kernel FlashDecoding split-KV path** that spreads each row's key reduction across many blocks to fill the idle SMs. ## Design Behind `ONNX_GENAI_ATTN_SPLITKV` (**default ON**; `=0` reproduces the monolithic baseline for A/B): - **`attention_split`** — grid `(total_rows, num_splits)`. Each block reduces one KV slice with a chunk-local online softmax (running max + rescale), writing an *unnormalized* partial P·V plus per-split `max/sum` meta. A **single-split fast path** (`total_seq <= chunk`) reproduces `attention_row` exactly and writes the *final normalized* output with sentinel meta (`max=0, sum=1`). - **`attention_combine`** — grid `(total_rows)`. Uniform log-sum-exp merge of the per-split partials. The single-split sentinel makes it a **bit-exact pass-through**, so any context `<= chunk` (default 256) stays **byte-identical** to the monolithic kernel. **Capture-safety:** `num_splits` and the engage decision derive **only from the fixed KV capacity (`cap`)**, never live seqlen, so eager and capture make identical launch decisions → the `eager==capture` long-context lock holds. Split scratch is drawn from the internal **capture-warmed ws pool** (WS_SPLIT slot), leaving the governed workspace layout (and its exact-byte unit tests) untouched. Multi-split reorders fp32 partials, so wide context is validated to the **#1150 f64 tolerance** (`2e-4.max(·*2e-5)`), not byte-identity. Greedy token output is unchanged in practice. ## Numerics gates (H200, GPU3, pinned, single-threaded) | Gate | Result | |---|---| | V2-Lite golden **24-tok** lock (`matches_golden`) | ✅ **byte-identical** (single-split path) | | V2-Lite **340-tok** long-context lock (`eager==capture` + golden prefix) | ✅ pass (multi-split engaged, capture-safe) | | `standard_attention_capture_gpu` / `_fp16_gpu` / `_bf16_gpu` | ✅ pass | | `standard_attention_gpu` | ✅ 23/24 (the 1 failure `..._requires_homogeneous_floating_input_dtypes` is **pre-existing on clean origin/main**, unrelated) | | layout unit tests (`--lib standard_attention`) | ✅ 14/14 | | Dense qwen2.5-0.5b token md5 | ✅ byte-identical (no regression) | ## Perf A/B (H200 GPU3, pinned, medians-of-5) **Wide-context kernel (nsys, same binary, split ON vs `=0`):** | | decode attention kernel (median) | DRAM throughput | |---|---|---| | baseline (`attention_row`) | 147.6 µs | 0.79% | | split (`attention_split`+`combine`) | ~62.5 + 2.5 µs (**~2.3× faster**) | 9.26% | **E2E throughput (tok/s):** | context | split OFF (`=0`) | split ON (default) | Δ | |---|---|---|---| | **wide** (~390-word prompt, 128 tok) | 51.64 | **56.90** | **+10.2%** | | short "Hello" eager (128 tok) | 62.25 | 63.04 | +1.3% (neutral) | | short "Hello" capture (128 tok) | 128.48 | 128.17 | −0.2% (neutral) | | dense qwen2.5-0.5b capture | 610.36 | 610.25 | neutral | `num_splits = cap.div_ceil(chunk).clamp(1, 64)` with `chunk` default 256 (tunable via `ONNX_GENAI_ATTN_SPLIT_CHUNK`). At the true long-context regime where `attention_row` reaches ~40% of decode (cap≈2048), this yields ~8 splits → grid ~128, filling the machine; smaller chunks extract a bit more at moderate context at the cost of engaging split earlier. ## Verdict: **GO**, default-ON Short context stays byte-identical (comfortable margin below `chunk`); wide context gets +10.2% E2E / ~2.3× kernel; capture-safe; no dense regression. All hard gates green. Do **not** self-merge — reporting to coordinator for admin-merge after review. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
…log (PR #1337, #1340) (#1343) ## Summary Scribe consolidation of two Deckard decision inbox notes into the canonical `.squad/decisions.md` ledger. ### Decisions logged **PR #1337 — attention_row block-width 128→256 (GO, byte-identical, merged 37bdefe)** - Profile evidence: grid=16, achieved occupancy 6.23%, waves/SM=0.01, 62.7% barrier-stall cycles — memory-latency bound, not roofline-limited. - attention_row context scaling: 12.4% GPU decode share at short ctx → 40% at deep ctx → **#1 decode kernel at realistic/wide context** (QMoE levers exhausted). - Numerics gate: BYTE-IDENTICAL (same accumulation order, more warps). - H200 A/B: −14% attention_row kernel latency, **+5.5% E2E** on V2-Lite ~500-tok prompt; neutral on short/dense. **PR #1340 — split-KV FlashDecoding for attention_row (GO, default-ON, merged 763d81f)** - Profile evidence: grid=16, waves/SM=0.02, ~116 idle SMs, DRAM 0.79%, occupancy 9.8% — machine starved. - Single-split sentinel fast path: contexts ≤ chunk (default 256) → **byte-identical pass-through**. - Wide-ctx (multi-split): reordered fp32 partials → **f64-tol** (#1150 oracle); capture-safe (num_splits derived from fixed KV cap, not live seqlen). - H200 A/B: ~2.3× kernel speedup, DRAM 0.79%→9.26%, **+10.2% E2E** wide ctx; neutral short/dense. ### Archive gate All entries in decisions.md are dated 2026-08-18 (active campaign arc) — nothing >30 days old. No archiving performed per the size-not-age rule. ### File scope - Modified: `.squad/decisions.md` (+86 lines) - Deleted (untracked, rm only): `.squad/decisions/inbox/deckard-attention-row.md`, `.squad/decisions/inbox/deckard-attention-splitkv.md` Do NOT self-merge — awaiting coordinator review. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
… constant Review follow-ups on the degree-6 pure-Horner `exp8`, all documentation and test-domain only - the kernel's arithmetic is unchanged. * `UNDERFLOW` is now a single module constant that both `exp8`'s clamp and the accuracy test's domain endpoint read, so the two cannot drift apart. * The accuracy test swept `[-87.336, -0.0]` and therefore left the ~70-pattern sliver down to the real flush threshold `-87.336544` uncovered. It now sweeps to exactly one ULP inside that threshold - the most negative argument `exp8` still approximates. At and below the threshold `exp8` returns exactly `0` by contract, which is not an approximation error this bar applies to, so the endpoint stops one ULP short rather than swallowing the flush point. * Corrected two comment claims: the stride being coprime with 8 is not what makes every lane traverse the domain (each lane marches by `8 * stride`, which holds for any stride), and the offline exhaustive sweep covers the nonzero domain `(-87.336544, 0]`, not `[-87.336, 0]`. Re-verified after the change: the documented `c2` mutation (`4.999_999_1e-1` -> `4.999_970_0e-1`) still fails only this test, with `max_ulp=6 at x=-0.32134318`. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
Owner
Author
Re-validated on latest
|
| artifact | evidence |
|---|---|
bench_generic (--no-default-features --features bench-native,cuda-13000), both arms |
nm -C | grep -ci mlas = 0 |
libonnx_runtime_ep_cpu_plugin.so, default features, release |
nm -CD | grep -ci mlas = 0 |
default dep trees of onnx-genai-cli, onnx-genai-server, onnx-genai-capi, onnx-genai-python, onnx-runtime-session, onnx-runtime-ep-cpu-plugin |
cargo tree -e normal | grep -ci mlas = 0 each |
mlas remains opt-in (onnx-runtime-ep-cpu: default = ["full"], mlas = ["dep:mlas-sys"]),
and the --features mlas lane still compiles.
Numerical contract — independent exhaustive re-run
I wrote my own harness that transcribes both exp8 forms verbatim (old Cephes from
origin/main, new pure-Horner from this branch), uses the real intrinsics, and sweeps
every f32 bit pattern in the nonzero softmax domain (-87.336544, -0.0] —
1,118,743,631 values — against a correctly-rounded f64 reference:
| arm | exhaustive max ULP |
|---|---|
origin/main (Cephes) |
1 |
| this PR (degree-6 pure Horner) | 1 |
Special-value contract, PASS on both arms: exp8(-inf) = +0.0 bitwise ·
exp8(NaN) = NaN · exp8(-1e30) = +0.0 · exp8(-87.336544) (threshold, inclusive <=)
= +0.0 · threshold+1 ULP finite, nonzero, ≤1 ULP · exp8(±0.0) = 1.0 bitwise ·
mixed NaN/-inf/finite lane block shows no cross-lane contamination · strided monotonicity holds.
The exponent build is safe by construction over that domain: x·log2e ≥ -126, so
biased = k+127 ≥ 1 and no nonzero lane ever builds a zero/negative exponent field;
exp8 never emits a subnormal (threshold+1 ULP = 1.1755086e-38, just above the smallest normal).
Softmax matrix — production-native, latest main, 40 reps total
Two bench_generic binaries differing only in softmax.rs, md5-distinct, both MLAS-free.
Arms interleaved innermost in one driver process, arm order alternating between reps,
thread-matched --native-threads 1 --ort-intra-threads 1. Metric = native_min/ort_min
(lower is better, <1.000 beats ORT).
Pass 1 (24 reps) — before → after, min / median ratio:
| model | before | after | Δmin | Δmed |
|---|---|---|---|---|
| sm_bert_b8_s128 | 1.445 / 1.486 | 1.287 / 1.340 | −10.97% | −9.89% |
| sm_decode_h32_kv1024 | 1.471 / 1.471 | 1.294 / 1.348 | −12.00% | −8.35% |
| sm_decode_h32_kv2048 | 1.484 / 1.533 | 1.318 / 1.383 | −11.17% | −9.78% |
| sm_decode_h32_kv4096 | 1.429 / 1.475 | 1.311 / 1.328 | −8.20% | −10.00% |
| sm_decode_h32_kv8192 | 1.472 / 1.534 | 1.266 / 1.329 | −13.99% | −13.33% |
| sm_prefill_h32_s512 | 1.377 / 1.396 | 1.185 / 1.247 | −13.96% | −10.63% |
| sm_whisper_cross | 1.227 / 1.240 | 1.102 / 1.115 | −10.20% | −10.03% |
Pass 2 (16 reps, rebuilt from the final tree):
| model | before | after | Δmin | Δmed |
|---|---|---|---|---|
| sm_bert_b8_s128 | 1.386 / 1.530 | 1.294 / 1.344 | −6.63% | −12.18% |
| sm_decode_h32_kv1024 | 1.389 / 1.471 | 1.294 / 1.305 | −6.82% | −11.26% |
| sm_decode_h32_kv2048 | 1.468 / 1.533 | 1.323 / 1.383 | −9.91% | −9.78% |
| sm_decode_h32_kv4096 | 1.427 / 1.532 | 1.306 / 1.328 | −8.43% | −13.30% |
| sm_decode_h32_kv8192 | 1.448 / 1.572 | 1.283 / 1.335 | −11.39% | −15.09% |
| sm_prefill_h32_s512 | 1.379 / 1.403 | 1.203 / 1.256 | −12.75% | −10.50% |
| sm_whisper_cross | 1.147 / 1.240 | 1.107 / 1.123 | −3.50% | −9.38% |
parity=PASS on every cell of every rep in both passes. 7/7 models improve on both
min and median in both passes; no model regressed anywhere. Combined range
−3.5% to −15.1%, i.e. this is a slightly larger win on current main than the
4–10% this PR opened with — but the host is shared with other agents (load average ~23
on 32 cores), so treat the magnitudes as ±a few points and the direction (7/7, twice,
40 reps) as the claim. Softmax is still 1.10–1.38x ORT; this does not close the gap.
Checks on the final tree
cargo test -p onnx-runtime-ep-cpu --lib: 1418 passed, 0 failed, 17 ignoredcargo clippy -p onnx-runtime-ep-cpu --lib --all-targets: clean (0 warnings)cargo fmt -p onnx-runtime-ep-cpu -- --check: cleancargo check -p onnx-runtime-ep-cpu --lib --features mlas: compiles
Independent review (Opus) — APPROVE-WITH-NITS, no blocking findings
The reviewer re-derived the exponent-boundary argument, checked the coefficients are a
genuine Remez fit rather than a typo (c2/c3 on Taylor, c4/c5/c6 deviating as Remez does;
polynomial relative error ≈1.7e-8 ≈ 0.29 f32-eps), and confirmed the op-count claim exactly:
old = 6 FMA + 1 MUL + 1 ADD = 8 vector ops, new = 6 FMA = 6.
Four nits were raised; 4150867ba addresses the two that were actionable:
- The test swept to
-87.336, leaving the ~70-pattern sliver down to the true flush
threshold-87.336544uncovered. It now sweeps to exactly one ULP inside the threshold,
and both the clamp and the test endpoint read one sharedUNDERFLOWconstant so they
cannot drift apart. - "stride coprime with 8 so every lane sees the whole domain" was a bogus justification
(each lane marches by8*stride, true for any stride). Comment corrected, as was the
[-87.336, 0]→(-87.336544, 0]domain claim.
Not changed, and why: the ≤2 ULP bar stays above the measured 1 ULP (double-rounding
headroom in the f64 reference), and the test is necessarily absent under --features mlas
/ non-x86 and vacuous without AVX2 — unavoidable for a #[target_feature] intrinsic.
Falsification re-run after the edit: the documented c2 mutation
(4.999_999_1e-1 → 4.999_970_0e-1) still fails only this test —
max_ulp=6 at x=-0.32134318 — while all 21 row-level softmax tests pass.
I also probed a c4 mutation (4.166_822_5e-2 → 4.166_600_0e-2, 5.3e-5 relative);
the test passed, and my exhaustive harness confirms that variant genuinely stays at
1 ULP — so that is coefficient slack, not a hole in the test.
justinchuby
marked this pull request as ready for review
August 19, 2026 00:13
justinchuby
added a commit
that referenced
this pull request
Aug 19, 2026
… 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>
justinchuby
added a commit
that referenced
this pull request
Aug 19, 2026
…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>
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
1 participant
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
What
Replaces the Cephes
1 + r + r^2*P5(r)form ofexp8inmod softmax_avx2with adegree-6 Remez minimax polynomial evaluated in pure Horner form, whose two leading
coefficients are pinned to exactly
1.0so the1 + rhead stays exact.Net effect per 8-lane
exp: one fewer FMA-unit op and one fewer FADD (the separater*rmultiply and the trailing+1both disappear into the Horner chain). TheCody-Waite two-part
ln2range reduction is unchanged - see "what I rejected".Why this, and why now
#1234 vectorised this
expand took the softmax family from 9.3-10.9x ORT to1.29-1.63x. Softmax is now the most-lost kernel family in the CPU EP.
I had attributed that residual to the three-pass memory structure (max / exp+sum /
normalize). That hypothesis is falsified. A two-pass online/streaming softmax
measured 2.4x worse across all 7 models, with the loss growing at the largest rows -
the opposite of a bandwidth-bound signature. The streaming recurrence needs an
expfor both the increment and the max-rescale, ~3 exp/element against the three-pass
form's 1. That negative result is being written up separately.
So the residual is exp throughput, and this PR attacks exactly that. It is a
modest win, and I am not going to dress it up as more.
Measurements - production-native only
Build:
--no-default-features --features bench-native,cuda-13000.Falsifier:
nm -C bench_generic | grep -ci mlas= 0 for both arms (binariesverified distinct by md5). Thread-matched
--native-threads 1 --ort-intra-threads 1,arms interleaved innermost in one driver process with arm order alternating between
reps. Ratio = native_min/ort_min; lower is better, below 1.000 beats ORT.
Two independent 8-rep passes of my own (16 reps total), on top of the implementing
agent's 12- and 14-rep passes. Medians:
parity=PASSon every cell of every rep. Pass 2 improves 7/7 on both min and median;pass 1 improves 7/7 on min and 6/7 on median (kv1024 median flat at 1.471 -> 1.477,
which pass 2 resolves to 1.412). Roughly a 4-10% ratio improvement, no model
regressed.
One honesty note on methodology: the absolute native times are not comparable across
arms on this shared, bimodal box - one pass showed bert at 0.890 "native speedup" while
its ratio vs ORT clearly improved, because contention moved under both arms. Since ORT
is byte-identical code in both binaries and is timed in the same process invocation
as native, the paired ratio is the internally-controlled metric, and it is the one I am
quoting.
Accuracy
Exhaustive sweep over all 1.118e9 f32 values in the softmax domain
[-87.336, 0]against a correctly-rounded f64 reference:
exp8: 1 ULPexp8: 1 ULPNo accuracy is spent to buy the speed.
New test, and what it catches
exp8_stays_within_two_ulp_of_the_f64_referenceprobesexp8directly, because therow-level tests structurally cannot police it: softmax divides every exponential by the
row sum, so relative error shared across lanes cancels in the ratio. An
exp8variant measured at 64 ULP against f64 still matched the scalar softmax to 2.4e-7.
I falsified it with my own mutation, distinct from the implementing agent's: perturbing
c2by ~6e-6 relative (4.999_999_1e-1->4.999_970_0e-1).That mutation fails only the new test - all 21 pre-existing row-level softmax tests
pass, including
out_of_place_softmax_is_bit_identical_on_pathological_rows,large_and_masked_logits_stay_finiteanda_fully_masked_row_reproduces_the_reference_runtimes_nan. So the test closes a realhole rather than restating existing coverage.
What I rejected, and the measurement that decided it
VROUNDPS)t -> kchainexp2formulation (fast, 1-part log2e)f = x*log2e - kdrops low bitsexp2formulation (accurate, 2-part log2e)The
exp2result is the informative one: cutting the two Cody-Waite range-reductionFMAs is worth ~15%, but only by destroying accuracy; doing it accurately gives the
FMAs back and erases the gain. The Cody-Waite reduction is load-bearing and I left
it alone. The only accuracy-free lever left was the polynomial's shape.
Semantics preserved
-infand deepunderflow -> exactly
0.mlasremains non-default.Checks
cargo test -p onnx-runtime-ep-cpu --lib: 1348 passed, 0 failed, 16 ignoredcargo clippy -p onnx-runtime-ep-cpu --lib: cleancargo fmt -p onnx-runtime-ep-cpu -- --check: cleancargo check -p onnx-runtime-ep-cpu --lib --features mlas: compiles - explicitlyre-verified on the final tree, because in perf(cpu-ep): vectorise the production softmax exp on AVX2+FMA #1234 I shipped a
const {}assert thatbroke this exact lane after I had validated an earlier state.
Standing
This does not close the softmax gap - it is still 1.17-1.50x ORT. It removes ~4-10% of
it without spending accuracy, and it narrows the remaining cause: with the memory-
structure hypothesis falsified and the range reduction proven load-bearing, what is
left is ORT's exp/vectorisation quality or per-call overhead. Still not complete.