Skip to content

fix(cli): render structured model responses - #1161

Closed
justinchuby wants to merge 20 commits into
split/native-backend-diagnosticsfrom
split/cli-structured-responses
Closed

justinchuby wants to merge 20 commits into
split/native-backend-diagnosticsfrom
split/cli-structured-responses

Conversation

@justinchuby

Copy link
Copy Markdown
Owner

Summary

  • apply every metadata-declared EOS token to pipeline generation
  • render ATEM channels and special tokens without dropping debug-visible output
  • keep only final answer content in conversation history

Dependency

Stacked on the native diagnostics/context PR because it uses that pipeline context plumbing.

Validation

  • 116 CLI tests
  • 68 pipeline unit tests

@github-actions

Copy link
Copy Markdown

🔴 Benchmark Regression Detected

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

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

Status Scenario Base PR Change
🔴 matmul/large_generic_f32_threads=8/32x1024x1024 4.26 ms 7.43 ms +74.5%
🔴 matmul/medium_generic_f16_threads=8/32x512x512 30.29 µs 45.78 µs +51.1%
🔴 matmul/medium_generic_bf16_threads=8/32x512x512 376.42 µs 513.99 µs +36.5%
⚠️ matmul/large_generic_f16_threads=1/32x1024x1024 81.74 µs 99.56 µs +21.8%
⚠️ qwen3_sampling_processors/top_k_full_sort_baseline 1.96 ms 2.33 ms +19.0%
⚠️ tokenization/encode_tokens_per_second 370.90 µs 440.59 µs +18.8%
⚠️ matmul/large_generic_bf16_threads=8/32x1024x1024 1.43 ms 1.66 ms +15.7%
✅ kv_cache/alloc_dealloc_pages 36.16 µs 41.39 µs +14.4%
✅ add/small_f32_threads=1-internal/1024 228.0 ns 259.8 ns +14.0%
✅ tokenization/decode_tokens_per_second 5.85 ms 6.63 ms +13.4%
✅ add/large_bf16_threads=1-internal/4194304 1.62 ms 1.83 ms +12.9%
✅ sampling_latency/top_p_per_token 353.84 µs 397.97 µs +12.5%
✅ sampling_latency/min_p_per_token 198.69 µs 222.10 µs +11.8%
✅ block_quantized_moe_cached_dense/mxfp4_uncached_expert_dequant_each_call/rows=1,H=256,I=256,E=4,top_k=1 426.30 µs 472.60 µs +10.9%
✅ add/small_bf16_threads=1-internal/1024 473.0 ns 522.3 ns +10.4%
✅ add/large_f32_threads=1-internal/4194304 582.76 µs 639.09 µs +9.7%
✅ add/medium_bf16_threads=1-internal/262144 101.52 µs 109.42 µs +7.8%
✅ qwen3_sampling_processors/top_k_top_p_full_sort_baseline 6.27 ms 6.75 ms +7.6%
✅ matmul/medium_generic_bf16_threads=1/32x512x512 603.76 µs 647.92 µs +7.3%
✅ add/large_f16_threads=1-internal/4194304 1.66 ms 1.77 ms +6.9%
✅ block_quantized_matmul_cached_dense/mxfp4_preexpanded_dense_oncelock_like_proxy/1x1024x1024 56.32 µs 60.02 µs +6.6%
✅ qwen3_sampling_processors/top_k_partial_selection 139.66 µs 148.52 µs +6.3%
✅ add/small_f16_threads=1-internal/1024 437.8 ns 462.6 ns +5.7%
✅ qwen3_sampling_processors/top_p_fast_after_top_k 535.49 µs 565.68 µs +5.6%
✅ sampling_latency/greedy_per_token 3.00 µs 3.16 µs +5.3%
✅ add/medium_f16_threads=1-internal/262144 102.22 µs 106.90 µs +4.6%
✅ sampling_latency/top_k_per_token 49.17 µs 50.50 µs +2.7%
✅ reduce_mean/small_f32_threads=1-internal/4096 14.65 µs 14.67 µs +0.1%
✅ qwen3_sampling_processors/top_k_top_p_fast 715.65 µs 709.71 µs -0.8%
✅ add/medium_f32_threads=1-internal/262144 25.69 µs 25.40 µs -1.1%
✅ matmul/large_generic_f16_threads=8/32x1024x1024 106.71 µs 105.26 µs -1.4%
✅ logit_processing/seven_processor_chain_per_step 348.87 µs 341.50 µs -2.1%
✅ gather/small_bf16_threads=1-internal/4096 455.0 ns 438.9 ns -3.5%
✅ gather/small_f32_threads=1-internal/4096 655.5 ns 631.5 ns -3.7%
✅ matmul/small_generic_bf16_threads=1/1x256x256 30.59 µs 29.28 µs -4.3%
✅ block_quantized_matmul_cached_dense/mxfp4_cached_dense_repeated_call/1x1024x1024 49.73 µs 47.35 µs -4.8%
✅ matmul/large_generic_f32_threads=1/32x1024x1024 10.94 ms 10.37 ms -5.2%
✅ matmul/large_generic_bf16_threads=1/32x1024x1024 2.06 ms 1.93 ms -6.4%
✅ matmul/medium_generic_f32_threads=1/32x512x512 2.33 ms 2.18 ms -6.6%
✅ gather/medium_f32_threads=1-internal/32768 3.71 µs 3.44 µs -7.2%
✅ block_quantized_moe_cached_dense/mxfp4_cached_dense_expert_repeated_call/rows=1,H=256,I=256,E=4,top_k=1 182.02 µs 168.05 µs -7.7%
✅ grammar_masking/llguidance_compute_mask/32 91.14 µs 84.02 µs -7.8%
✅ matmul/small_generic_f16_threads=1/1x256x256 30.86 µs 28.42 µs -7.9%
✅ reduce_mean/large_f32_threads=1-internal/262144 997.12 µs 918.07 µs -7.9%
✅ qwen3_sampling_processors/top_p_full_sort_after_top_k_baseline 4.14 ms 3.79 ms -8.4%
✅ gather/medium_f16_threads=1-internal/32768 2.47 µs 2.26 µs -8.5%
✅ gather/medium_bf16_threads=1-internal/32768 2.48 µs 2.25 µs -9.3%
✅ gather/large_bf16_threads=1-internal/131072 12.72 µs 11.54 µs -9.3%
✅ gather/small_f16_threads=1-internal/4096 498.0 ns 450.2 ns -9.6%
✅ block_quantized_matmul_cached_dense/mxfp4_uncached_dequant_each_call/1x1024x1024 657.04 µs 583.41 µs -11.2%
✅ gather/large_f32_threads=1-internal/131072 29.41 µs 25.90 µs -12.0%
✅ reduce_mean/medium_f32_threads=1-internal/65536 286.24 µs 250.33 µs -12.5%
✅ matmul/medium_generic_f16_threads=1/32x512x512 32.72 µs 28.09 µs -14.1%
🟢 matmul/small_generic_f32_threads=1/1x256x256 40.63 µs 34.24 µs -15.7%
🟢 matmul/small_generic_f16_threads=8/1x256x256 35.79 µs 29.91 µs -16.4%
🟢 gather/large_f16_threads=1-internal/131072 14.04 µs 11.41 µs -18.8%
🟢 matmul/small_generic_f32_threads=8/1x256x256 39.82 µs 32.33 µs -18.8%
🟢 matmul/medium_generic_f32_threads=8/32x512x512 1.10 ms 887.07 µs -19.0%
🟢 matmul/small_generic_bf16_threads=8/1x256x256 36.78 µs 28.54 µs -22.4%

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

Host info
CPU: Apple M1 (Virtual)
Cores: 3
OS: Darwin 25.5.0 arm64
Rust: rustc 1.97.1 (8bab26f4f 2026-07-14)
Load avg: { 6.32 4.44 5.05 }
What this cannot catch
  • Regressions in code paths not covered by these benchmarks (e.g., end-to-end decode with a real model)
  • Sub-threshold regressions that compound over multiple PRs
  • Performance changes that only manifest under GPU execution
  • Latency changes in the ORT integration path (these benchmarks exercise the native Rust kernels)

@codecov

codecov Bot commented Aug 18, 2026 •

Copy link
Copy Markdown

Codecov Report

❌ Patch coverage is 59.64126% with 90 lines in your changes missing coverage. Please review.
✅ Project coverage is 80.60%. Comparing base (266a6fe) to head (7c7dc0c).
⚠️ Report is 6 commits behind head on split/native-backend-diagnostics.

Files with missing lines Patch % Lines
crates/onnx-genai-cli/src/output.rs 63.00% 68 Missing and 6 partials ⚠️
crates/onnx-genai-cli/src/interactive.rs 35.71% 3 Missing and 6 partials ⚠️
...es/onnx-runtime-ep-cpu/src/kernels/matmul_nbits.rs 0.00% 4 Missing ⚠️
crates/onnx-genai-cli/src/generate.rs 40.00% 1 Missing and 2 partials ⚠️
Additional details and impacted files

Impacted file tree graph

@@                         Coverage Diff                          @@
##           split/native-backend-diagnostics    #1161      +/-   ##
====================================================================
+ Coverage                             80.58%   80.60%   +0.01%     
====================================================================
  Files                                   370      370              
  Lines                                162350   162560     +210     
  Branches                             162350   162560     +210     
====================================================================
+ Hits                                 130827   131026     +199     
+ Misses                                26786    26784       -2     
- Partials                               4737     4750      +13     
Flag Coverage Δ
cli-ort-linux 82.86% <60.73%> (-0.94%) ⬇️
cli-ort-windows 82.39% <60.73%> (-0.92%) ⬇️
mlas 85.67% <ø> (-0.24%) ⬇️
offline 80.41% <0.00%> (+0.05%) ⬆️

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

Files with missing lines Coverage Δ
crates/onnx-genai-cli/src/generate.rs 90.45% <40.00%> (-0.58%) ⬇️
...es/onnx-runtime-ep-cpu/src/kernels/matmul_nbits.rs 83.19% <0.00%> (+1.07%) ⬆️
crates/onnx-genai-cli/src/interactive.rs 75.54% <35.71%> (-0.79%) ⬇️
crates/onnx-genai-cli/src/output.rs 77.04% <63.00%> (-6.80%) ⬇️

... and 2 files with indirect coverage changes

🚀 New features to boost your workflow:
  • ❄️ Test Analytics: Detect flaky tests, report on failures, and find test suite problems.

@justinchuby

Copy link
Copy Markdown
Owner Author

Reviewed and validated. Panic safety checked at the boundary that matters: load_response_config and its helpers are Option/?-based, so a malformed or missing
esponse_template falls back to plain rendering rather than unwrapping; the slice boundaries sit immediately after the ASCII markers |>/<| so they are always valid UTF-8 boundaries. Chunked streaming is safe because ResponseStream::push dispatches on oken_id, not on text, so a structure split across chunks still transitions. Falsified: flipping the Content arm's include_in_reply to alse reddens the formatting test (left: , right: answer). 116 CLI lib tests pass, clippy clean. One non-blocking follow-up filed separately. Merging.

justinchuby and others added 14 commits August 17, 2026 19:46
## Summary
- keep the default INFO log level
- buffer tracing output during token streaming and flush after the reply
- prevent diagnostics from overwriting or splitting generated text

## Validation
- 115 CLI unit tests

---------

Co-authored-by: justinchuby <223556219+Copilot@users.noreply.github.com>
Copilot-Session: 0190e2eb-abe4-451f-b36d-44a035a99b7e
…ditive masks (V2-Lite) (#1171)

## Summary
Topology-gated extension of the executor capacity policy so a
**vestigial/global** attention mask that flows through the standard
**additive causal-mask builder** into a **capacity-form `Attention`**
(DeepSeek-V2-Lite / MLA) is classified **padded-capacity-safe**,
clearing the CUDA-graph-capture decline
`attention_mask_consumers_are_capacity_aware`. This is the classifier
half of Wallace's V2-Lite graph-capture GO brief.

**Not a blanket flip.** The builder op set (`CumSum, Unsqueeze, Cast,
GreaterOrEqual, And, Where, Slice, Sub`, derived from the real V2-Lite
mask cone) excludes broadcasting logical-width combiners — notably
GLM-5.2's indexer `Add` — so GLM-5.2 keeps exposing its logical valid
length. Dense GQA masks keep the existing `Shape`/`ReduceSum` fast path.
No kernel edits; still behind opt-in `ONNX_GENAI_CUDA_GRAPH`.

## Change surface
- `executor/geometry.rs`: `is_additive_mask_builder_op`,
`is_capacity_form_attention_mask_input`,
`mask_binding_feeds_capacity_form_attention` (structural forward cone
walk).
- `executor/build.rs`: `binding_consumers_use_padded_capacity` keeps the
fast path, then tries the topology-gated path.
- `deepseek_v2_tiny_qmoe_native_e2e.rs`: flips the `=1` expectation —
the mask reason must no longer decline capture.

## Tests / validation
- 7 new unit tests lock the vestigial-window reroute **and** the GLM-5.2
exclusion (+ is_causal / no-Attention / non-builder negatives).
`only_capacity_aware_inputs_keep_physical_capacity` untouched, still
green. 176/176 session lib tests pass; `fmt`/clippy clean on touched
crates.
- Tiny DeepSeek-V2 fixture, `ONNX_GENAI_CUDA_GRAPH=1` (GPU5):
`captures=1 replays=6 decline=None`, tokens == CPU golden ⇒ **capture
now engages, byte-identical**. `=0` still declines via the opt-in
predicate.
- Real V2-Lite int4: eager `=0` unaffected (57 tok/s).

## ⚠️ Second, pre-existing blocker (not fixed here — flagged for
follow-up)
Real 27-layer V2-Lite `=1` capture does **not** yet engage: the additive
mask-bias `v_model.Unsqueeze_18` has a **query-length axis `_d1`** (axis
2) that is an unresolved internal symbol, so capture-prep workspace
planning cannot reserve the `Attention` workspace. **This predates and
is independent of this PR** — baseline `main` already crashes under `=1`
on the same node (`workspace invariant mismatch: requires 2648584,
prepared 57344`, a silent under-reservation). My change only makes the
failure earlier/principled (the planner refuses to guess `_d1`). Needs a
shape/symbol-resolution + capacity-reservation fix (Leon /
shape-inference) before the tok/s win lands. Full analysis in
`.squad/decisions/inbox/deckard-v2lite-graphcapture.md`.

⚠️ **Flagged as capacity/parity-adjacent — @rachael must review before
merge. Do not self-merge.** Wallace runs the hardware
byte-identity/tok-s A/B once the `_d1` blocker lands.

Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

---------

Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
… layout (#1175)

## What

Two changes to the CPU EP's `RotaryEmbedding` kernel:

1. **`rotary_units_per_task`** — Rayon tasks are now sized by *work*,
not by layout. The kernel used to fan out one task per layout unit: one
`[B,H,S,D]` plane or one `[B,S,H·D]` row. On `rope_llama3_s128` that is
4096 elements (16 KiB) per task, and the scheduling cost swamped the
work. Tasks now aggregate whole units until each clears 16 Ki elements,
targeting ~4 tasks per worker for balance, and collapse to a single
serial pass when the pool has one worker.
2. **Explicit AVX2 arms** for both rotation conventions
(`rotate_split_half`, `rotate_interleaved`). The inner loop was left to
the autovectoriser, which did not take it.

Plus a test-only change: the A/B transform grid now generates
`interleaved=1` (GPT-J) RoPE models, which were previously uncovered.

## Why

The old fan-out made the parallel path **slower than the serial path at
every thread count**. Measured on `rope_llama3_s128`,
`exec_kernel.compute` p50, native-only:

| threads | parallel (before) | serial (before) |
|---|---|---|
| 1 | 162 µs | 126 µs |
| 8 | 142 µs | 126 µs |
| 16 | **195 µs** | 127 µs |

Fanning 128 tasks out over a one-worker pool alone cost 36 µs of pure
Rayon bookkeeping.

After both changes, same measurement: **242 → 85 µs (t=1), 174 → 67 µs
(t=8), 138 → 59 µs (t=16)**.

### Why the vector arms avoid FMA

`_mm256_fmsub_ps` would drop the intermediate rounding of `cos·x1`, so
the 8-wide body and the scalar tail would disagree in the last ulp on
the same tensor. The same graph would then produce different bits
depending on how a thread count happened to split it, which flips the
A/B parity check on scheduling alone. Separate `_mm256_mul_ps` +
`_mm256_sub_ps` is bit-identical to the scalar loop. This kernel moves 8
bytes per 3 flops — it is bandwidth-bound, so fusing the multiply is not
measurable anyway.

## Measured: paired interleaved A/B vs ORT 1.27 CPU EP

One driver invocation, `scripts/ort_ab/ab.py`, 7 trials × 80 runs × 20
warmups per cell, `--threads` sets both `--native-threads` and
`--ort-intra-threads`. `base` = merge-base `8219fa250`, `new` = this
branch. Ratio is native/ORT p50, **lower is better**; bracket is the
min–max band across trials.

| model | t | base ratio | **new ratio** | new native p50 |
|---|---|---|---|---|
| rope_llama3_s1 | 1 | 1.347 [1.334–1.366] | **1.278** [1.251–1.291] |
0.006 ms |
| rope_llama3_s1 | 8 | 1.349 [1.307–1.404] | **1.273** [1.221–1.311] |
0.006 ms |
| rope_llama3_s1 | 16 | 1.399 [1.342–1.441] | **1.320** [1.275–1.327] |
0.008 ms |
| rope_llama3_b8_s1 | 1 | 1.227 [1.185–1.258] | **1.032** [1.007–1.054]
| 0.011 ms |
| rope_llama3_b8_s1 | 8 | 1.310 [1.253–1.718] | **1.145** [1.092–1.307]
| 0.012 ms |
| rope_llama3_b8_s1 | 16 | 1.028 [0.719–1.169] | **0.881** [0.797–0.961]
✅ | 0.017 ms |
| rope_llama3_s128 | 1 | 1.642 [1.631–1.653] | **0.800** [0.791–1.046] ✅
| 0.091 ms |
| rope_llama3_s128 | 8 | 4.372 [3.526–6.329] | **3.890** [3.149–4.139] |
0.127 ms |
| rope_llama3_s128 | 16 | 7.209 [6.217–9.948] | **6.111** [4.966–6.830]
| 0.189 ms |
| rope_llama3_s512 | 1 | 1.258 [1.163–1.286] | **0.929** [0.853–0.945] ✅
| 0.431 ms |
| rope_llama3_s512 | 8 | 4.485 [4.413–5.668] | **3.437** [2.221–4.180] |
0.296 ms |
| rope_llama3_s512 | 16 | 7.036 [6.439–8.881] | **5.574** [4.888–5.908]
| 0.348 ms |
| rope_gptj_il_s1 | 1 | 1.400 [1.375–1.429] | **1.301** [1.287–1.313] |
0.006 ms |
| rope_gptj_il_s1 | 8 | 1.385 [1.296–1.432] | **1.299** [1.269–1.358] |
0.006 ms |
| rope_gptj_il_s1 | 16 | 1.424 [1.351–1.453] | **1.300** [1.267–1.329] |
0.008 ms |
| rope_gptj_il_s128 | 1 | 1.911 [1.900–1.920] | **0.874** [0.861–0.909]
✅ | 0.098 ms |
| rope_gptj_il_s128 | 8 | 3.724 [3.369–5.509] | **3.628** [2.779–5.536]
| 0.124 ms |
| rope_gptj_il_s128 | 16 | 7.486 [7.140–8.042] | **6.146** [5.976–6.773]
| 0.195 ms |
| rope_gptj_il_s512 | 1 | 1.530 [1.494–1.560] | **0.964** [0.954–0.974]
✅ | 0.434 ms |
| rope_gptj_il_s512 | 8 | 4.533 [4.336–5.130] | **3.008** [2.418–3.450]
| 0.264 ms |
| rope_gptj_il_s512 | 16 | 7.655 [7.503–9.390] | **5.808** [5.037–7.063]
| 0.377 ms |

**21 / 21 cells improved. 0 regressed. 5 cells now beat ORT outright.
Parity 294/294 PASS.**

## This does not finish RoPE

Five cells win; sixteen still lose. The remaining loss is **not** in
this kernel — it is thread-pool wake latency, and it is the same root
cause behind every "wins at t=1, loses at t=8/16" row in the master
benchmark doc (§23.7: Softmax, Transpose, Concat, MoE all have that
shape).

Isolated measurement, 16 workers, one 524 288-element elementwise region
that costs 388 µs on one thread:

| gap between regions | region cost |
|---|---|
| none (back-to-back) | 67 µs |
| 20 µs | 226 µs |
| 100 µs | 230 µs |
| 500 µs | 236 µs |

A 20 µs gap — less than the serial glue between two graph nodes — turns
a 5.8× speedup into a 1.7× *slowdown*. Rayon parks its workers at the
end of every region; ORT's intra-op pool spins first
(`ALLOW_INTRA_OP_SPINNING`, default on), so its workers are still hot
when the next op arrives and ours are not. `exec_kernel.compute` for
`rope_llama3_s128` at t=16 is 59 µs against ORT's ~30 µs, but the paired
wall is 189 µs — the difference is wake latency, not arithmetic.

This repo already has the right mechanism: `decode_spmd`'s persistent
pool with a bounded `KMP_BLOCKTIME`-style active-spin window. It is
wired only to the decode GEMM path. A prototype routing this kernel's
fan-out through it takes `rope_llama3_s128` t=16 from **79 µs to 25 µs
native-only, past ORT's 30 µs** — but it oversubscribes when an ORT
session shares the process, so it needs SMT-aware worker placement to
land. That is the follow-up, tracked separately; this PR is the kernel
half and stands on its own.

## Tests

New, each falsified before being kept:

* `avx2_split_half_is_bit_identical_to_the_scalar_loop` — sweeps `half`
across 1…128 so the tail runs at every residue mod 8; compares
`to_bits()`, not a tolerance. *Falsified:* swapping one
`_mm256_mul_ps`/`_mm256_sub_ps` pair for `_mm256_fmsub_ps` fails it at
`half=16`.
* `avx2_interleaved_is_bit_identical_to_the_scalar_loop` — the
de-interleave uses a lane permutation that is its own inverse; a wrong
index vector still yields plausible magnitudes, so this compares every
element. *Falsified:* replacing the permutation with the identity fails
it at `half=8`.
* `task_chunking_covers_every_unit_exactly_once` — the derived tiling
must be a partition for 4 × 4 × 5 shape/pool combinations, because it
feeds `par_chunks_mut` where an off-by-one silently drops or
double-writes a row.
* `a_single_worker_pool_never_splits_the_rotation`
* `tasks_clear_the_minimum_work_bar_when_the_pool_is_wide` —
*falsified:* removing the minimum-task-size floor fails it with "task of
8192 elements is below the 16384-element bar".

Existing suites, all green after merging `origin/main` (`0bd8397f7`):

* `cargo test -p onnx-runtime-ep-cpu --lib` → **1328 passed, 0 failed**
* `cargo test -p onnx-runtime-ep-cpu-plugin --test plugin_ort_e2e` →
**53 passed, 0 failed** (real ORT 1.27 session,
`session.disable_cpu_ep_fallback=1`)
* `cargo fmt --all --check`, `cargo clippy -p onnx-runtime-ep-cpu
--all-targets` → clean

## No defer

No node is handed to ORT's CPU EP by this change. The claim policy is
untouched; `no_supported_node_is_ever_left_to_the_ort_cpu_ep` and
`every_fixture_loads_with_cpu_fallback_disabled` still pass.

## Host

AMD EPYC 9V74, 16 physical cores / 32 threads, one NUMA node, 32 MiB
shared L3, AVX2+FMA (no AVX-512). Shared and contended, so only paired
ratios from a single driver invocation are quoted; absolutes drift by
>4×.

---------

Co-authored-by: Leon <leon@squad.local>
Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
…#1178)

`onnx-genai-ort::cuda_rt` listed only CUDA 12 `cudart` names, while the
CUDA EP's own loader
(`onnx-runtime-ep-cuda::dynamic_library::candidates`) already listed
CUDA 13. Two lists for one job, drifted apart.

**Why it looked like something else.** On a CUDA 13 host every candidate
in this list fails, and the error reports the *last* one tried: `could
not load CUDA runtime (cudart): libcudart.so: LoadLibraryExW failed` --
a Linux name failing on Windows. That reads as *this machine has no
CUDA*, not *this list is stale*. It is why native CUDA appeared
unavailable on a box with a working RTX 4060, driver 591.55, CUDA 13.1,
and a complete CUDA 13 runtime present.

**Verified on that host.**
- Before, with only CUDA 13 on PATH: load fails at `cudart`.
- After, same PATH: gets past `cudart` and commits real device memory
--- `cuda_ep: VMM arena closing: committed 2097152 B of 68719476736 B
reserved`.
- Control: adding a CUDA **12** runtime directory to PATH produced that
same progression *before* the fix, which isolates the cause to the
missing name rather than to anything else in the environment.

**Guard, falsified.**
`every_cuda_major_the_ep_loads_is_also_resolvable_here` pins every major
the EP accepts. Dropping `cudart64_13.dll` reddens it: `missing
cudart64_13.dll; the CUDA EP loads this major but cuda_rt would not`.

**Follow-up, not in this PR.** The real fix is one table, not two
agreeing tables. Both crates already depend on
`onnx-genai-cuda-version-guard`, which is the natural home for it. Filed
separately; this PR unblocks CUDA 13 hosts now.

Note that a *second*, unrelated defect sits immediately behind this one
and is now reachable: native CUDA decode fails at `node 14 (op
'::Attention') reached execution without prepared SessionPersistent
workspace`. Filed separately.

Co-authored-by: justinchuby <223556219+Copilot@users.noreply.github.com>
…default-OFF) (#1167)

## QMoE gate/up SwiGLU expert-GEMV occupancy lever

Ports the dense #1139 gate/up occupancy win into the fused gate/up
SwiGLU expert GEMV (the MoE analog). New flag `ONNX_GENAI_QMOE_OCC`,
**default-OFF**.

### ncu evidence (spill-verified)
| | regs | theo occ | achieved | duration |
|---|---|---|---|---|
| default | 54 | 50% | 43.3% | 42.3us |
| `_occ (256,6)` | **40** (spill-free) | **75%** | **63.8%** | **37.8us
(−10.6%)** |

Register-granularity trap avoided: `(256,8)` → 32 reg/100% theoretical
but spills 1.22 MB → regression. Shipped block config 6.

`qmoe_linear` (fc2/down) already at 32 reg/100% (no lever); grouped-GEMM
M>1 path already dequants-once with TILE_M=8. Gate/up is the sole
remaining occupancy target.

### Parity — byte-identical
- `qmoe_occ_is_bit_identical`: 0 differing bits (f16/bf16/f32 ×
rows{1,4,6,8})
- Full `qmoe_gpu` suite: **32/32 pass** (re-verified post-rebase onto
current main)
- E2E DeepSeek-V2-Lite tokens identical OFF vs ON, both equal the merged
golden `[…207,16,24,1012,1712,5075…]`

### E2E
V2-Lite decode ~−1.0% median (56.6→57.2 tok/s, runs=10 warmups=2, 3
rounds). Small because gate/up ≈8% of V2-Lite decode (MLA/Attention ≈21%
dominates a different lane); bigger-intermediate MoE would convert more.

### Review
- 🧪 Rachael: 🟡 **approve-with-notes** — byte-identity + ncu occupancy +
spill-free config all independently verified. Note: flag selects `_occ`
for any QMoE QuantLayout when enabled (not literally
symmetric-int4-only); affine OCC smoke tests pass; default-OFF keeps it
non-blocking.

Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
)

## Why

Two problems with the master benchmark doc as it stands on `main`:

1. **§21.2/§21.3 are stale.** They describe the input-binding copy as an
open gap and explicitly decline to say what closing it would be worth.
#1146 closed it — `bind_inputs` on `rope_llama3_s128` went from ~601 µs
to 0.7 µs. Every ratio in the document that predates that merge is being
read against the wrong baseline, and nothing in the document says so.
2. **§23.7 frames the remaining losses as five independent kernel
problems.** They are mostly one problem, and it is not in any of those
kernels.

## What this adds

### §24 — Phase 6 (#1146)
What the borrowed `DeviceBuffer` does and the exact preconditions; the
sequence-storage choke point (`read_seq_element` *moves* a buffer into
cross-run `shared_buffers`, so guarding the slot was not enough); the
57-cell result (56/57 improved, 342/342 parity PASS); and precisely
which §21 claims it supersedes.

### §25 — Phase 7 (#1175)
The two independent `RotaryEmbedding` defects, with the measurement that
made them visible once #1146 left `exec_kernel.compute` at 96% of the
run:

* the fan-out was **net-negative at every thread count** (t=16: 195 µs
parallel vs 127 µs serial; 36 µs of pure Rayon bookkeeping to fan 128
tasks over a *one-worker* pool);
* the inner loop was scalar at ~1.4 cycles/element.

Plus why the new AVX2 arms deliberately avoid FMA (a tensor must not
change value because a different thread count split it in a different
place), and the 21-cell matrix: 21/21 improved, 0 regressed, 5 now beat
ORT.

### §26 — the root cause behind the `t=8/16` column
An isolated 524 288-element region on 16 workers, varying only the idle
gap between regions:

| gap | region cost |
|---|---|
| none | 67 µs |
| 20 µs | 226 µs |
| 500 µs | 236 µs |

A 20 µs gap — less than the serial glue between two graph nodes — costs
~160 µs. Rayon parks its workers; ORT's intra-op pool spins first, so
its workers are hot when the next op arrives and ours are not.
`rope_llama3_s128` at t=16 is 59 µs of kernel against ORT's ~30 µs, but
a 189 µs paired wall.

`decode_spmd` already implements the right policy (bounded
`KMP_BLOCKTIME` spin) and is wired only to the decode GEMM path. A
prototype routing RoPE through it takes t=16 from **79 µs to 25 µs
native-only, past ORT's ~30 µs**.

It is documented rather than merged, for two reasons that are both worth
recording:

**(a) A latent correctness hazard.** `SharedState::publish` writes the
job into a single `UnsafeCell` slot (`decode_spmd.rs:243`); `wait`'s doc
comment states "the dispatcher is a single, never-idle thread" as an
assumption, not an invariant. Two concurrent dispatchers overwrite each
other's job pointer and both sets of workers run whichever closure
landed last — **silently wrong tensors, not a crash**. Found
empirically: new fan-out unit tests reported 6 tasks run when 5 were
dispatched, and separately 0, depending on what else in the 1350-test
binary was dispatching; `--test-threads=1` passes. A per-module CAS
guard is *not* sufficient because the 11 existing dispatch sites do not
take it. Latent today (the decode engine dispatches inline on one
thread, and the plugin EP calls `disable_persistent_decode_pool()`),
live as soon as two sessions in one process both run decode GEMM — which
is exactly the concurrency the paged-KV work has to preserve.

**(b) Two spinning pools oversubscribe.** With the prototype on and an
ORT session alternating in-process, t=16 is bimodal — 0.030 ms on one
run, 0.441 ms on the next, same binary, same command.
`ONNX_GENAI_CPU_DECODE_THREADS=16` confines the process to CPUs 0–15,
which on this host is **8 physical cores** (siblings are adjacent). A
blocktime sweep over 500/200/50/20/5/0 µs gave ratios from 0.233 to 72.3
with no monotone trend. The section also records that the compact mask
is *correct* and must not be "fixed" to spread — one thread per core
across 0,2,…,30 straddles two CCXs and measured worse (0.133 vs 0.079
ms).

### §26.3 — the consequence for the work queue
Softmax, Transpose, Concat and the MoE cells should not be attacked as
separate kernel problems until (a) is fixed; their t=8/16 numbers are
mostly measuring thread wake-up, so per-kernel tuning against that
column will have to be redone. The stated exception is §11.5's
paged/appendable KV cache, which is a real algorithmic cost no scheduler
change touches.

## Scope

Documentation only — one file, append-only. No code, no test, no
behaviour change.

Co-authored-by: Leon <leon@squad.local>
Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
…UDA-graph workspace planning (#1181)

## What

Fixes the second (real) blocker to V2-Lite CUDA-graph capture: the
additive-mask-bias **query axis `_d1`** was an unresolved internal
symbol, so prepare-only workspace planning refused to reserve the
Attention workspace. Stacks on the topology-gated mask classifier
(#1171).

`v_model.Unsqueeze_18` (additive mask → every layer's `Attention[3]`)
has shape `['batch',1,'_d1','past+seq']`. Axis 3 (key/context) was
correctly bound to KV capacity; axis 2 `_d1` (query length) was a fresh
symbol shape-inference did not unify with `sequence_len`.

**Fix:** exact prepare-only recovery through the V2 mask cone
(`Shape/Sub/Slice/Unsqueeze/Where/Cast`), resolving `_d1` to the current
`sequence_len` — no max-seq guess, no over-reservation.

## Results (real 27-layer V2-Lite int4, H200)

- **eager 57.38 → CUDA-graph 101.76 tok/s (1.77x)**, `captures=4
replays=504 fallbacks=0`
- 128 generated token IDs **byte-identical** to eager golden
`…207,16,24,1012…`
- `deepseek_v2_lite_decode_lock` under `ONNX_GENAI_CUDA_GRAPH=1` →
**PASS** byte-identical to golden

## Safety

- Behind opt-in `ONNX_GENAI_CUDA_GRAPH` — no default flip
- Eager (`=0`) path untouched (V2-Lite eager 57.38 tok/s unchanged)
- Qwen3 dense CUDA planning reservation unchanged (`persistent=983040`)
- 181/181 session lib tests pass; planner tests 7/7

⚠️ Reviewer gate: Rachael reviews the planner change; Wallace runs the
hardware byte-identity + tok/s A/B before merge.

Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
Scribe consolidation of the graph-capture arc. Merges 7 inbox drops into
decisions.md (15636->19014B), marks the V2-Lite `_d1` fix #1181 MERGED
and the graph-capture top lever FULLY LANDED (classifier #1171 + `_d1`
#1181 = 1.79x V2-Lite MoE capture, byte-identical, opt-in). Squad state
only — no code changes.

Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
## Summary

- establish an Obsidian-compatible `wiki/` with stable English note
titles and paths
- add a Chinese beginner-oriented explanation of allocator, virtual
backing, shared mapping, governance, EP ordering, and process-level
ownership
- document note conventions and distinguish proposed architecture from
current implementation

The formal sources under `docs/memory/` remain authoritative; this wiki
is explanatory.

Co-authored-by: Copilot App <223556219+Copilot@users.noreply.github.com>
Copilot-Session: c80f8522-983c-47f7-8241-2155a823aabe
… it (#1188)

Agents on this repo have repeatedly concluded *no CUDA runtime on this
machine* and used that to accept GPU performance and correctness claims
as unverified. It is false, and it degraded review quality on every CUDA
PR.

The box has an RTX 4060 (driver 591.55, CUDA 13.1) and a complete CUDA
13 runtime under anaconda site-packages. It is simply not on PATH, and
because this repo loads CUDA libraries dynamically by DLL name, adding
two directories is enough. The recipe is in the decision note.

Two defects kept the belief alive and hid each other: the stale `cudart`
name list (#1178, fixed) and the unprepared `::Attention` workspace
(#1179, open). Together they made native CUDA genuinely unrunnable here,
and that unrunnability was reported upward as absence of hardware.

The lesson worth keeping past both bugs: distinguish *absent* from
*misconfigured* before building on a negative result. A negative that
excuses you from verification deserves more scrutiny than a positive,
not less.

Co-authored-by: justinchuby <223556219+Copilot@users.noreply.github.com>
…chers (#1184)

## The defect

`SharedState` (`decode_spmd.rs`) has **one** job slot and **one**
barrier. `wait`'s doc comment states the constraint —

> The dispatcher is a single, never-idle thread…

— as an *assumption*. Nothing enforced it. Two threads reaching
`dispatch` concurrently overwrite each other's job pointer and each
other's per-node pending counts.

Both failure modes are reachable, and neither is a clean crash:

* **Silently wrong tensors.** A worker set runs the *other* thread's
closure. Each shard writes a disjoint output region and the barrier
still balances, so the run returns success with plausible garbage in one
of the two sessions.
* **A spinning hang.** More often the pending counters never drain.
`wait` spins by design (it wants results the instant they land), so this
burns every core until the process is killed.

**Latent on `main`, not live.** The decode engine runs its forward
inline on one thread, and the plugin EP calls
`disable_persistent_decode_pool()` in `CreateEpFactories` before ORT can
reach a kernel. It becomes live the moment two sessions in one process
both run decode GEMM — which is exactly the concurrency the paged-KV
work has to preserve, so this is a prerequisite for that work rather
than a side quest.

## The fix

`dispatch` is the sole funnel for `publish`/`wait` (verified: 7 call
sites, all through it), so the gate goes in **one place** and covers
every existing and future caller. No API change.

A compare-exchange claims the slot. Losing the race runs the same shards
**inline on the calling thread**. That is not a fallback approximation —
the shard closures already partition their output by global worker
index, so running all of them on one thread is the *identical
computation*, only serial. A second concurrent session degrades in
throughput instead of corrupting the first.

Re-entrant dispatch (a shard closure that dispatches again) now runs
inline too, where it previously deadlocked against the barrier it was
already inside.

Claim and release live together in `DispatchClaim` — the only
constructor is `try_claim`, and `Drop` releases while unwinding. A
leaked claim would silently send every later dispatch inline forever, so
the panic path is covered by a test.

## The flag needs its own cache line

This was not obvious and I nearly shipped it wrong. The dispatcher
writes the flag twice per op while **every worker polls `shutdown` in
its wait loop** — so an unpadded flag sharing that line makes each claim
invalidate all 16 workers' copies.

Empty 16-worker dispatch, best of 7 × 20 000, three alternating rounds:

| variant | ns/dispatch |
|---|---|
| ungated (pre-fix) | 908 / 903 / 907 |
| **gated, padded** | **898 / 881 / 898** |
| gated, unpadded | 1452 / 1497 / 1432 |

Padded, the gate costs **nothing measurable**. Unpadded it costs
**~60%**, which at ~400 barriers/token is ~230 µs/token of pure
coherency traffic. My first measurement of this was a cold run (1474 vs
1148) that got the *magnitude* wrong in both directions; the table above
is three alternating rounds.

## Falsification

Deleting the gate and re-running the new tests reproduces the original
defect directly.
`two_threads_dispatching_at_once_each_get_their_own_complete_partition`
**hangs** (killed at 10 min) rather than failing an assertion — which is
precisely why this hazard had not been noticed: the failure is a hang or
silent corruption, never a test-visible panic.

Five tests: claim exclusivity, release-on-unwind, complete partition per
dispatch, the two-thread race (64 dispatches × 2 threads, every shard
exactly 64), and re-entrancy.

## Validation

* `cargo test -p onnx-runtime-ep-cpu --lib` → **1333 passed, 0 failed**,
16 ignored (1328 + 5 new)
* `cargo test -p onnx-runtime-ep-cpu-plugin --test plugin_ort_e2e` →
**53 passed, 0 failed** (real ORT 1.27, `disable_cpu_ep_fallback=1`)
* `cargo fmt --all --check`, `cargo clippy -p onnx-runtime-ep-cpu
--all-targets` → clean

## Why now

§26 of the benchmark doc (merged in #1182) establishes that Rayon
worker-park latency — not kernel arithmetic — is the root cause behind
every "wins at t=1, loses at t=8/16" row in §23.7. A prototype routing
RoPE's fan-out through this pool takes `rope_llama3_s128` at t=16 from
**79 µs to 25 µs, past ORT's ~30 µs**. That routing cannot land while
the pool has one dispatcher by assumption. This is step one.

---

## Independent adversarial review (Opus)

Verdict: **MERGE, no blockers.** The reviewer did not take the claim on
trust — they built a faithful standalone replica of the ungated protocol
outside the repo and ran it 12×, observing **both** claimed failure
modes: counts ≠ 64 (corruption) and `wait` spinning forever (hang). They
further confirmed the Acquire/Release pairing is *load-bearing* by
weakening it to Relaxed and reproducing the same race.

Verified across the ten questions put to them:

* **Sole funnel — yes.** `publish`/`wait` are module-private and reached
only from `dispatch`; `SharedState` is private; all public
`dispatch_*`/`place_rows` methods funnel through it. `shutdown()` bumps
senses but never publishes or touches `node_pending`.
* **Ordering — correct.** The worker copies `Job` into a local *before*
running and decrements `node_pending` (AcqRel) *after*, never touching
`shared.job` post-decrement. `wait`'s Acquire-load-of-0 happens-after
every worker read via the RMW release chain; `drop(claim)` Release →
next `try_claim` Acquire → next `publish` orders the next non-atomic
slot write after all prior workers' reads. No previous-dispatch worker
can read `job` during a new overwrite.
* **`dispatch_inline` equivalence — yes, all 7 call sites.** Five index
`segments[global_index]` from
`worker_row_segments`/`output_segments_aligned`, which produce exactly
`total_workers` disjoint entries covering the output — serial execution
is bit-identical. The two work-stealing closures drain an atomic cursor,
so index 0 drains all and the rest break: correct serially. No closure
waits on another shard, uses thread-locals, or needs a pinned thread for
correctness. `place_rows` loses NUMA first-touch locality inline —
performance only, acceptable for a rare fallback.
* **mlas bypass — safe.** `WorkStealingThreadPool::parallel_for` holds
its own `dispatch_lock: Mutex<()>`, serializing concurrent callers
independently.
* **Padding — sound.** `Padded<AtomicBool>` empirically size 128 / align
128, so it gets its own block regardless of Rust field reordering and
cannot share a line with `shutdown`/`ready`/`poisoned_worker`.
* **Poison, shutdown, re-entrancy, ABA, liveness** — all preserved. The
CAS is unfair, so a thread could persistently lose, but inline makes
full correct progress, so there is no starvation-into-blocking and no
convoy. The claim is held only across a bounded spin barrier, never
unbounded blocking.
* **`pools()` is `Some(16)` under `cargo test`** — independently
confirmed, so the pool-backed tests are not vacuous.

Two cosmetic findings, both applied in the follow-up commit: `wait`'s
doc still stated the single-dispatcher property as an assumption rather
than an enforced invariant, and the test module did not say which tests
are the actual falsifiers (only the two contended ones fail if the gate
is deleted).

## Final validation, on `main` merged into this branch

* `cargo test -p onnx-runtime-ep-cpu --lib` → **1333 passed, 0 failed**,
16 ignored
* `cargo test -p onnx-runtime-ep-cpu-plugin --test plugin_ort_e2e` →
**53 passed, 0 failed** (real ORT 1.27, `disable_cpu_ep_fallback=1`)
* `cargo test -p onnx-runtime-session --lib` → **181 passed, 0 failed**
* `cargo fmt --all --check`, `cargo clippy -p onnx-runtime-ep-cpu
--all-targets` → clean

---------

Co-authored-by: Leon <leon@squad.local>
Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
## Summary
- re-prepare native CUDA decode workspace when KV/mask capacity grows on
Engine single-token paths
- keeps Attention governed workspace sized to the current physical
context bucket before eager/capture resumes
- adds an ignored real V2-Lite Engine long-context eager-vs-capture
regression

## Validation
- Reproduced pre-fix V2-Lite Engine eager failure: node 38 Attention
required 33288 B, prepared 16904 B
- `profile_native --steady --tokens 340` V2-Lite eager: passes, 47.32
tok/s, golden prefix intact
- `ONNX_GENAI_CUDA_GRAPH=1 profile_native --steady --tokens 340`:
passes, 89.69 tok/s, token IDs byte-identical to eager, captures=2
replays=336 fallbacks=0
- `cargo test -p onnx-genai-engine --release --features
native-backend,cuda --test deepseek_v2_lite_decode_lock
deepseek_v2_lite_engine_long_context_workspace_survives_capacity_growth
-- --ignored --nocapture`
- `cargo test -p onnx-runtime-session --lib`
- `CUDA_VISIBLE_DEVICES=0
ONNX_GENAI_QWEN3_0_6B_FOUNDRY_DIR=/home/justinchu/.foundry/cache/models/Microsoft/qwen3-0.6b-generic-cpu-4/v4
cargo test -p onnx-genai-engine --features native-backend,cuda --test
qwen3_0_6b_foundry_native_cuda_lock -- --ignored --nocapture`

Rachael/Wallace gate requested by coordinator.

Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
…nt 66-cell grid (#1191)

Two gaps in the benchmark doc after #1184 merged.

## §26.2.1 — a negative result worth recording

#1184 closed blocker **(a)**: the SPMD pool is now genuinely
multi-dispatcher. So I re-wired `RotaryEmbedding`'s fan-out through it
and re-ran the paired grid, expecting the §26.2 prototype win (79 → 25
µs native-only, past ORT) to survive.

It did not. §26.2(b)'s bistability is intact, and the grid lands on
**both modes at once**:

| model | t | rayon (base) | hot pool |
|---|---|---|---|
| rope_llama3_s128 | 16 | 9.457 `[6.4–23.1]` | **72.704** `[20.7–86.6]`
|
| rope_gptj_il_s128 | 16 | 6.247 | **55.751** `[29.0–85.7]` |
| rope_gptj_il_s512 | 16 | 5.445 | **1.232** `[1.08–35.6]` |
| rope_llama3_b8_s1 | 16 | 0.807 | **0.622** |
| rope_llama3_s512 | 16 | 30.771 | **18.385** |

Dispersion bands wider than the medians, in one driver invocation. The
three improving rows are **not** a partial win — they are the same
bistable system landing on its other mode. Nothing here is publishable
as a ratio, so the wiring is reverted and RoPE stays on Rayon.

**This settles the ordering: (b) is the binding constraint, not (a).**
Capping *spinning* workers at one per physical core inside the compact
mask — which needs SMT sibling discovery `decode_affinity` does not
currently do — has to land before the hot-pool routing can be evaluated
at all, let alone merged.

Recording it so the next person doesn't repeat the experiment.

## §27 — where the CPU EP actually stands

The document had no single place stating the current position. Full
transform + KV grid on `main` at `f2e6ad97d`, paired interleaved vs real
ORT 1.27, 5 trials × 60 runs × 15 warmups: **8 of 66 cells win, 58 lose,
parity 330/330 PASS.**

Read by **column**, which is the whole point of §26:

| threads | wins | median ratio |
|---|---|---|
| 1 | 8 / 22 | ~1.02 |
| 8 | 0 / 22 | ~3.0 |
| 16 | 1 / 22 | ~3.6 |

At t=1 the kernels are at parity — median 1.02, eight outright wins,
worst cell 1.706. That column is a fair measure of the arithmetic and
says it is broadly competitive. At t=8/16 the median is 3–3.6× with one
win between them. **The kernels did not get worse; the scheduling did.**

Two caveats are stated in the section rather than buried: the host is
shared, so absolute numbers drift between runs (`rope_llama3_s128` at
t=16 reads 17.557 in this grid and 9.457 in the §26.2.1 grid taken later
the same session) and only ratios from a single invocation compare —
which is why the arms interleave. And sub-100 µs cells are
overhead-dominated, which is why `rope_*_s1` sits at a flat ~1.3 at
every thread count instead of scaling.

Closes with a ranked blocker list: SMT-aware spin capping (blocks 58 of
the 58 losing cells), then paged KV (§11.5, the only structural item
actionable *right now* since it is independent of the scheduler), fused
transpose-with-consumer, grouped MoE GEMM, fused-attention session A/B,
and the output arena.

## Scope

Documentation only — one file, append-only. No code, no test, no
behaviour change.

Co-authored-by: Leon <leon@squad.local>
Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
Scribe consolidation: merges Leon #1189 Engine long-context fix,
Rachael's review gate, and Wallace's Native-vs-ORT DeepSeek-V2-Lite
measurement into decisions.md (archived older entries to 2026-08.md,
19014→6523 B). Updates now.md standing tables. Squad-state only.

Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
justinchuby and others added 3 commits August 18, 2026 04:09
…ion, not just assignment (#1169)

## What

Selecting our CPU EP means full local ownership of every node it claims.
The
performance-based decline (`assignment_policy.rs`) was already deleted,
but the
*mechanism* it plugged into was still standing — a deferral could have
been
reintroduced by overriding a trait default. This removes the mechanism,
and
then proves the property that replaced it.

Deleted:

| what | where |
|---|---|
| `ClaimPreference::{Claim, DeferToHost}` |
`onnx-runtime-ep-api/src/kernel.rs` |
| `ExecutionProvider::claim_preference` / `claim_preference_node` |
`onnx-runtime-ep-api/src/provider.rs` |
| the `apply_claim_preference` gate in `GetCapability` |
`onnx-runtime-ep-plugin/src/ep.rs` |
| `NativeEp::host_fallback_available` / `with_host_fallback` |
`onnx-runtime-ep-plugin/src/ep.rs` |
| `factory::host_cpu_fallback_disabled` (69 lines) |
`onnx-runtime-ep-plugin/src/factory.rs` |
| the CPU EP's `claim_preference_node` override |
`onnx-runtime-ep-cpu/src/provider.rs` |

`supports_op` — a *capability* answer, not a preference — is now the
only
claim-time question, and there is no default left to override. The
`session.disable_cpu_ep_fallback` special case went with it: the gate it
disabled no longer exists, so the option no longer changes our
behaviour.

## Why a counter, and not just numerics

ORT's node→EP attribution is a session-build fact. Output equality
against ORT
**cannot** distinguish "our kernel ran" from "ORT's kernel ran", because
agreeing with ORT is exactly what a correct kernel does. Assignment and
execution therefore get two separate counters:

- `nxrt_ep_compiled_node_count` — assignment/compile time (already
existed)
- `nxrt_ep_executed_node_count` — **new**, incremented after each node
kernel
  this EP actually executes, on both the routed multi-node path and the
  single-node path in `compute.rs`

`every_assigned_node_is_also_executed_by_this_ep` runs each assignment
fixture
with `session.disable_cpu_ep_fallback=1` and asserts that one `Run`
advances the
execution counter by exactly the number of nodes ORT assigned to
`cpu_ep`. It
reads the counter by `dlopen`ing the **same absolute cdylib path** ORT
registered, so it observes the same statics rather than a second
mapping.

The assertion is kept non-vacuous by a baseline: the identical graph and
inputs,
on a session where the EP is never appended, must leave the counter at
zero.

## Evidence

`NXRT_REQUIRE_ORT_TESTS=1` (skips become hard failures), ORT 1.27.0
prebuilt,
release build, AMD EPYC 9V74 (AVX2/FMA/F16C, no AVX-512).

```
$ cargo test --release -p onnx-runtime-ep-cpu-plugin --no-fail-fast
test result: ok. 54 passed; 0 failed; 1 ignored   (plugin_ort_e2e)
test result: ok.  9 passed; 0 failed              (plugin_export_abi)
test result: ok.  7 passed; 0 failed              (shape_inference_coverage)
... all suites ok, 0 failed

$ cargo test --release -p onnx-runtime-ep-plugin -p onnx-runtime-ep-api --no-fail-fast
test result: ok. 246 passed; 0 failed
test result: ok.  56 passed; 0 failed
test result: ok.  12 passed; 0 failed
... 0 failed
```

New test, over the 38 entries of `ASSIGNMENT_FIXTURES`:

| outcome | count |
|---|---|
| assigned to `cpu_ep` **and** executed here, counter matched exactly |
**32** |
| of those, also proved the zero baseline (no EP appended ⇒ counter
stays 0) | 29 |
| ORT alone cannot run the graph, so no baseline exists | 3 |
| needs semantic inputs noise cannot synthesise | 6 |

The 3 without a baseline are real ORT gaps, reported rather than
asserted: bf16
`Add` has no ORT CPU kernel at all, and ORT's own `QMoE` CPU kernel
rejects the
fixture's `fc2_experts_scales` shape `{4,32,2}` (it wants `{4,32,1}`).
The 6
needing semantic inputs — dynamic dims, position ids, sequence lengths —
stay
covered by `rope_and_gqa_execute_on_our_ep_and_match_ort_numerics`,
which drives
them with real values and checks numerics.

Pre-existing conformance tests still pass unchanged:
`no_supported_node_is_ever_left_to_the_ort_cpu_ep` and
`every_fixture_loads_with_cpu_fallback_disabled`.

## Not in scope

No kernel, no scheduling and no numerics change — `cargo test` timings
and the
activation/MatMul ratio tables are untouched by this PR. The performance
work
these docs list continues separately.

Docs updated: `CPU_ACTIVATION_GAPS.md` (Filter 1 is now
*unrepresentable*, plus
the assignment-is-not-execution note) and `CPU_MATMUL_ASSIGNMENT.md`.

Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
Wallace scoped GLM-4-9B int4 graph capture: engages cleanly TODAY with
zero code change (V2-Lite capture stack #1171+#1181+#1189 already covers
GLM's dense GQA topology). captures=3 replays=185 fallbacks=0,
byte-identical 0/256 tokens, capture 211.74 vs eager 128.82 = 1.64x.
Updates now.md standing table + GLM support row + decisions.md.
Squad-state only.

Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
…oss the activation family) (#1174)

## What

Pin `codegen-units = 1` for `onnx-runtime-ep-cpu` in the workspace
release
profile.

The AVX2 elementwise kernels only vectorise when the compiler can see
the
dispatcher, the chunk loop and the per-vector body together. The default
sixteen codegen units split that chain and the loops come out
**scalar**.

#1136 found one *cause* of the repartition — instantiating the generic
`run_chunked` from another module — and fixed it by making `run_chunked`
private. That closed one door. The compiler was still free to split the
crate
on its own, and it does.

## Why this is a regression, not a tuning knob

`main` has been measurably slower than the ratios this repo publishes in
`docs/performance/CPU_ACTIVATION_GAPS.md`, with no source change to
account for
it. `Tanh` at 1 Mi is published at 0.82 of ORT; measured on `main` today
it is
**0.41**. With the pin it is 0.70. The published numbers came from a
build whose
partition happened to be favourable and were not reachable on `main` at
all.

## Evidence — kernel level

`activation_bench`, two binaries built from **identical source**
(default vs
`CARGO_PROFILE_RELEASE_CODEGEN_UNITS=1`), run interleaved 5 rounds,
medians,
`taskset -c 8-15`. 105 cases (7 ops × 5 shapes × 3 dtypes).

| case | 16 CGUs | 1 CGU | ratio |
|---|---|---|---|
| `Sqrt` f32 / 4096 | 0.639 ns/elem | 0.241 | **2.65x** |
| `Sqrt` f32 / 3072 | 0.642 | 0.244 | 2.63x |
| `Tanh` f32 / 3072 | 0.731 | 0.381 | 1.92x |
| `Tanh` f32 / 2 Mi | 0.189 | 0.099 | 1.90x |
| `Sigmoid` f32 / 3072 | 0.735 | 0.389 | 1.89x |
| `Sqrt` f16 / 4096 | 0.915 | 0.499 | 1.83x |
| `Sigmoid` f32 / 2 Mi | 0.196 | 0.108 | 1.81x |
| `QuickGelu` f32 / **16** (control) | 8.556 | 8.657 | 0.99x |
| `Tanh` bf16 / **16** (control) | 21.149 | 21.305 | 0.99x |

**77 of 105 cases move by more than 1.15x.** The ones that do not are
the
16-element shapes — dispatch-bound, so a de-vectorized loop cannot show
up in
them, and nothing else should. That split is the signature of
de-vectorization,
not of noise.

Reproduced independently through the per-package override actually being
landed
here (`[profile.release.package.onnx-runtime-ep-cpu]`), same 77/105 and
same
2.65x worst case.

## Evidence — session level, against ORT

Through the ORT session API, our EP vs ORT's own CPU EP in one process,
interleaved iteration by iteration, `intra_op = 1` on **both** sides
(the harness refuses a half-pinned comparison), 31 iterations after 3
warmups,
p50/p90 of whole-`Run`, `taskset -c 8-15`.

Columns are `ours_ms / ort_ms`, so **lower is better and below 1.00
means we
win**.

### float32, 1 Mi, one thread

| case | before p50 | after p50 | before p90 | after p90 |
|---|---|---|---|---|
| `Tanh` | 2.449 | **1.433** | 3.135 | 1.381 |
| `Sigmoid` | 2.342 | **1.383** | 2.286 | 1.372 |
| `Erf` | 2.058 | **1.565** | 2.049 | 1.573 |
| `Gelu` (tanh) | 2.020 | **1.397** | 2.017 | 1.379 |
| `Gelu` (exact) | 1.951 | **1.507** | 1.934 | 1.497 |
| `Exp` | 2.408 | **1.285** | 2.333 | 1.271 |
| `FastGelu` | 2.003 | **1.402** | 1.996 | 1.394 |
| `QuickGelu` | 1.450 | **0.953** | 1.428 | 0.945 |
| `Sqrt` | 1.530 | **0.718** | 1.508 | 0.734 |
| `Relu` (control) | 1.038 | 1.033 | 1.060 | 1.058 |

`Sqrt` and `QuickGelu` cross from loss to **win**. `Relu` is the
control: it is
memory-bound at 1 Mi, so the codegen partition cannot move it, and it
does not.

### float32 4 Ki and float16 1 Mi, one thread

| case | before p50 | after p50 |
|---|---|---|
| `Tanh` f32 4 Ki | 2.058 | **1.523** |
| `Sigmoid` f32 4 Ki | 2.051 | **1.503** |
| `Erf` f32 4 Ki | 1.905 | **1.598** |
| `Sqrt` f32 4 Ki | 1.616 | **1.150** |
| `Tanh` f16 1 Mi | 2.060 | **1.359** |
| `Exp` f16 1 Mi | 2.015 | **1.332** |

Reproduce:

```sh
NXRT_MM_BENCH=1 NXRT_MM_BENCH_THREADS=1 ONNX_GENAI_MLAS_THREADPOOL_THREADS=1 \
NXRT_MM_BENCH_CASE=f32_1m NXRT_MM_BENCH_ITERS=31 taskset -c 8-15 \
cargo test --release -p onnx-runtime-ep-cpu-plugin --test plugin_ort_e2e \
  plugin_path_ab -- --nocapture --ignored
```

## Also in this PR

- **`unary_bench_cases()`** — extends the session A/B harness, which
until now
covered only the matmul family, to the elementwise grid the activation
doc is
written against. Same harness, same refusal to report a half-pinned
ratio, so
a number from it is directly comparable to a matmul number. `#[ignore]`d
like
  the rest of the harness, so no CI cost.
- **`codegen_units_are_pinned`** — reads the setting back out of the
workspace
manifest. A build setting is exactly what a rebase drops silently: no
test
fails, no path changes, everything just gets slower. That is how this
class of
  regression reached `main` in the first place, so it gets a test.
- Doc update recording the before/after grid and the reproduce line.

## Numerics

Unchanged — this is a compiler partitioning setting, not a code change.
Full
suites green: `onnx-runtime-ep-cpu` 1324 passed / 0 failed, and every
`onnx-runtime-ep-cpu-plugin` suite passed with
`NXRT_REQUIRE_ORT_TESTS=1`
(so ORT-dependent tests are hard failures rather than skips).

## Cost

`onnx-runtime-ep-cpu` builds in ~65 s instead of ~19 s. Scoped to this
one
package, so nothing else in the workspace changes.

Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
#1158) (#1200)

Same content as #1158, reopened on a branch based directly on `main`.

#1158 was reviewed, rebased and validated, but could not be merged: it
is the base of #1161, and GitHub refuses both `gh pr merge` and a base
change for stacked pull requests, directing callers to an asynchronous
merge endpoint this repository does not expose (404). Rather than leave
two reviewed PRs stuck behind an API we cannot reach, the commit is
cherry-picked onto `main` here and #1158 is closed. #1161 is then free
to rebase onto `main` normally.

The rebase produced four conflicts in `pipeline/mod.rs`, all the same
shape as the earlier ones in this family: this branch adds
`max_sequence_length` and `effective_max_context` while `main` has since
added `generation_defaults`. They are independent additions in the same
struct, constructor and accessor block, so both sides are kept; the
accessor hunk needed a closing brace reinserted because the two sides
shared one.

Validation on the rebased tree: `pipeline_executor` 2 passed, and the
engine lib suite 508 passed / 0 failed / 1 ignored, matching the
baseline on `main`.

Copilot-Session: 0190e2eb-abe4-451f-b36d-44a035a99b7e
justinchuby and others added 2 commits August 17, 2026 21:27
Apply all metadata EOS token IDs in pipeline generation so ATEM models stop at end-of-turn. Parse tokenizer response templates to preserve and format generated channel markers while keeping only answer content in conversation history.

Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

Copilot-Session: 0190e2eb-abe4-451f-b36d-44a035a99b7e
`load_response_config` succeeds from `tokenizer_config.json` alone, but the
marker strings are only turned into ids later, against the tokenizer. When a
template names a marker the tokenizer does not encode as a single special
token, `message_token_id` stayed `None` -- and `ResponseStream` starts in
`Header`, whose only exit compares `message_token_id == Some(token.token_id)`.
`None` never equals `Some(_)`, so the parser was trapped in a state where every
output carries `include_in_reply: false`. The caller got an empty string and no
error: a good generation reported as silence.

Withdraw structured mode entirely instead of degrading in place. Raw output
shows the markers as literal text -- a poorer presentation, but an honest one.
Silence is neither.

Make it a mechanism rather than a patch: `bind_response_tokens` is the single
guarded entry point, and `bind_token_ids` is now module-private so no future
caller can bind without the check. Only `message_token_id` is load-bearing;
withdrawing on the optional markers would drop structured mode for templates
that parse fine.

Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
Copilot-Session: d60eb808-7cc6-4abc-b48d-2a6dd3841624
@justinchuby
justinchuby force-pushed the split/cli-structured-responses branch from 7c7dc0c to 7607fa1 Compare August 18, 2026 05:38
@justinchuby

Copy link
Copy Markdown
Owner Author

Rebased onto main and merged after independent verification.

Rebase. This branch was stacked on #1158, which is now in main as #1200; the duplicated commit was skipped and the remaining pipeline/mod.rs conflict was two independent free functions folded into one region -- pipeline_metadata_prefill_chunk_size and pipeline_metadata_generation_defaults on one side, pipeline_metadata_eos_token_ids on the other. All three reconstructed in full, no side dropped.

A silent-empty-reply bug found while reviewing, fixed in 7607fa11. load_response_config succeeds from tokenizer_config.json alone, but the marker strings are resolved to ids later, against the tokenizer. If a template names a marker the tokenizer does not encode as a single special token, message_token_id stays None -- and ResponseStream starts in Header, whose only exit is message_token_id == Some(token.token_id). None never equals Some(_), so the parser is trapped in a state where every output carries include_in_reply: false. The caller gets an empty string and no error.

The fix withdraws structured mode entirely rather than degrading in place: raw output shows the markers as literal text, which is a poorer presentation but an honest one. Silence is neither. It is a mechanism, not a patch -- bind_response_tokens is the single guarded entry point and bind_token_ids is now module-private, so no future caller can bind without the check. Only message_token_id is treated as load-bearing; withdrawing on the optional markers would drop structured mode for templates that parse fine.

Validation. onnx-genai-cli lib: 121 passed / 0 failed. Engine lib: 508 / 0 / 1, matching main's baseline. Both new tests falsified -- forcing is_parseable to true turns them RED on the exact assertion they claim to guard. The two repl_e2e attachment failures (missing required pipeline input 'encoder.pixel_values') reproduce identically on unmodified main and are unrelated to this branch.

justinchuby added a commit that referenced this pull request Aug 18, 2026
…as silence (#1214)

Unstacked replacement for #1161, which GitHub will not let us merge: it
was stacked on #1158, and the stack metadata survives even though that
PR is now closed and its content is in main via #1200. `gh pr merge`,
`gh pr edit --base`, and both async merge endpoints all refuse or 404.
Same branch content, off `main`, no stack.

## What it does

Formats structured model responses in the CLI, and fixes a
silent-empty-reply bug found while reviewing it.

`load_response_config` succeeds from `tokenizer_config.json` alone, but
the marker strings are only resolved to token ids later, against the
tokenizer. If a template names a marker the tokenizer does not encode as
a single special token, `message_token_id` stays `None` -- and
`ResponseStream` starts in `Header`, whose only exit is
`message_token_id == Some(token.token_id)`. `None` never equals
`Some(_)`, so the parser is trapped in a state where every output
carries `include_in_reply: false`. **The caller gets an empty string and
no error: a good generation reported as silence.**

The fix withdraws structured mode entirely rather than degrading in
place. Raw output shows the markers as literal text -- a poorer
presentation, but an honest one. Silence is neither.

It is a mechanism, not a patch: `bind_response_tokens` is the single
guarded entry point, and `bind_token_ids` is now module-private so no
future caller can bind without the check. Only `message_token_id` is
treated as load-bearing; withdrawing on the optional markers would drop
structured mode for templates that parse fine.

## Validation

- `onnx-genai-cli` lib: **121 passed / 0 failed**
- engine lib: **508 / 0 / 1**, matching main's baseline
- Both new tests falsified: forcing `is_parseable` to `true` turns them
RED on the exact assertion they claim to guard
- The two `repl_e2e` attachment failures (`missing required pipeline
input 'encoder.pixel_values'`) reproduce identically on unmodified main
and are unrelated to this branch

Supersedes #1161.

---------

Co-authored-by: justinchuby <223556219+Copilot@users.noreply.github.com>
Copilot-Session: 0190e2eb-abe4-451f-b36d-44a035a99b7e
Copilot-Session: d60eb808-7cc6-4abc-b48d-2a6dd3841624
@justinchuby

Copy link
Copy Markdown
Owner Author

Superseded by #1214, which is the same content off main with the silent-empty-reply fix included. Closing only because GitHub will not merge this one: the stack metadata from #1158 survives even though that PR is closed and its content reached main via #1200, and gh pr merge, gh pr edit --base, and both async merge endpoints all refuse or 404. Nothing was dropped.

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants