Sync with Microsoft ONNX Runtime - 02112026 by ai-fw-intg · Pull Request #1379 · intel/onnxruntime

ai-fw-intg · 2026-11-01T20:39:41Z

Automated daily backmerge from ORT main to ovep-develop. No conflicts detected. Do NOT squash or rebase - use merge commit only.

…e and 4-lane warp-shuffle decode epilogue (microsoft#33738) ## Description Optimizes `HyperConnectionLowRankMix` for both multi-row prefill and small-M decode on CUDA: 1. **Prefill (`FusedPreMixInjectPostGateB4H2560Kernel<T>`)**: Fuses the pre-mix projection (`Norm(streams) * sigmoid_gate * (2 * alpha_pre)`), the `alpha_inj * streams + y` injection, and the `2 * sigmoid(post_logits) * post_scale` post-gate into a single 320-thread register-only kernel for the 4-branch, `hidden = 2560` shape (`num_rows >= 8, rank == 320`), eliminating two intermediate `(num_rows, 4, 2560)` global-memory round-trips and reducing prefill HyperConnection mix latency by **25%** (`41.7 -> 31.4 us/launch` at `num_rows = 8192` on H200). 2. **Decode (`HcDownKernel` and `HcUpPreKernel<T, 2, 4, 320>`)**: Adds a 4-lane warp-shuffle epilogue (`kFastEp4x320`) that eliminates the `sg` shared-memory staging and third CTA barrier in `HcUpPreKernel`, prefetches `norm_w` / `streams` before `griddepcontrol.wait`, caches `w_reg8` / `w_reg9` in `HcDownKernel`, and unrolls `p.splits == 8`, reducing combined `HcDown + HcUpPre` critical-path time in Qwen3.8-Flash-Next decode by **94 us/cycle**. ## Summary of Changes | File | Change | |------|--------| | `onnxruntime/contrib_ops/cuda/bert/hyper_connection_low_rank_mix_impl.cu` | Add `FusedPreMixInjectPostGateB4H2560Kernel<T>` for prefill and `kFastEp4x320` 4-lane warp-shuffle epilogue + weight/stream prefetching in `HcDownKernel` and `HcUpPreKernel<T, 2, 4, 320>` for decode | | `onnxruntime/test/contrib_ops/hyper_connection_low_rank_mix_op_test.cc` | Add `(num_rows=11, num_branches=4, hidden=2560, rank=320)` FP16/BF16 coverage to `UnfusedManyRows` | ## Testing - `onnxruntime_provider_test --gtest_filter='*HyperConnection*'`: **30 / 30 passed** (including new `(11, 4, 2560, 320)` prefill test cases and `FusedDecodeShapeRank320`). - End-to-end `engine_behaviours.py` on Qwen3.8-Flash-Next RC5 (`all_int4`, `top16`, `top08`): **27 / 27 PASS** on each package. ## Checklist - [x] Tests added/updated - [x] No breaking changes - [x] `lintrunner` clean --------- Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

…soft#33404) ### Description  Refines the extensible runtime fusion design around EP-registered, ordered patterns and the existing capability path: - EPs own pattern registration and priority; ORT selects compatible candidates in that order, subject to semantic, ownership, kernel, and eventual memory-feasibility checks. - Reuse `GetCapability`, `ComputeCapability`, `nodes_to_optimize`, and `optimization_func` through validated adapters. Distinguish existing application hooks from the additional recipe-description, ordering, capability-update, and external-registration support still needed. - Separate candidate description, selection, application, and reconciliation. Baseline ordered selection comes first; target memory estimation can connect to the same recipes before final budgeted selection, without requiring the general placement planner. - Replace the parallel post-partition hook and proposed `OrtRF*` function table with logical recipe/lifetime contracts. Preserve optional independent package delivery; defer exact public ABI additions until the capability integration is established. - Treat gramalingam#16 as an optional matcher implementation to evaluate with small Qwen-derived fixtures and the model demo, not a system-design prerequisite. - Update the qualification plan, approval ledger, and milestones consistently; retain CUDA mapping, state, numerical, and resource requirements. Design-only; no runtime implementation, build, GPU allocation, or model qualification is included. Public interfaces and implementation contracts remain subject to maintainer approval. Validation: `git diff --check` passed. Documentation fences and stale hook/API references were checked; no runtime tests were run. ### Motivation and Context  Keep canonical ONNX artifacts unchanged while allowing optional matcher/adapter updates to reuse existing kernels. Align fusion with capability-based EP ownership and the recipe/planning direction in microsoft#33372 instead of introducing a separate selection pipeline. Qwen3.8-27B remains the model-independent mechanism's pilot, not its API. Fixes microsoft#33403. --------- Signed-off-by: copilot-swe-agent[bot] <198982749+Copilot@users.noreply.github.com> Co-authored-by: copilot-swe-agent[bot] <198982749+Copilot@users.noreply.github.com> Co-authored-by: titaiwangms <18010845+titaiwangms@users.noreply.github.com> Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> Copilot-Session: 38c05fba-777a-447c-af24-f930f13e4021

## Description Adds a `com.microsoft` `SwiGLU` contrib operator with a CUDA kernel, plus a CUDA graph fusion that rewrites the decomposed SiLU-gated activation that exporters emit today into the new operator. `SwiGLU` exposes the clamped SwiGLU activation that `MoE`/`QMoE` already apply internally, for dense feed-forward and shared-expert paths that do not run through a grouped expert GEMM. > This PR previously also added `MoERouter`. That operator is now in its own PR: microsoft#32935. ## Summary of Changes ### `SwiGLU` operator | File | Change | |------|--------| | `onnxruntime/core/graph/contrib_ops/contrib_defs.cc`, `ms_opset.h` | Schema, shape inference and registration. | | `onnxruntime/contrib_ops/cuda/math/swiglu.{h,cc}` | CUDA `OpKernel`: validates shapes and handles the packed single-input form. | | `onnxruntime/contrib_ops/cuda/math/swiglu_impl.{h,cu}` | Fused clamp + gate kernel. | | `onnxruntime/contrib_ops/cuda/cuda_contrib_kernels.cc` | Registration for float, float16 and bfloat16. | ``` G = clamp(gate, max=limit) L = clamp(up, min=-limit, max=limit) output = G * Sigmoid(activation_alpha * G) * (L + activation_beta) ``` - Arithmetic is done in float regardless of `T`. A `limit` of zero or less disables both clamps. - The result is rounded to `T` once, after the multiply. The optional INT attribute `round_activation` (default 0) additionally rounds `G * Sigmoid(activation_alpha * G)` to `T` before the multiply by `(L + activation_beta)`, as a `QuickGelu` followed by a `Mul` in `T` would. - Uses the same `limit` / `activation_alpha` / `activation_beta` contract as `MoE`/`QMoE`. - `up` is optional. When omitted, `gate` holds both halves of one `[.., 2 * inter]` projection (gate first, then up), so a fused sibling GEMM needs no `Split`. - With `up` supplied and default `limit`/`activation_beta`, it matches ONNX `SwiGLU` (opset 28) with `alpha = activation_alpha`. ### `SwiGluFusion` (CUDA EP, level 2) | File | Change | |------|--------| | `onnxruntime/core/optimizer/swiglu_fusion.{h,cc}` | Rewrites `QuickGelu(gate, alpha) * up` to `SwiGLU(gate, up, activation_alpha=alpha, round_activation=1)`. | | `onnxruntime/core/optimizer/graph_transformer_utils.cc` | Registers it for CUDA right after `QuickGeluFusion`. | `QuickGeluFusion` already turns `gate * Sigmoid(gate)` into `QuickGelu`, so existing models with a decomposed SiLU-gated MLP now run the fused kernel without being re-exported. The fusion only fires when `QuickGelu` has a single consumer, the element type is float, float16 or bfloat16, and `gate` and `up` provably share one shape (`Mul` broadcasts but `SwiGLU` does not). `round_activation=1` keeps `QuickGelu`'s rounding to `T` before the multiply, so the fused graph is bit-identical to the unfused `QuickGelu` + `Mul` for float, float16 and bfloat16, including the negative tail where the activation underflows in `T`. ### Shared helper - `_Sigmoid(float)` added to `onnxruntime/core/providers/cuda/cu_inc/common.cuh` next to `_Exp`/`_Log`. It uses the same numerically stable form as `OP_Sigmoid`. ## Testing - `onnxruntime/test/python/transformers/test_swiglu.py`: PyTorch reference over float32/float16/bfloat16, with and without `up`, with and without the clamp, non-default `activation_alpha`/`activation_beta`, packed vs. split parity, and shape validation errors. Fusion-on vs. fusion-off outputs must be exactly equal for float32/float16/bfloat16 on negative-tail and bulk inputs, and both `round_activation` modes are pinned. - `onnxruntime/test/optimizer/swiglu_fusion_test.cc`: fusion fires in either `Mul` operand order and carries `alpha` and `round_activation=1`. It is skipped for broadcast `Mul`, a shared `QuickGelu` output and CPU-assigned nodes. End-to-end CUDA parity (Level1 vs. Level2) is exact (tolerance 0) for float, float16 and bfloat16, including negative-tail inputs. ``` python -m pytest onnxruntime/test/python/transformers/test_swiglu.py # 12 passed onnxruntime_test_all --gtest_filter='*SwiGlu*:*SwiGLU*:*QuickGelu*' # 27 passed, 2 skipped (WebGPU) onnxruntime_test_all --gtest_filter='GraphTransformationTests*:*Fusion*:*SwiGlu*:*SwiGLU*:*QuickGelu*' # 505 passed, 17 skipped (WebGPU) # Two MobileClip CUDA attention-fusion tests need libcudnn on the library path; they pass when it is visible. ``` ## Checklist - [x] Tests added - [x] No breaking changes - [x] Documentation updated --------- Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

) ### Description  Adds a subgroup-matrix Conv path for the WebGPU EP that computes the convolution as an implicit GEMM. Covers F16, channels-last, `group == 1` convs on devices reporting an 8x16x16 subgroup-matrix config, and falls back to the existing Conv paths otherwise. ### Motivation and Context  Conv dominates runtime in vision models, but the existing WebGPU Conv paths do not use the subgroup-matrix units that the MatMul path already exploits on Intel Xe2/Xe3. This brings the same acceleration to Conv.

…soft#33663) ### Description `QDQSelectorActionTransformer` rewrites `DQ -> op -> Q` node groups in two ways: it replaces them with a `QLinear*`/`QGemm` op, or it drops the Q/DQ nodes around ops such as `MaxPool`, `Resize` and `Slice`. Neither rewrite is valid for float8/float4 quantized types: - The `QLinear*`/`QGemm` replacement ops only accept int8/uint8. The schema of `QLinearConv`, `QLinearAdd`, `QLinearConcat`, `QGemm` and the other ops rejects float8. `QLinearMatMul`-21 allows float8 in its schema, but the CPU EP only registers int8/uint8 kernels for it. - After the Q/DQ nodes are dropped, the target op receives the float8 tensor directly. Most of these ops (for example `MaxPool`, `Resize` and `Slice`) either do not accept float8 in their schema or have no float8 kernel on the CPU EP. So for a float8 QDQ model, session creation fails at `ORT_ENABLE_EXTENDED` and `ORT_ENABLE_ALL`, while `ORT_DISABLE_ALL` and `ORT_ENABLE_BASIC` run fine, because the CPU EP has float8 `QuantizeLinear`/`DequantizeLinear` kernels. This PR makes `BaseSelector::Select` leave a node group unchanged when any of its DQ inputs or Q outputs has a float8 or float4 type. The check sits in `BaseSelector`, which only `QDQSelectorActionTransformer` uses. It does not affect the `NodeGroupSelector` classes that EPs reuse for QDQ node unit grouping. One trade-off: for opset >= 21, `Transpose` and `Reshape` do support float8, so dropping their Q/DQ nodes would still have been valid. Their Q/DQ nodes are now kept for float8 as well. That costs only the lossless drop-QDQ optimization for those two ops, and avoids encoding a per-op, per-opset allow list in the selector. ### Motivation and Context Fixes microsoft#33761. Minimal repro (opset 21, CPU EP): ```python import numpy as np import onnxruntime as ort from onnx import TensorProto, helper F8 = TensorProto.FLOAT8E4M3FN inits = [ helper.make_tensor("s", TensorProto.FLOAT, [], [0.125]), helper.make_tensor("zp", F8, [], [0]), helper.make_tensor("w", F8, [4, 4, 3, 3], [0.25] * 144), helper.make_tensor("ws", TensorProto.FLOAT, [], [0.0625]), ] nodes = [ helper.make_node("QuantizeLinear", ["x", "s", "zp"], ["xq"]), helper.make_node("DequantizeLinear", ["xq", "s", "zp"], ["xdq"]), helper.make_node("DequantizeLinear", ["w", "ws", "zp"], ["wdq"]), helper.make_node("Conv", ["xdq", "wdq"], ["y"]), helper.make_node("QuantizeLinear", ["y", "s", "zp"], ["yq"]), helper.make_node("DequantizeLinear", ["yq", "s", "zp"], ["out"]), ] graph = helper.make_graph( nodes, "fp8_qdq_conv", [helper.make_tensor_value_info("x", TensorProto.FLOAT, [1, 4, 8, 8])], [helper.make_tensor_value_info("out", TensorProto.FLOAT, None)], inits) model = helper.make_model(graph, opset_imports=[helper.make_opsetid("", 21)], ir_version=10) x = np.random.uniform(-2, 2, (1, 4, 8, 8)).astype(np.float32) for level in [ort.GraphOptimizationLevel.ORT_DISABLE_ALL, ort.GraphOptimizationLevel.ORT_ENABLE_EXTENDED]: so = ort.SessionOptions() so.graph_optimization_level = level sess = ort.InferenceSession(model.SerializeToString(), so, providers=["CPUExecutionProvider"]) sess.run(None, {"x": x}) ``` With `ORT_ENABLE_EXTENDED`, this fails with: ``` INVALID_GRAPH : This is an invalid model. Type Error: Type 'tensor(float8e4m3fn)' of input parameter (xq) of operator (QLinearConv) in node () is invalid. ``` The same failure happens for float8 QDQ `Add`, `Mul`, `Concat`, `AveragePool`, `GlobalAveragePool`, `Sigmoid`, `LeakyRelu`, `Softmax`, `Gemm`, `MaxPool`, `Resize`, `Slice`, `Gather`, `Expand`, `Tile`, `DepthToSpace` and `Squeeze`/`Unsqueeze`/`Flatten`. `MatMul` fails with `NOT_IMPLEMENTED` for `QLinearMatMul(21)` instead. The current workaround is to disable `QDQSelectorActionTransformer` through `disabled_optimizers`, which also disables all int8 QDQ fusions in the session. ### Testing - Added `QDQTransformerTests.Float8QDQNodeGroupsAreNotFused`. It builds float8 QDQ graphs for `Conv`, `MatMul`, `Add`, `Mul`, `AveragePool`, `Softmax`, `Concat`, `MaxPool`, `Resize` and `Slice`, runs them at Level1 and Level2, checks that the outputs match and that no `QLinear*`/`QGemm` op is created. Without the selector change, all 10 cases fail with the errors above. - Extended `GetZeroPointBytes` in the test graph builder to support `FLOAT8E4M3FN` zero points. - Ran locally on Linux x64 (Release, CPU EP), all passing: `*QDQ*:*Qdq*:*Quant*:*NodeUnit*:*TransposeOpt*:GraphRuntimeOptimizationTest.*` (409 tests) and `*Transform*:*Optimiz*:*Fusion*:*Selector*` (1044 tests, of which 25 are skipped). Co-authored-by: Cursor <cursoragent@cursor.com>

…oft#33745) ### Description  Guard the three disk-backed constant tests with `!defined(__wasm__)`, matching existing optimizer external-data test coverage. Preserve all native tests and WASM in-memory constant tests. ### Motivation and Context  Fix Web CI build failure. --------- Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

@edgchen1

### Description Adds a WebGPU `QuantizeLinear` kernel and int8 support in `Cast`, and fixes byte-lane packing in the segmented (>`maxStorageBufferBindingSize`) output setters. All gaps were found by the W3C WebNN conformance (WPT) suite running on the WebGPU EP with CPU fallback disabled. Commits: 1. **Add WebNN regression test for 1-D blocked uint8 DequantizeLinear.** Adds the WebNN vector "dequantizeLinear uint8 1D tensor with 1D scale, implicit block_size = 2" (microsoft#32730) for f32 and f16 scales, with `block_size` both explicit and inferred. microsoft#32734 fixed this path but only tested int8 with block_size 4. 2. **Add QuantizeLinear kernel.** - f32/f16 input, int8/uint8 output. Registered for opsets 10-12, 13-18, 19-20, 21-22 and 23-24. From opset 23, x / y_scale is computed in the `precision` type (default: the y_scale type). The kernel divides in f32, so opset 23/24 is registered for f32 `y_scale` only. Both GetCapability paths (in-tree `webgpu_execution_provider.cc` and plugin `ep/ep.cc`) leave nodes with an explicit non-f32 `precision` to another EP. - Each invocation writes one packed u32 word (4 outputs), so output bytes are never shared between invocations. - DequantizeLinear's per-tensor/per-axis/blocked classification and shape validation (including block_size inference and the element-wise axis 0, block_size 1 form) is factored into a helper used by both kernels. - Adds an `Int8x4` case to `ShaderVariableHelper::SetByOffset` that packs the low byte of each lane. - WGSL f32 division is only accurate to 2.5 ULP, which flips results near .5 (e.g. `QuantizeLinearOpTest.Int8`: -5 / 0.039215686 = -127.4999...). The shader recomputes the quotient with integer long division when 0.25 <= |q| < 1024, so round-half-to-even matches the CPU kernel bit-exactly. 3. **Support int8 in Cast.** Adds int8 to T1/T2. Input is unpacked with `unpack4xI8` (sign-extending) on both the generic and to-int64 paths; output uses the new Int8x4 packing, which keeps each lane's low byte like a narrowing conversion. 4. **Pack byte lanes in segmented output setters.** Outputs larger than `maxStorageBufferBindingSize` are written through a generated per-segment setter that stored the packed u32 value unchanged. Uint8x4/Int8x4 writers (Cast, Expand, Gather, QuantizeLinear) pass one value per lane, so those shaders failed to compile. The setter now takes `vec4<u32>`/`vec4<i32>` lanes and packs them like the unsegmented `SetByOffset`. This also fixes a latent uint8 bug on main for Cast/Expand/Gather. Tests (WebGPU-pinned, CPU EP fallback disabled) are in `onnxruntime/test/providers/webgpu/tensor/{quantize_linear_test,cast_op_test}.cc`. They cover the WebNN conformance vectors, ties, saturation, per-axis, blocked (first/last axis, partial block), all opsets, `precision` routing, near-tie quotients against the CPU reference, Q -> DQ round trip, int8 casts to/from f32/f16/int32/int64/uint8/bool, and outputs spanning buffer segments. Validation (Windows ARM64, NVIDIA RTX Spark N1X / D3D12, Release): - `build.py --use_webgpu ... --target onnxruntime_provider_test`, then `onnxruntime_provider_test --gtest_filter="*WebGpu*:CastOpTest.*:*QuantizeLinear*:*DequantizeLinear*:ExpandOpTest.*:GatherOpTest.*"`: 1159 tests, 1147 passed, 0 failed, 12 skipped. The skips existed before this change and are unrelated: `QuantizeLinearOpTest.*_DML*` (DML only) and hardware/build-feature skips in `WebGpuMatMulAlgorithmTest`, `WebGpuMatMulNBitsTest.DP4A*` and `WebGpuContextTest.AdapterIndexRejectsUnsupportedBuild`. - New tests that ran and passed: `WebGpuQuantizeLinearTest.*` (including `OpsetVersions` 10-24, `Opset23PrecisionOtherThanFloat`, `Opset23ExplicitFloatPrecision`, `SegmentedOutput`), `CastOpTest.Int8ToOtherTypes_WebGpu`, `CastOpTest.OtherTypesToInt8_WebGpu`, `CastOpTest.ByteOutputsAcrossBufferSegments_WebGpu` and `WebGpuDequantizeLinearTest.*`. - The existing CPU `QuantizeLinearOpTest.*` suite also runs on WebGPU and passes, including `Int8`, which depends on the exact-quotient rounding. - Plugin EP build: `build.py --use_webgpu shared_lib --cmake_extra_defines onnxruntime_BUILD_UNIT_TESTS=OFF --target onnxruntime_providers_webgpu` succeeds. This compiles the `ep/ep.cc` GetCapability change. - Sensitivity check: without the fixes, `Opset23PrecisionOtherThanFloat` returned 5 and 3 (expected 4 and 2), and `SegmentedOutput` and `CastOpTest.ByteOutputsAcrossBufferSegments_WebGpu` failed. Related work: this overlaps with @edgchen1's branch [`edgchen1/webgpu_quantizelinear`](https://github.com/microsoft/onnxruntime/tree/edgchen1/webgpu_quantizelinear) (last updated 2026-04-10, no PR yet). That branch also adds int4/uint4 QuantizeLinear output, which this PR does not cover; int4/uint4 is planned as a follow-up PR. ### Motivation and Context WebNN's `quantizeLinear`, qdq subgraphs and int8 `cast` lower to ONNX `QuantizeLinear`/`Cast`, which had no WebGPU kernel (or no int8 type), so graph build failed with CPU fallback disabled. WPT results (webgfx/webnn-test full suite, Edge Canary 158 arm64, NVIDIA RTX Spark N1X (D3D12), ORT 1.28.3 runtime, WebGPU plugin EP). Before = plugin built from main at 8120434; after = plugin with this change and four companion WebGPU PRs applied, attributed by op. This PR's ops fix 67 subtests (57 required): | WPT test | Before | After | Notes | |---|---|---|---| | quantizeLinear | 4/29 | 15/29 | remaining 14 are optional int4/uint4/int32 outputs | | qdq_subgraph | 1/43 | 43/43 | | | dequantizeLinear | 25/33 | 27/33 | "quantizeLinear then dequantizeLinear" cases | | cast | 38/51 | 50/51 | all int8 casts; remaining failure is a separate f32->f16 rounding issue | For context, with this PR and its four companions applied, failing WPT subtests drop from 343 to 68 (required failures 275 -> 11). Co-authored-by: Copilot App <223556219+Copilot@users.noreply.github.com>

### Description  Two WebGPU EP gaps found by running the W3C WebNN conformance (WPT) suite on the WebGPU EP. **[WebGPU] Support channels-last bilinear Resize** - A 4-D linear `Resize` can scale dims 1 and 2 with channels last (scales `[1, h, w, 1]`). Scale validation accepts this, and WebNN emits it for `resample2d` with `axes = [1, 2]`. The WebGPU EP keeps Resize NCHW-only, so this layout reaches the NCHW bilinear kernel, which always interpolated the last two dims. - `ResizeBilinearProgram` now takes explicit `height_axis` / `width_axis`. For a rank-4 input with `scales[1] != 1` it interpolates dims 1 and 2; otherwise it keeps interpolating the last two dims. The layout is part of the program cache key. - New `onnxruntime/test/providers/webgpu/tensor/resize_op_test.cc`: f32 and f16 channels-last linear upsample (WebNN conformance vector), pinned to WebGPU with CPU EP fallback disabled. **[WebGPU] Add LpPool and GlobalLpPool kernels** - Adds an Lp mode to the shared WebGPU pool program. It accumulates `sum(|x|^p)` in f32 so float16 windows don't overflow, then applies `^(1/p)`. Padding contributes 0. `p = 1` (abs-sum) and `p = 2` (`sqrt` of the sum of squares) are specialized, and other `p` use `pow`. Parallel reduction for large kernels uses an f32 shared buffer in Lp mode. - `PoolProgram` now takes a `PoolType` and `p` instead of an `is_max_pool` flag. Average-pool-only paths (`count_include_pad`, count tracking) are now gated on `kAveragePool`. - Registered in both the ONNX domain and the internal NHWC domain: `LpPool` 11–17, 18–21 and 22+, and `GlobalLpPool` 2–21 and 22+. The internal NHWC domain only defines LpPool 11/18 and GlobalLpPool 2, so LpPool before opset 11 stays on the CPU EP. - New tests in `onnxruntime/test/providers/webgpu/nn/pool_op_test.cc` cover p=1/2/3, strides/pads/dilations, opsets 11/18/22, GlobalLpPool, a large kernel (parallel reduction) and a float16 no-overflow case. Each runs NCHW and NHWC against a double-precision reference, pinned to WebGPU with CPU EP fallback disabled. #### Testing Unit tests: Windows ARM64, NVIDIA RTX Spark N1X (D3D12), Release, built with `build.py --use_webgpu ... --target onnxruntime_provider_test` on this branch (main + these two commits). - `onnxruntime_provider_test --gtest_filter="*WebGpu*:PoolTest.*:ResizeOpTest.*:UpsampleOpTest.*"`: 912 tests, **899 passed, 0 failed**, 13 skipped. - Skipped tests target other EPs (`PoolTest.*_Dml`, `ResizeOpTest.*_CudaRegression*`) or need hardware/build features this machine lacks (`WebGpuMatMulAlgorithmTest.ForcedSubgroupMatrix*` / `ForcedPackedSplitK*`, `WebGpuMatMulNBitsTest.DP4A*`, `WebGpuContextTest.AdapterIndexRejectsUnsupportedBuild`). None are related to this change. - The existing shared LpPool tests (`PoolTest.LpPool*`, `GlobalLpPool`, `LpPool_*` failure cases) now also run on WebGPU and pass. WebNN conformance (webgfx/webnn-test, full WPT suite), Edge Canary 158 arm64, NVIDIA RTX Spark N1X (D3D12), ORT 1.28.3 runtime, private WebGPU plugin EP. Before = plugin built from main; after = plugin built with this PR plus the four related WebGPU PRs (subtest fixes attributed by op). This PR fixes 29 subtests, all of them required: | WebNN op | Before | After | Cause | |---|---|---|---| | `l2Pool2d` | 2/30 | 30/30 | lowers to `LpPool`, which had no WebGPU kernel | | `resample2d` | 13/14 | 14/14 | linear upsample with `axes = [1, 2]` (channels last) | ### Motivation and Context  WebNN `l2Pool2d` and channels-last `resample2d` either fell back to the CPU EP or gave wrong results on the WebGPU EP. This is one of five PRs from the WebNN conformance sweep of the WebGPU EP; the others cover elementwise kernels, precision fixes, quantization and f16 GRU/LSTM. With all five applied, failing WPT subtests drop from 343 to 68 and required failures from 275 to 11, compared with main. Co-authored-by: Copilot App <223556219+Copilot@users.noreply.github.com>

…icrosoft#33652) ### Description  Draft pilot for EP-owned pattern rewriting through the existing capability optimization callback. **Prerequisite integrated:** microsoft#33631 merged as `9d3dc81268ab91a925b103d8bc4cee9807799148`. This branch is rebased onto `main` at `f809e08b988cc7613f2ed4ea40cdbc5c4c98aea4`; the aggregate diff no longer includes the Rama prerequisite. Review the two pilot commits: - `9aa06ce4d2` — internal region-scoped rewriting and test-only EP capability handoff. - `13d14f0f6e` — function-expansion lifecycle documentation and integration coverage. The rebase preserves upstream constant-loading tests and semantics; it does not revert the final merged Rama implementation. - Add an internal region-scoped `PatternRewriter::Run` overload. Only fully contained matches are eligible; existing rule ordering, provider checks and iterative behavior remain intact. - Track removed and newly emitted node IDs in the mutable region. Preserve boundary values and document nontransactional failure behavior. - Add a test-only CPU EP that owns ordered Transpose-composition patterns, discovers a region without changing the graph, and attaches an optimization callback to its capability. After ORT accepts the capability, the callback rewrites the selected region, updates its node list, and uses existing CPU kernels. - Verify mixed graphs, unselected overlapping capabilities, priority/fallback, outside consumers, scoped iteration and replacement DAG tracking, invalid/deleted IDs, missing ONNX domain imports, match-only mode and failures. - Add function-expansion lifecycle coverage using a fixed model-local function and a test-only CPU adapter. With AOT inlining enabled and disabled, verify that discovery does not apply rewrites, a claimed function survives until accepted application, and an unclaimed function expands before fresh body-pattern discovery. This addition changes no production code or expansion/partitioning policy. #### Rebased-head validation Qualified the pilot-only source at `13d14f0f6e05f342a4f42c99f70bddcc981d64cb`, independently of the CUDA fusion commits, in the existing Linux Release CUDA-enabled build. The capability adapters under test use the CPU EP. With the existing virtual environment, CUDA/cuDNN paths and build flags retained: ```sh timeout --signal=TERM --kill-after=30s 900s /usr/local/bin/cmake -S /home/titaiwang/workspace/ort-runtime-fusion-h200/cmake -B /mnt/data/titaiwang/runtime-fusion/build-cuda/Release -DFETCHCONTENT_FULLY_DISCONNECTED=ON -DFETCHCONTENT_UPDATES_DISCONNECTED=ON -DPython_EXECUTABLE=/mnt/data/titaiwang/runtime-fusion/venv/bin/python timeout --signal=TERM --kill-after=30s 7200s /usr/local/bin/cmake --build /mnt/data/titaiwang/runtime-fusion/build-cuda/Release --config Release --parallel 8 --target onnxruntime_test_all onnxruntime_providers_cuda -- -v ``` Results: both commands passed. The build completed 1,460 Ninja actions, compiling core/CUDA sources and linking the CUDA implementation provider; this was not a cached-object/link-only qualification. From `/mnt/data/titaiwang/runtime-fusion/build-cuda/Release`: ```sh timeout --signal=TERM --kill-after=30s 900s ./onnxruntime_test_all '--gtest_filter=FunctionTest.*:CheckAndBindComponentTests.*:MatchComponentTests.*:PatternCapabilityTests.*:PatternComponentTests.*:PatternRewriterTests.*:ReplacementComponentTests.*' --gtest_output=xml:/mnt/data/titaiwang/runtime-fusion/rama-rebase-pilot-engine.xml ``` Result: **283/283 passed**, zero failures/errors/disabled/skips, with no exclusions. This includes both tests previously excluded in the older CPU-only environment. XML counts and log/XML hashes were independently checked. `git diff --check f809e08..integration/ep-pattern-capability-rebase` also passed. Exact commands, build provenance, configuration and artifact hashes are retained locally in `rama-rebase-qualification-report.json`. #### Historical review and validation The original pilot commit `826601e64f14cee31ec35d72faca2e09fc1601a0` completed readability, implementation, critical, deep/spec and integration reviews. No Critical/Major findings were reported. Minor documentation/diagnostic clarification findings were resolved; the missing-domain capability fallback was reproduced with a failing regression test and fixed. The function-lifecycle addition and rebase integration received direct review, not another five-reviewer round. Before this rebase, the CPU-only build and scoped lint passed: ```sh cmake --build build/cpu-capability/Release --target onnxruntime_test_all --parallel 24 lintrunner --take CLANGFORMAT onnxruntime/core/optimizer/pattern_rewriter.h onnxruntime/core/optimizer/pattern_rewriter.cc onnxruntime/test/optimizer/pattern_rewriter_components_test.cc git diff --check ``` Results: all three commands passed on the old pilot revision. From that build directory: ```sh ./onnxruntime_test_all --gtest_filter='PatternCapabilityTests.*:PatternComponentTests.*:PatternRewriterTests.*:FunctionTest.*-PatternRewriterTests.GroupSignaturesAreIndependentOfGlobalLocale:FunctionTest.TestInlinedLocalFunctionNotRemoved' ``` Result: **248/248 passed**. This was an older Linux CPU Release `-O1 -DNDEBUG` build with contrib/ML ops disabled and local-only GCC warning exceptions; it is historical evidence, not qualification of the rebased head. Two tests were explicitly excluded: the locale-signature test segfaulted on both the pilot and a diagnostic executable with the unchanged prerequisite engine; the retained-local-function test failed on both the function-lifecycle addition and its unchanged prior component-test source. Those outcomes are not counted as passes. The new qualification above uses an unexcluded selector. **Limits:** the generic scoped rewriter tracks multi-node replacement results, but the capability adapters are test-only and deliberately handle N-to-1 Transpose chains and one fixed model-local function definition. Recognizing a function name alone is not production function-equivalence checking. No production EP registration, general N-to-M placement, resource accounting/memory estimation, external registration ABI or Qwen qualification is included. The function probe does not compare whole-function versus expanded-body placement alternatives or change cross-EP priority. Full CI and non-Linux platforms have not been validated locally. **Open question (integration review, non-blocking):** production integration must define rewrite policy during ORT-format `kAssignOnly` export. The current optimization callback does not receive the partitioner mode; this pilot does not establish or validate export behavior. ### Motivation and Context  Pilot implementation toward microsoft#33403 and the EP-capability direction in microsoft#33404. An EP should be able to supply ordered patterns while ORT retains placement decisions. The existing `nodes_to_optimize` / `optimization_func` hook provides the acceptance-to-application seam; the missing local primitive was restricting iterative rewriting to the accepted region and returning its updated node membership. This pilot reuses the merged microsoft#33631 matcher and mutation engine rather than introducing a second rewrite pipeline or requiring ORT to select every individual match. It is intended to collect feedback on the region-scoped interface and capability handoff before adding production patterns or memory planning. --------- Co-authored-by: copilot-swe-agent[bot] <198982749+Copilot@users.noreply.github.com> Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> Copilot-Session: 38c05fba-777a-447c-af24-f930f13e4021

…microsoft#33736) ### Description Skip the cuDNN include dir when it is one of the compiler's implicit include directories; the compiler searches it anyway. ### Motivation and Context When cuDNN is installed under the system prefix (e.g. --cudnn_home=/usr), CUDA_PLUGIN_CUDNN_INCLUDE_DIR is /usr/include. CMake drops that implicit directory from the plugin target's own compile lines, but onnxruntime_unittests.cmake copies the target's INCLUDE_DIRECTORIES onto moe_gemm_int2_test.cc as a source property, which is not filtered. The file was then compiled with -isystem /usr/include ahead of the libstdc++ headers, and <cstdlib>'s #include_next <stdlib.h> failed with "stdlib.h: No such file or directory".

### Description  Add standard ONNX `BitShift` to the CUDA EP using the existing elementwise broadcasting dispatcher: - Unsigned 8/16/32/64-bit integers for opsets 11–27; all eight signed/unsigned integer types from opset 28. - Both directions, arithmetic right shifts for signed inputs, and wrapping left shifts implemented with unsigned intermediates and a bit-preserving conversion. - Counts are checked before shifting: negative or oversized counts produce `-1` only for right shifts of negative signed inputs, otherwise `0`, matching the current ONNX schema and CPU implementation. Negative counts do **not** reverse direction. - CUDA-only tests disable CPU fallback and cover integer extrema/high 64-bit patterns, zero/width-boundary/negative counts, scalar and multidirectional broadcasts, per-channel paths, initializer counts, empty tensors, and 2,051-element tails. - Update only the CUDA `BitShift` rows in `OperatorKernels.md`. This change is independent of the other bitwise kernel PRs. **Validation:** Passed on an H200 (`sm90`), Linux, CUDA 13.3, cuDNN 9.10.2, Python 3.12, ORT 1.32.0 Release/QUICK_BUILD. The standalone validation source was base `812043410eebd29c4bdd46b906447e2e5b77f7b2` plus this branch through `a94f033211605b6352f75a371790e803cb648cb8`; no XOR/AND/OR/NOT changes were present. All five compiled C++/CUDA source files match the committed content after CRLF normalization. Build/wheel command, run from the validation source directory — **passed**, including nvcc compilation and device `memcpy`: ```bash ../.venv/bin/python tools/ci_build/build.py --config Release --build_dir /tmp/ort-cuda-bitwise-xor-8d05709c/build --update --build --build_wheel --skip_tests --skip_submodule_sync --use_cuda --cuda_home /usr/local/cuda --cudnn_home /usr/lib/python3.12/site-packages/nvidia/cudnn --cmake_generator Ninja --parallel 16 --disable_contrib_ops --disable_ml_ops --cmake_extra_defines CMAKE_CUDA_ARCHITECTURES=90 FETCHCONTENT_FULLY_DISCONNECTED=ON onnxruntime_QUICK_BUILD=ON onnxruntime_USE_FLASH_ATTENTION=OFF onnxruntime_USE_MEMORY_EFFICIENT_ATTENTION=OFF onnxruntime_USE_TRT_FUSED_ATTENTION=OFF onnxruntime_USE_XQA=OFF --targets onnxruntime_provider_test onnxruntime_pybind11_state ``` From the build's `Release` directory — **42/42 passed, no skips**: 18 new CUDA tests and 24 existing `BitShiftOpTest` tests: ```bash CUDA_VISIBLE_DEVICES=4 LD_LIBRARY_PATH=/usr/local/cuda/lib64:/usr/lib/python3.12/site-packages/nvidia/cudnn/lib ./onnxruntime_provider_test '--gtest_filter=*CudaBitShiftTest*:*BitShiftOpTest*' ``` Additional ad hoc numerical/profile probes, from `validated-shift` with the same environment — **230/230 models passed**, CPU fallback disabled, exact integer results, and all 230 profiled `BitShift` executions assigned to `CUDAExecutionProvider`: ```bash CUDA_VISIBLE_DEVICES=4 LD_LIBRARY_PATH=/usr/local/cuda/lib64:/usr/lib/python3.12/site-packages/nvidia/cudnn/lib ../.venv/bin/python ../probe-remaining-bitwise.py --group shift ``` The probes include six exhaustive 8-bit pair models (all 256×256 value/count pairs for signed opset 28 and unsigned opsets 11/28, both directions: 393,216 scalar results). The ad hoc probe script is validation infrastructure, not an added repository test. The initial 224-model baseline also matched CPU before applying this patch. Operator documentation generation, from the validation task root — **passed**; generated `BitShift` rows match the committed rows: ```bash PYTHONPATH=/tmp/ort-cuda-bitwise-xor-8d05709c/validated-shift/wheel-runtime LD_LIBRARY_PATH=/usr/local/cuda/lib64:/usr/lib/python3.12/site-packages/nvidia/cudnn/lib .venv/bin/python src/tools/python/gen_opkernel_doc.py --providers cuda --output_path validated-shift/OperatorKernels.generated.md ``` Validation is limited to this configuration; the full suite, other GPU architectures, and performance benchmarks were not run. The archive build does not embed a Git revision; provenance was checked using explicit source hashes. Validated wheel SHA256: `791fbb2278876c4dba5df7cb78d43d795209a67edf4629a23cf2de22fb22f132`. ### Motivation and Context  CUDA currently lacks a standard ONNX `BitShift` kernel. Native integer operations avoid CPU fallback and preserve exact wide-integer results without floating-point emulation or hardware-modulo shift-count behavior. --------- Signed-off-by: Justin Chu <justinchuby@users.noreply.github.com> Co-authored-by: Copilot App <223556219+Copilot@users.noreply.github.com>

### Description  Remove `continue-on-error: true` / `continueOnError: true` from browser WebGPU EP tests in GitHub Actions and Azure Pipelines, and from C# delegate validation in the Windows WebGPU workflow. Update the migration release-gate documentation. ### Motivation and Context  Make test failures errors. --------- Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

### Description  Schedule `Python - CUDA13 - Combined WHLs` (Lotus definition 2353) daily at **09:00 UTC** from `main`. Set `always: true` so wheels are built even when the source has not changed. Push and PR triggers remain disabled, and the build continues to publish pipeline artifacts only. The existing Package pipeline health automation has also been updated separately to monitor definition 2353 explicitly. ### Motivation and Context  Generate combined CUDA 13 wheels every morning for downstream publishing. The companion [release-pipelines PR 56020](https://aiinfra.visualstudio.com/Lotus/_git/release-pipelines/pullrequest/56020) repoints the existing 1ES publisher to this pipeline's unified wheel artifacts, preserving its feed options and `main`/`rel-*` completion triggers. Validation: YAML parsing and assertions passed for the exact cron, main-only branch filter, and `always: true`; comparison against the original YAML confirmed all other pipeline configuration is unchanged. `git diff --check` passed. No build or publishing pipeline was queued. The schedule takes effect after merge to `main`; the publishing pipeline's new source resource still requires authorization. Co-authored-by: Copilot App <223556219+Copilot@users.noreply.github.com>

microsoft#28019) ### Description Add CACHE when unsetting `CMAKE_CUDA_ARCHITECTURES` to ensure it's empty. ### Motivation and Context Recently, when building passing `CMAKE_CUDA_ARCHITECTURES=all` configuring fails because `CMAKE_CUDA_ARCHITECTURES` is not being correctly unsetted and then set to the curated list of architectures. Error message: ``` -- Setting CMAKE_CUDA_ARCHITECTURES to all enables a list of architectures OnnxRuntime optimized for, not all architectures CUDA compiler supports. CMake Error at external/cuda_configuration.cmake:154 (message): Unrecognized CUDA architecture: all Call Stack (most recent call first): CMakeLists.txt:715 (setup_cuda_architectures) -- Configuring incomplete, errors occurred! ``` I got this error both on Alma 9 and Ubuntu 24.04 and after this change the build starts to compile. --------- Co-authored-by: Juan Miguel Carceller <jmcarcell@users.noreply.github.com>

…nce (microsoft#33680) ### Description  Adds WebGPU EP kernels and integer type support for ops that the W3C WebNN conformance (WPT) suite lowers to ONNX. Without them, graph build fails when CPU EP fallback is disabled. Each commit includes WebGPU-pinned unit tests with CPU fallback disabled. | Commit | Change | |---|---| | [WebGPU] Add Or and Xor kernels | Registers `Or`/`Xor` (opset 7+, bool) next to `And`. WGSL defines `\|` on `vec4<bool>` but not `^`, so Xor uses `!=`. Tests cover scalar, element-wise (partial vec4), non-vectorized broadcast (including 4-D to 5-D) and both vectorized broadcast paths. | | [WebGPU] Add IsNaN and IsInf kernels | Registers `IsNaN` (9-12, 13-19, 20) and `IsInf` (10-19 f32, 20 f32/f16). Values are classified on raw IEEE-754 bits because WGSL lets implementations assume floats are finite, so `x != x` checks can be folded away. f16 lanes are split out of the raw words instead of being widened to f32. IsInf honors `detect_positive`/`detect_negative`. | | [WebGPU] Register Slice for int32 and uint32 | The Slice shader is type-agnostic, so `T` widens from float types to f32/f16/i32/u32. | | [WebGPU] Add Round, Sign and Softsign kernels | `Round` (11-21, 22; f32/f16): WGSL `round()` rounds half to even. `Sign` (9-12, 13; f32/f16/i32): WGSL `sign()` has no u32 overload. `Softsign` (1-21, 22; f32/f16) is evaluated in f32 and returns ±1 once \|x\| >= 2^24. Dividing directly returned 0 for ±FLT_MAX because GPUs flush the subnormal reciprocal of the divisor. | | [WebGPU] Accept an int64 exponent in Pow | Allows int64 for Pow's `T1` (exponent) constraint when int64 is enabled. The shader already converts the exponent to f32. An int64 base remains unsupported (existing TODO). | | [WebGPU] Support int32 in Abs and Neg | `Abs` accepts all supported number types (u32 is the identity). `Neg` accepts f32/f16/i32. INT_MIN wraps as in two's complement. | #### Testing Built on Windows ARM64 (Release, `build.py --use_webgpu ... --target onnxruntime_provider_test`) and run on an NVIDIA RTX Spark N1X GPU (D3D12): ``` onnxruntime_provider_test --gtest_filter="*WebGpu*:MathOpTest.*:ActivationOpTest.*:SliceTest.*:IsNaNOpTest.*:IsInfTest.*:RoundTest.*:SignTest.*:SoftsignTest.*:AbsTest.*:NegTest.*" ``` Result: 1214 tests, 1200 passed, 0 failed, 14 skipped. All 14 skips already exist on main and are unrelated to this change: - `MathOpTest.MatMulInt32Type`/`MatMulUint32Type`/`MatMulZeroK*`/`MatMulFloat16SharedInitializerWithKernelOwnedPrepack`: known WebGPU MatMul limitations. - `SliceTest.Slice1D_WithNegativeSteps_EndOutOfBounds_2`: an unconditional WebGPU skip on a float test. - `WebGpuMatMulAlgorithmTest.ForcedSubgroupMatrix*`/`ForcedPackedSplitK*`, `WebGpuMatMulNBitsTest.DP4A*`, `WebGpuContextTest.AdapterIndexRejectsUnsupportedBuild`: skipped because the hardware or build doesn't support the feature. `lintrunner` (clang-format, ruff) reports no issues on the changed files. ### Motivation and Context  The W3C WebNN conformance suite found these ops failing on the WebGPU EP. WebNN `logicalOr`/`logicalXor`, `isNaN`/`isInfinite`, `roundEven`, `sign`, `softsign`, `abs`/`neg` (int32) and `slice` (int32) lower to the ONNX ops above. Chromium lowers `reduceSumSquare` over no axes, or over a 0-D input, to `Pow(x, int64 2)`. WPT results (webgfx/webnn-test full suite, Edge Canary 158 arm64, NVIDIA RTX Spark N1X (D3D12), ORT 1.28.3 runtime, WebGPU plugin EP). "Before" is the plugin built from main at 8120434. "After" is the plugin built with this PR plus four related WebGPU WebNN-conformance PRs; results are attributed by op. This PR fixes 106 subtests (105 required): | WebNN op | Before | After | Notes | |---|---|---|---| | logical_or | 2/17 | 17/17 | | | logical_xor | 2/17 | 17/17 | | | is_nan | 1/15 | 15/15 | | | is_infinite | 1/18 | 18/18 | | | round_even | 1/11 | 11/11 | | | sign | 2/8 | 6/8 | remaining 2 are optional int8/int64 | | softsign | 2/19 | 19/19 | | | subgraph | 49/52 | 52/52 | conv2d/convTranspose2d/batchNormalization + softsign | | abs | 18/21 | 19/21 | int32; remaining are optional int8/int64 | | neg | 17/20 | 18/20 | int32; remaining are optional int8/int64 | | slice | 20/21 | 21/21 | | | reduce_sum_square | 37/45 | 45/45 | via Pow(x, int64 2) | With all five PRs applied, failing subtests drop from 343 to 68 compared with main, and required failures from 275 to 11. --------- Co-authored-by: Aditya Rastogi <adityar@ntdev.microsoft.com> Co-authored-by: Copilot App <223556219+Copilot@users.noreply.github.com>

### Description  Add CUDA execution for standard ONNX `BitwiseXor` (opset 18+) for `int8`, `uint8`, `int16`, `uint16`, `int32`, `uint32`, `int64`, and `uint64`. Reuse the existing binary elementwise implementation for multidirectional broadcasting and empty tensors. Boolean `Xor` is unchanged; this does not add other bitwise operators. Add CUDA-only regressions with CPU fallback explicitly disabled, covering all eight integer types at opsets 18 and 24, signed/unsigned boundaries, scalar operands, multidirectional broadcasting in both operand orders, empty outputs, and the INT64 `state XOR (state / 2048)` RNG pattern. Update the CUDA operator support table. Validation on Azure Linux 3 / NVIDIA H200 (physical GPU4), CUDA 13.3, cuDNN 9.10.2, Python 3.12, Release/sm90; source based on main `997d2d1b26` plus this patch: - `python tools/ci_build/build.py --config Release --build_dir /tmp/ort-cuda-bitwise-xor-8d05709c/build --build --build_wheel --skip_tests --skip_submodule_sync --use_cuda --cuda_home /usr/local/cuda --cudnn_home /usr/lib/python3.12/site-packages/nvidia/cudnn --parallel 16 --disable_contrib_ops --disable_ml_ops --targets onnxruntime_provider_test onnxruntime_pybind11_state` — **passed**, producing an ORT 1.32.0 GPU wheel. The build was configured with Ninja, `CMAKE_CUDA_ARCHITECTURES=90`, `onnxruntime_QUICK_BUILD=ON`, and Flash/MEA/TRT-fused-attention/XQA disabled. - From the build output directory: `CUDA_VISIBLE_DEVICES=4 LD_LIBRARY_PATH=/usr/local/cuda/lib64:/usr/lib/python3.12/site-packages/nvidia/cudnn/lib ./onnxruntime_provider_test --gtest_filter=*Xor*` — **45 tests passed**, no skips, including existing boolean XOR tests. - Additional ad hoc wheel probe: **146/146** models produced bit-identical NumPy results, and **146/146** profiled `BitwiseXor` executions used `CUDAExecutionProvider`, with CPU fallback disabled. This covers all integer types, opsets 18/24, full-width random values (including upper 64-bit lanes), scalar/multidirectional/empty cases, and integrated `Div → BitwiseXor` graphs over 4,103 INT64 RNG states. Before installing the patched wheel, the same probes on ORT 1.28.0 (`0368187f84`) produced 146 CPU XOR executions; GPU-only initialization failed as expected. - `clang-format --dry-run --Werror` on all six changed C++/CUDA files — **passed**; `git diff --check` — **passed**. The initial untargeted contrib-disabled build encountered a pre-existing `-Werror=unused-function` in `onnxruntime/test/optimizer/nhwc_transformer_test.cc:71`; the relevant provider-test and wheel targets above built successfully without changing unrelated code. Full-suite, non-sm90, full-model fidelity, and performance validation were not performed. ### Motivation and Context  Current main has CPU and DirectML `BitwiseXor` kernels but no CUDA registration. Integer RNG hashing therefore falls back to CPU, even for INT32/UINT32, introducing repeated device transfers in GPU inference. This was observed while adding fishaudio/drama-3 support in onnxruntime/mobius#793 (2,862 CPU XOR executions in one sampled profile). Supporting the standard operator preserves the model's RNG algorithm and exact integer outputs without a custom operator or graph rewrite. Signed-off-by: Justin Chu <justinchuby@users.noreply.github.com> Co-authored-by: Copilot App <223556219+Copilot@users.noreply.github.com>

## Description Ports the architecture-table part of NVIDIA/TensorRT-LLM#12170 ("Compile XQA with sm_120f") to ORT's XQA decode kernels: `__CUDA_ARCH__ == 1210` (sm_121 — GB10 / DGX Spark and other Blackwell Arm devices) is now listed next to sm_120 in the two compile-time arch tables. sm_121 has the same 99 KB opt-in shared-memory limit as sm_120 and uses the same K/V tile configuration. This is explicit-parity hardening, not a user-visible bug fix: sm_121 already built and ran through the "unknown architecture" `#else` default, which has the same values as sm_120. See Testing for the evidence that the generated code is identical before and after. ## Summary of Changes ### XQA arch tables | File | Change | |------|--------| | `onnxruntime/contrib_ops/cuda/bert/xqa/mha_impl.cuh` | Add `__CUDA_ARCH__ == 1210` to the `860 / 890 / 1200` branch that selects `preferedKHeadPartBytes = 64` and `cacheVTileSeqLen = 32` (previously reached only via the `#else` default). Mention sm_121 in the related comment and in the shared-memory-limit hint text. | | `onnxruntime/contrib_ops/cuda/bert/xqa/utils.cuh` | Add `__CUDA_ARCH__ == 1210` to the 99 KB `kMAX_SMEM_SIZE` branch. Before, 1210 fell through to the 48 KB default. `kMAX_SMEM_SIZE` is currently only used by a `#if 0` `static_assert`, so this has no runtime effect today; it just keeps the table correct. | ### Not ported from TRT-LLM microsoft#12170 (not applicable) The rest of microsoft#12170 changes TensorRT-LLM's NVRTC JIT wrapper (`-arch=sm_120f`, dropping the compute_89 PTX to sm_120 two-stage compile through nvPTXCompiler, and the `nvptxcompiler_static` link). ORT compiles XQA ahead of time with nvcc, so there is nothing to port. In TRT-LLM the default arch branch was an `#error` and `kMAX_SMEM_SIZE` was undefined for 1210, which is why sm_121 failed to compile there; ORT's `#else` fallbacks mean it never did. ## Do the shipped libraries include sm_121? Yes, so no packaging or CMake change is needed: - Linux aarch64 CUDA 13 wheel: `87-real;89-real;90-real;100-real;103-real;110-real;120-real;121-real` (`py-cuda13-packaging-pipeline.yml`). - Unified Linux aarch64 wheel (`py-cuda13-unified-packaging-pipeline.yml`): passes no `-a`, so it takes the `build_linux_python_package.sh` default for CUDA 13.0 on aarch64, which is the same list. - CUDA plugin (`plugin-cuda13-pipeline.yml` / `plugin-cuda-packaging-stage.yml`): Linux aarch64 uses the same list; Windows ARM64 uses `120-real;121-real`. - x86_64 lists only carry `120-real` (sm_121 parts are Arm). - CMake normalizes `121-real` to `121a-real`, and XQA sources are compiled for every selected arch >= 80, so XQA gets sm_121a SASS (no PTX). The CUDA plugin's XQA target (`onnxruntime_providers_cuda_plugin_xqa`) uses the same `MIN_SM 80` filter, with no virtual replacement (that option is only used for the plugin's separate LLM target), so it gets the same SASS on Linux and Windows. ## Testing Built the XQA object library of the CUDA provider (`onnxruntime_providers_cuda_xqa`, CUDA 13.0) with `CMAKE_CUDA_ARCHITECTURES='120-real;121-real'`. The configure log shows `CMAKE_CUDA_ARCHITECTURES: 120a-real;121a-real`. - All 44 XQA objects contain both `sm_120a` and `sm_121a` cubins (`cuobjdump -lelf`), with 158 `kernel_mha` kernels each. Per-kernel register / shared / local / constant usage is identical between sm_120a and sm_121a. - A/B against `main` (without this change): SASS is identical for all 44 objects on sm_120a and sm_121a (600,232 instructions compared; all 88 cubins byte-identical), and also identical on sm_90a for 8 representative translation units (loader, paged, spec-dec, fp8/int8/int4, head sizes 64-512). So existing archs are unaffected and sm_121 behaves exactly as before. - The XQA kernels are present in the final `libonnxruntime_providers_cuda.so` (checked on an sm_90a build: 158 `kernel_mha` kernels). - `lintrunner` and `clang-format` are clean. - Not run on real sm_121 hardware (none available), so a smoke test on a DGX Spark would be welcome. To reproduce the cubin check: ``` python tools/ci_build/build.py --build_dir build/xqa_sm121 --config Release --use_cuda \ --cuda_home <cuda13> --cudnn_home <cudnn> --build_shared_lib --skip_tests --update \ --cmake_extra_defines 'CMAKE_CUDA_ARCHITECTURES=120-real;121-real' ninja -C build/xqa_sm121/Release onnxruntime_providers_cuda_xqa cuobjdump -lelf <obj dir>/xqa_loader_fp16_128.cu.o # lists sm_120a and sm_121a cubins ``` ## Motivation and Context Keeps ORT's XQA arch tables in sync with TensorRT-LLM (NVIDIA/TensorRT-LLM#12170) so sm_121 is handled explicitly rather than by the unknown-arch fallback, and so the `kMAX_SMEM_SIZE` table is right if that check is re-enabled. Independent of the XQA race fixes in microsoft#33703. Overlap with microsoft#33373: its commit "xqa: treat SM121 like SM120 in the per-architecture constants" makes the same two `|| __CUDA_ARCH__ == 1210` edits with identical spelling. This PR carries the XQA part on its own so it does not have to wait for that larger CMake change. The two PRs merge cleanly in either order (checked with `git merge-tree` against the head of microsoft#33373), so whichever lands second just has a no-op hunk. The extra edits here (the sm_121 mention in the comment and the shared-memory-limit hint text) are in different hunks. ## Checklist - [ ] Tests added/updated (N/A: arch-table-only change with identical generated code; verified by cubin comparison above) - [x] No breaking changes - [x] Documentation updated (N/A: only an error-message hint and comments changed) Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

### Description - **Output ownership:** Only the final reducing CTA publishes multi-block attention output; other CTAs cannot overwrite it with partial results. - **Synchronization:** All threads latch the final-CTA flag before its shared memory is reused. Four more `cp.async` drains / warp barriers on the same path (after GEMM0, before `mergeAndSaveOutTile`, at the start of `loadBufAsync` and after the `waitGroup<1>()` in the multi-block merge loop) were compiled only for `XQA_PAGED_INT4`; they are now enabled for every variant. Two more `__syncwarp()`s, after the K-tile wait in the GEMM0 loop and after the V-tile wait in `syncVTileLoad`, are ported from [NVIDIA/TensorRT-LLM#18141](NVIDIA/TensorRT-LLM#18141): `cp.async.wait_group` only waits for the copies of the executing thread, but every lane reads tile data that other lanes of the warp copied. - **Regression coverage:** `XqaMultiBlockTest` (fp16 and bf16) runs 24 independent decode problems per launch on 4 concurrent streams, each with its own workspace and output, with K/V in mapped host memory. That is what makes the races observable; a single decode problem on one stream passes on the unpatched kernel. It covers head sizes 64, 128 and 256 (256 has a single GEMM1 warp group, so the CTA output is staged in a different shared-memory buffer), group sizes 1, 2 and 4, batch 1, 2 and 4, a sliding window with per-head attention sinks, 1 and 4 sub-sequences, direct launches and CUDA graph replay. Every variant captures a 4-stream fork/join graph and asserts that the multi-block grid is launched. The test is skipped below sm_80. ### Motivation and Context GQA decoding with aliased past/present KV caches can intermittently corrupt attention and change generated tokens, especially under CUDA graph capture. Eligible shared-cache decoding selects XQA independently of `sdpa_kernel`, so requesting Flash does not exclude the affected path. ### Testing The A100 section below was measured on the first revision of this PR (8 tests, without the two K / V `__syncwarp()`s and the head size 256 tests); the H200 section covers the current revision. #### A100-SXM4-80GB (sm_80), first revision CUDA 13.0, Release, built with `CMAKE_CUDA_ARCHITECTURES=80`, `onnxruntime_BUILD_UNIT_TESTS=ON`, `onnxruntime_ENABLE_CUDA_EP_INTERNAL_TESTS=ON`, targets `onnxruntime_provider_test` and `onnxruntime_providers_cuda_ut`: ```bash cd build/<config>/Release GTEST_FILTER='XqaMultiBlockTest*' ./onnxruntime_provider_test --gtest_filter='CUDA_EP_Unittest.All' ``` **The test detects the bug.** The previous version of the test passed in 24/24 processes on the unpatched kernel. The current version, 24 independent processes per row with all 8 tests in each: | `mha_impl.cuh` | Result | |---|---| | `main` (both race-fix hunks reverted) | **0/24 pass**: all 8 tests fail in every run, typically within the first few launches per stream | | the two race-fix hunks only | 24/24 pass | | this PR | 24/24 pass (192 test runs, none skipped) | **The four extra sync sites matter.** `compute-sanitizer --tool racecheck` on the fp16 head size 64 test, one repeat per configuration: | | errors | racing instruction pairs | pairs involving `cp.async` | |---|---|---|---| | the two race-fix hunks only | 634 | 347 | 328 | | this PR | 168 | 19 | 0 | The 19 remaining pairs (`LDSM`/`STS`, `LDS`/`STS`) are also present without this change. I believe they are the GEMM0 → GEMM1 tile hand-off, which the `xBar` / `ctaRowMax` mbarriers order and racecheck does not model, but I have not proven that. **Cost of the four sync sites.** CUDA graph replay of 100 back-to-back launches, median of 3 runs on one A100, 12 shapes (head size 64/128, group size 1/4, batch 1/8, length 505–2503): +0.0% to +1.1%, for example 15.57 → 15.68 µs and 25.16 → 25.42 µs. **Regression suites** (no failures; skipped cases need WebGPU, Flash Attention or a PagedAttention EP that this build does not enable, SM89+ for FP8, or SM90+ for cuDNN paged SDPA and H512 MEA): | Suites | Passed | Skipped | |---|---|---| | `GroupQueryAttention{,GrowingCache,Position,ShapeInference}Test.*` | 107 | 25 | | `PagedAttention.*`, `SparsePagedAttention.*` | 45 | 53 | | `*PagedAttentionXqa*`, `*SparsePagedAttentionInt8*` (INT8 XQA) | 12 | 2 (FP8) | | internal `Xqa*` (8 multi-block, 2 H512) | 10 | 0 | #### H200 (sm_90a), current revision CUDA 13.0, Release, built for `sm_90a` with the same options and targets; Flash Attention and memory-efficient attention are not compiled in. ORT runs the same `mha_impl.cuh` kernel on sm_90 (the Hopper `gmma_impl.cuh` is not built). | `mha_impl.cuh` | `XqaMultiBlockTest*` (10 tests) | |---|---| | `main` | **0/10 pass**: every test fails, including both head size 256 tests | | this PR before the K / V `__syncwarp()`s | 10/10 pass | | this revision | 10/10 pass in each of 7 parallel processes (one per GPU), none skipped | The two K / V `__syncwarp()`s are hardening, not a failure I can reproduce: the tests also pass without them on H200, and racecheck does not model intra-warp `cp.async` visibility. They follow from the `cp.async.wait_group` semantics and match [NVIDIA/TensorRT-LLM#18141](NVIDIA/TensorRT-LLM#18141), which includes them in its fix for the same race. The head size 256 tests cover the original race fix for a kernel configuration that was not tested before. **Cost of the two K / V `__syncwarp()`s.** CUDA graph replay of 50 back-to-back launches (min of 9 trials per run), median of 4 alternating A/B rounds, 10 shapes (head size 64/128/256, batch 1–32, length 1024–8192), fp16 and bf16: median change +0.02%, range −0.43% to +0.86%, within the round-to-round noise (about ±1–2%). For example 14.87 → 14.83 µs (head size 128, batch 1, length 8192), 12.01 → 12.06 µs (head size 64) and 20.30 → 20.47 µs (head size 256). **Regression suites on this revision:** | Suites | Passed | Skipped | Failed | |---|---|---|---| | `GroupQueryAttention{,GrowingCache,Position,ShapeInference}Test.*` | 108 | 24 | 0 | | `PagedAttention.*`, `SparsePagedAttention.*` | 47 | 46 | 5 | | `*PagedAttentionXqa*`, `*SparsePagedAttentionInt8*` (INT8 and FP8 XQA) | 14 | 0 | 0 | | internal `Xqa*` (10 multi-block, 2 H512) | 12 | 0 | 0 | The 5 `PagedAttention` failures (`Cuda_SkipInputSanitizerCudaGraphReplayChangingLengths`, `Cuda_SkipInputSanitizerRaggedSmallBatchParity`, `Cuda_XqaNativeFp16CacheMultiTokenBoundFallsBack`, `Cuda_XqaSpecDecInt8CacheHeadSize128NonCausalFallsBack`, `Cuda_XqaSpecDecInt8CacheHeadSize128UnsupportedWidthFallsBack`) are not caused by this PR: they expect a Flash Attention / memory-efficient fallback (`PagedAttention requires FlashAttention ...`), which this build does not compile in, and they fail the same way on `main` and on the PR before this revision. **Not covered:** GPUs other than sm_80 and sm_90a, Windows (the `_WIN32` environment-variable handling in the test was kept from the earlier revision but not compiled here), builds with Flash Attention enabled, and the Python GQA tests. ### Comparison with TensorRT-LLM Compared with [NVIDIA/TensorRT-LLM#18141](NVIDIA/TensorRT-LLM#18141) and every later change to `cpp/kernels/xqa` in TensorRT-LLM (the numbers in the table are TensorRT-LLM pull requests): | TensorRT-LLM change | Status here | |---|---| | [TRT-LLM#18141]: only the final reducing CTA writes the output, `isLastCta` is latched by all threads before shared memory is reused, GEMM0 `cp.async` drain | In this PR | | [TRT-LLM#18141]: `__syncwarp()`s inside the multi-block merge loop | Covered by the barriers at the start of `loadBufAsync` and after `waitGroup<1>()` that this PR enables for every variant | | [TRT-LLM#18141]: `__syncwarp()` after the K-tile and V-tile `cp.async` waits | **Added in this revision** | | [TRT-LLM#18141]: `gemm0ToPostProcBarrier` mbarrier | Not needed: the `__syncthreads()` at the end of GEMM0 pairs with the first `__syncthreads()` of the GEMM1 post-processing (`needMergeGlobal` or `mergeAndSaveOutTile`), which orders the same hand-off; warps that already exited are not waited for | | [TRT-LLM#18141] (sink handled by CTA 0 only) and [TRT-LLM#19220] (sink in the initial row max, zero-denominator guard) | Not applicable: the sink is still added once in the final reduction, with the merged row max. The sliding-window tests check per-head sinks of 4.0–6.0 against the analytic reference with 1 and 4 sub-sequences, and a sink applied twice would move the result by at least 13%, far above the tolerance | | [TRT-LLM#14459] (paged sliding window with unmapped pages) | Not needed: `LaunchSanitizeBlockTable` maps page ids outside `[-1, num_blocks)` to -1, `HeadPtr::operator+` turns them into `nullptr`, and `copyPartialHeadsAsync` zero-fills them | | [TRT-LLM#15739] (speculative-decoding sinks) | Already equivalent here | | [TRT-LLM#17976] (SM120 performance) | `writesOutput` is part of this PR. The V-page cursor rewrite is not needed: ORT only instantiates 128-token pages, and for every head size and `cacheVTileSeqLen`, `nbVItersPerXIter > 1` only occurs with `xIterSeqStride <= tokensPerPage`, where the old and the new cursor stepping are identical | | [TRT-LLM#12170] (`__CUDA_ARCH__ == 1210` in the XQA per-architecture tables) | Not needed for correctness: ORT's unknown-architecture `#else` default already gives sm_121 the sm_120 values. Ported as explicit parity in microsoft#33762 | | [TRT-LLM#10335] (performance), [TRT-LLM#11066] and the Hopper part of [TRT-LLM#14459] (`gmma_impl.cuh`, not built by ORT), [TRT-LLM#15312] (host runner) | Not applicable | [TRT-LLM#10335]: NVIDIA/TensorRT-LLM#10335 [TRT-LLM#11066]: NVIDIA/TensorRT-LLM#11066 [TRT-LLM#12170]: NVIDIA/TensorRT-LLM#12170 [TRT-LLM#14459]: NVIDIA/TensorRT-LLM#14459 [TRT-LLM#15312]: NVIDIA/TensorRT-LLM#15312 [TRT-LLM#15739]: NVIDIA/TensorRT-LLM#15739 [TRT-LLM#17976]: NVIDIA/TensorRT-LLM#17976 [TRT-LLM#18141]: NVIDIA/TensorRT-LLM#18141 [TRT-LLM#19220]: NVIDIA/TensorRT-LLM#19220  - Fixes microsoft#33562 --------- Co-authored-by: copilot-swe-agent[bot] <198982749+Copilot@users.noreply.github.com> Co-authored-by: tianleiwu <30328909+tianleiwu@users.noreply.github.com> Co-authored-by: Tianlei Wu <tlwu@microsoft.com> Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

…s, FileExpertStore, and direct CPU QMoE (microsoft#33450) ### Description Implements the QMoE shared-layout CPU kernel refinement plan in 3 modular commits: 1. **Commit 1 (`a3bc0acc70`) — Shared-Layout INT4 (`qmoe_shared_layout_int4_avx512.cc`, `qmoe_shared_layout_int4_avx2.cc`), Byte-Preserving `FileExpertStore`, and Direct `QMoECPU`**: - Splits INT4 AVX-512 and AVX2+FMA+F16C kernels into dedicated translation units in `onnxruntime_mlas`. - Removes all `UnpackSm80Int4Weights` calls from `QMoECPU` (`moe_quantization_cpu.cc`) and bypasses `PrePack` (`is_packed = false`), routing directly to `ComputeSharedQMoE` (`qmoe_shared_cpu.cc`). - Removes all scale transposition (`transpose_rows`/`transpose_cols`) from `FileExpertStore` so cold-expert bank reads preserve exact source bytes (`scales_are_block_major = false`). 2. **Commit 2 (`6a0b73bec8`) — Shared-Layout INT3 (`qmoe_shared_layout_int3_avx512.cc`, `qmoe_shared_layout_int3_avx2.cc`) and Direct `QMoECPU`**: - Adds AVX-512 and AVX2+FMA+F16C INT3 kernels supporting `block_size` in `{32, 64, 128}` with FP16 and BF16 scales/biases, consuming raw LSB-packed 3-bit expert rows (`3 * K / 8` bytes/row) directly. - Adds direct CPU/CUDA parity tests (`test_qmoe_shared_direct.py`, `test_qmoe_shared_direct_edges.py`). 3. **Commit 3 (`add6ae40ae`) — Shared-Layout NVFP4/MXFP4 (`qmoe_shared_layout_fp4_avx512.cc`) and Shared-Layout Tests**: - Adds `qmoe_shared_layout_fp4_avx512.cc` (`Avx512Nvfp4GemmRange`, `Avx512Mxfp4GemmRange`), direct `QMoECPU` FP4 support, `qmoe_shared_layout_test.cc`, `test_qmoe_shared_layout.py`, and `docs/contrib_ops/cpu/qmoe_shared_layout.md`. --------- Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

## Description Adds a `com.microsoft` `RouterGateUpProj` contrib operator with a CUDA kernel. It is the fused first projection of a MoE layer that has a shared expert and an exporter that concatenates the router weight with the shared expert's gate/up weight into one matrix: ``` Y = X * W W: [K, router_size + 2 * inter] router = T(Y[:, :router_size]) activated = activation(T(Y[:, router_size : router_size + inter]), T(Y[:, router_size + inter :])) ``` This is `MatMul` + `Split` + `SwiGLU` today. For decode (S <= 8) those are three launches with launch gaps between them, and Split/SwiGLU add their own memory traffic; the fused kernel is one launch. ## Summary of Changes | File | Change | |------|--------| | `core/graph/contrib_ops/contrib_defs.cc`, `ms_opset.h` | Schema, shape inference, registration. | | `contrib_ops/cuda/math/router_gate_up_proj.{h,cc}` | CUDA `OpKernel` (bfloat16, float16): validation, output shapes, path selection. | | `contrib_ops/cuda/math/router_gate_up_proj_kernel.cuh` | Skinny tensor-core GEMV kernel (rows <= 8, sm_90+). | | `contrib_ops/cuda/math/router_gate_up_proj_impl.{h,cu}` | Launchers; generic path = cuBLAS GEMM + one split/activation kernel. | | `contrib_ops/cuda/math/swiglu_activation.cuh`, `swiglu_impl.cu` | The SwiGLU element math moves into a shared `__device__` helper used by both ops (no numerical change). | | `test/contrib_ops/router_gate_up_proj_op_test.cc`, `swiglu_op_test.cc` | CUDA op tests. | | `docs/ContribOperators.md`, `docs/OperatorKernels.md` | Operator and CUDA kernel documentation. | ### Operator Attributes: `router_size` (required), `activation` (default `"swiglu"`, the only defined value; other values are rejected), and the SwiGLU operator's `limit`, `activation_alpha`, `activation_beta`. Gate and up are rounded to `T` before the activation, exactly as the unfused graph does; the activation is evaluated in fp32. ### CUDA implementation - **rows <= 8 on sm_90+: one kernel.** The weight is the MMA "A" operand and the tokens are the 8 "B" columns of `mma.m16n8k16`, so one instruction covers 16 output columns x 8 tokens x 16 K. The MMA may sum its 16 K slots in any order as long as A and B agree, so slots are mapped to rows that each lane loads as plain 16-byte vectors (8 columns of 4 rows) and packs with byte permutes: no shared-memory staging or transpose, and every warp row segment is 128 contiguous bytes. A cluster of 8 CTAs splits K and reduces through distributed shared memory in a fixed order, so the result is deterministic. Supported shape: K a multiple of 512 and at most 2560, `router_size` a multiple of 128, `inter` a multiple of 64. W must be 16-byte aligned and X 4-byte aligned, and the kernel image loaded for the device must be compiled for sm_90+ (`cudaFuncGetAttributes`, `ptxVersion >= 90`); a build that only carries older PTX would otherwise run the empty stub. - **Everything else** (prefill, other shapes, sm < 90): cuBLAS GEMM into scratch plus one split/activation kernel. This also covers operands that are not aligned for the fused kernel's vector loads (a view bound through IOBinding can start at any 2-byte offset), a build without an sm_90+ kernel image, and K = 0 (the projection is zeros; the activation still runs). - Programmatic dependent launch (weight prefetch before `griddepcontrol.wait`) is implemented but opt-in (`ORT_ROUTER_GATE_UP_PDL=1`); on H200 it slowed the producer kernel more than it gained. ## Testing - `RouterGateUpProjTest.*` (22 tests): bf16 and fp16, 1/3/5/8 rows, a realistic 2560 x (512 + 2 x 640) shape, clamp/alpha/beta, dynamic (non-initializer) weight, explicit `activation`, the generic path (40 rows, unaligned K), K = 0, and IOBinding views of X and W at aligned and misaligned offsets with prefilled outputs. - `SwiGLUCudaTest.*` (3 tests): separate and merged inputs, clamp/alpha/beta; covers the refactor of the SwiGLU kernel onto the shared helper. ``` onnxruntime_provider_test --gtest_filter='RouterGateUpProj*:SwiGLUCudaTest*' ``` Measured on H200 with Qwen3-Flash-Next (48 MoE layers, K = 2560, router 512, inter 640), speculative decode with 8 rows per step: | | unfused | fused | |---|---:|---:| | router prefix per layer (MatMul + Split + SwiGLU) | about 12 us | about 6 us | | decode, ms per 8-row cycle (2K context) | 14.75 | 14.49 | | decode, ms per cycle (8K context) | 14.88 | 14.66 | | TTFT at 8K (generic path) | 0.853 s | 0.843 s | | MMLU-Pro (800 items) | 83.125% | 83.50% | The fp32 accumulation order differs from cuBLAS, so greedy text can diverge from the unfused graph. ## Checklist - [x] Tests added - [x] No breaking changes (new operator; the SwiGLU refactor keeps the arithmetic identical) - [x] Documentation updated Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

…3760) ### Description `SkipLayerNormalization` / `SkipSimplifiedLayerNormalization` on WebGPU indexed their storage arrays directly (`x[offset]`, `skip[...]`, `output[...]`, `input_skip_bias_sum[...]`). When a tensor is bigger than `maxStorageBufferBindingSize` it gets bound as several segments (microsoft#25962), and direct indexing only ever reaches the first one, so reads past it return 0 and writes past it are dropped. This switches every read and write in the shader to `GetByOffset` / `SetByOffset`, same as `layer_norm.cc` after microsoft#32629. For tensors that fit in one binding the shader computes the same thing as before; the only difference is a `let output_value` before each store. Also adds four tests that use the test-only `max_storage_buffer_binding_size` (16 KiB) with `[3, 2560]` fp32 inputs, so input, skip, output and `input_skip_bias_sum` each split into 2 segments with the boundary inside a row. Each op is tested with and without `input_skip_bias_sum`; the cases with it need 10 / 9 storage bindings and skip on devices with fewer. Outputs are checked against a reference computed in the test. The tests skip when the test factory is not available (EP API adapter / plugin builds). Tested on Windows 11, RTX 3090 Ti (D3D12), Release build with `--use_webgpu`: - The new tests fail without the kernel change and pass with it. With `maxStorageBuffersPerShaderStage` set to 8, the two cases without `input_skip_bias_sum` pass and the other two skip. - `SkipLayerNormTest.*` and `LayerNormTest.*` in onnxruntime_provider_test: 69 passed. - `MatMulNBitsEmptyWeightsFusion.*`, `WebGpu/MatMulNBitsEmptyWeights*` and the WebGPU MoE tests: 52 passed, 1 skipped (skips itself on this GPU). - `GraphTransformationTests.MatMulNBitsMlp*` and `MatMulNBitsQkv*` in onnxruntime_test_all (the fused ops call `RunSkipLayerNormProgram`): 28 passed. ### Motivation and Context Fixes microsoft#33755. On devices with the default 128 MiB binding limit (a lot of phones), any model with fp32 activations got silently wrong output once `seq_len * hidden_size * 4` went over 128 MiB.

…ler, and memory budget ledger (microsoft#33451) ### Description Stacked on microsoft#33450, structured into 2 modular commits: 1. **Commit 1 (`f9507fa74b`) — Session-Scoped MoE/QMoE Offload Planner & Device/Pinned Memory Budget Ledger**: - Adds `session.moe.offload.*` configuration keys, `MoeOffloadSessionPolicy`, `MoeNodeExecutionPlan`, and the 5-stage deterministic admission planner (`moe_offload_plan.h`). - Adds Stage 1 `PlaceMoeOffloadInitializersOnCpu` in `SessionState` (`session_state.cc`) and `CudaDeviceMemoryBudget` / `BoundedPinnedPool` across CUDA EP and Plugin EP allocators/arenas. 2. **Commit 2 (`a4df250105`) — `QMoESessionController` & CUDA EP / Plugin EP Lifecycle Hooks**: - Adds `QMoESessionController` (`qmoe_session_controller.h`) for session-scoped MoE offload admission, shared landing ring management, Linux read-only file-mapping promotion/eviction, and run/graph-replay synchronization. - Wires `OnSessionInitializationEnd`, `OnRunStart`, `OnRunEnd`, and CUDA Graph capture/replay generation tracking in `CUDAExecutionProvider` and `CudaEp` (Plugin EP). ### Documentation - Added `docs/moe-cpu-offload/qmoe_offload.md` documenting the `session.moe.offload.*` configuration keys, external data layout & `PromoteFileMappingToSharedReadOnly`, 5-stage `SessionState` planning & admission pipeline, CUDA memory budget ledgers (`CudaDeviceMemoryBudget`, `SessionBudgetAccount`, `HostPinnedBudgetLedger`), and `QMoESessionController` lifecycle hooks. --------- Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

microsoft#33538) ## Description Adds a paged XQA specialization for multi-token speculative verification (query widths `2..8`) with: - **BF16** query and output (`__nv_bfloat16`) - **INT8** paged KV cache (`PER_TENSOR` or folded `PER_CHANNEL` K/V scales) - **Head size 256** (`H256`) - **GQA group size 6 or 16** (`XQA_PAGED_GROUP6_AND_16`) ### Motivation & Root Cause Models with BF16 activations, INT8 paged KV caches, `head_size = 256`, and `group_size = 16` (`num_heads = 32, kv_num_heads = 2`, such as `Ornith-1.5-122B`) or `group_size = 6` previously missed `xqa_spec_dec_candidate` during speculative verification (`max_query_len` in `2..8`), because `xqa_spec_dec_candidate` gated quantized KV caches on `std::is_same<T, MLFloat16>::value && group_size == 6`. As a result, every speculative verification step fell back to: 1. `GatherAndExpandPagedKVCache` (`3.68 ms/step` at 32K context across 12 full-attention layers, reading and dequantizing INT8 -> BF16 into a staging buffer), followed by 2. Unsplit `flash_fwd_kernel` varlen (`11.49 ms/step` at 32K context with only `batch_size * num_heads = 32` CTAs on a 148-SM H200). With this PR, `PagedAttention` dispatches `H256::LaunchXQAPagedSpecDecInt8KernelBF16` directly on the INT8 paged KV cache in place with multi-block split-KV (`M_TILESIZE = 32`), eliminating `GatherAndExpandPagedKVCache` and cutting 12-layer attention verification time at 32K context from **15.17 ms/step to 0.72 ms/step (21x speedup)**. ## Summary of Changes | File | Change | |------|--------| | `onnxruntime/contrib_ops/cuda/bert/xqa/xqa_paged_spec_dec_bf16_int8_256.cu` | New TU instantiating `H256::LaunchXQAPagedSpecDecInt8KernelBF16` (`HEAD_ELEMS=256`, `XQA_PAGED_CACHE_ELEM=1`, `XQA_PAGED_INPUT_FP16=0`, `XQA_PAGED_SPEC_DEC=1`, `XQA_PAGED_GROUP6_AND_16=1`) | | `onnxruntime/contrib_ops/cuda/bert/xqa/xqa_paged_loader_impl.cuh` | Support `XQA_PAGED_GROUP6_AND_16` (instantiating `grp6_` and `grp16_` with `M_TILESIZE = 32` under `XQA_PAGED_SPEC_DEC=1`) and group-aware `WorkspaceSize` | | `onnxruntime/contrib_ops/cuda/bert/xqa/xqa_paged_loader.h` / `.cu` | Route `is_bf16 && kv_quant_type == XqaQuantType::kInt8 && head_size == 256` at group sizes 6 and 16 to `H256::LaunchXQAPagedSpecDecInt8KernelBF16`, and thread `num_heads`, `kv_num_heads`, `is_bf16` through `GetXQAPagedSpecDecWorkspaceSize` and `GetXQAPagedSpecDecRequiredSharedMemoryBytes` | | `onnxruntime/contrib_ops/cuda/bert/paged_attention.cc` | Enable `xqa_spec_dec_candidate` for `BFloat16` query/output + `INT8` cache + `head_size == 256` + `group_size in {6, 16}` | | `docs/contrib_ops/cuda/paged_attention.md` | Update §8 and §13 documentation for H256 BF16-query/INT8-cache (`group_size in {6, 16}`) speculative XQA | | `onnxruntime/test/contrib_ops/paged_attention_op_test.cc` | Add unit tests for H256 BF16-query/INT8-cache at `group_size = 16` (query widths `2..8`, ragged batches, fragmented page tables, local window, `PER_CHANNEL` scales, CUDA graph replay) and `group_size = 6` | ## Testing & Performance Validated on NVIDIA H200 (`sm_90a`, CUDA 13.0): - `onnxruntime_provider_test --gtest_filter="*PagedAttention*XqaSpecDec*"`: **24/24 passed**. - `lintrunner`: no issues. - End-to-end 8-token DFlash2 speculative verification on `Ornith-1.5-122B` (`12` full-attention layers, `num_heads=32, kv_num_heads=2, head_size=256`, BF16 query + calibrated `PER_CHANNEL` INT8 KV cache): | Metric (32K Context, 12 Full-Attention Layers) | Before (Gather + Flash varlen) | With H256 BF16/INT8 SpecDec XQA | Improvement | |---|---:|---:|---:| | `GatherAndExpandPagedKVCache` (`ms/step`) | 3.68 ms | **0.00 ms** (eliminated) | -3.68 ms | | Attention kernel (`ms/step`) | 11.49 ms | **0.72 ms** | **16.0x faster** | | Total 12-layer attention verification (`ms/step`) | 15.17 ms | **0.72 ms** | **21.1x faster (-14.45 ms/step)** | --------- Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

…t#33770) ### Description Call `Graph::set_sm_version()` on each cuDNN frontend graph that `Conv` and `ConvTranspose` build in `CreateCudnnFeExecutionPlan`. The value comes from the EP's cached `cudaDeviceProp` (`GetDeviceProp()`). Starting with cudnn_frontend 1.24 (pinned since ORT 1.28; `main` pins 1.27.0), `Graph::check_support()` calls `Context::populate_sm_version_from_device()`. That function calls `cudaGetDeviceProperties()` unless an SM version was already set. Both kernels build a new graph per node whenever the input or weight shape changes, so a model with dynamic shapes pays one `cudaGetDeviceProperties()` call per Conv node every time the shape changes. When `target_sm_version > 0`, `populate_sm_version_from_device()` returns early. The value set here is the same one cudnn_frontend would compute for the device, so plan selection does not change. - Two files, five lines each. The commented-out bias-graph code in both files is unchanged. - contrib `FusedConv` uses the legacy cuDNN API (no frontend graph) and is unaffected. - The cuDNN SDPA graphs in `contrib_ops/cuda/bert/cudnn_fmha` are left as they are. They are cached per full graph configuration (per thread), so they rebuild only for new configurations; left out of scope. - This is neutral with cudnn_frontend >= 1.32.0, which reads the SM version with `cudaDeviceGetAttribute()` instead (NVIDIA/cudnn-frontend#1566). ### Motivation and Context Fixes microsoft#33769 On an RTX 5060 (driver 580), `cudaGetDeviceProperties()` costs about 0.9-1.0 ms per call, almost all of it sys time. Tested with a FastConformer speech encoder (77 Conv nodes) where every audio segment has a new length: - every encoder run whose input shape changed costs 77 extra calls, about 75 ms; - end-to-end transcription is 22-42% slower on 1.28.2 than on 1.27.1; - with input lengths alternating over 20 runs: 1,543 calls (1.53 s), and 136.6 ms per run against 64.8 ms on 1.27.1. When a process-level interposer returns a cached `cudaDeviceProp`, the 1.27.1 timings come back and the outputs are byte-identical. This change removes those queries from Conv/ConvTranspose graph construction; I expect the same effect, but it has not been measured in an ORT build. Details and permalinks are in the linked issue. This complements microsoft#33123 (an LRU cache of Conv plans), which helps when shapes recur but not when every shape is new. If microsoft#33123 lands first, the same two lines apply to `plan->graph`. ### Testing - **Not built locally.** The host where I measured has no CUDA toolkit (it has only the pip CUDA/cuDNN runtime wheels), 8 cores and 16 GB RAM. Building the CUDA EP (about 580 `.cu` files) plus core was not practical there. I rely on CI for build and test coverage. - **Unit tests.** The change does not alter results, and the existing CUDA Conv / ConvTranspose tests run this path. They are in `onnxruntime/test/providers/cpu/nn/conv_op_test.cc`, `conv_fp16_test.cc` and `conv_transpose_op_test.cc`. I did not add a test that counts `cudaGetDeviceProperties` calls, because ORT's test infrastructure has no clean way to observe them. I am happy to add one if you want. - **Standalone check of the mechanism, on the same GPU.** - Setup: a small program with cudnn_frontend v1.24.0, ORT's `cudnn_frontend_win_dynamic_loading.patch` applied, and `NV_CUDNN_FRONTEND_USE_DYNAMIC_LOADING` defined. - It builds 100 fp32 `conv_fprop` graphs with new shapes, in the same order as `CreateCudnnFeExecutionPlan`: validate, build_operation_graph, create_execution_plans, check_support, build_plans. - Counted with an `LD_PRELOAD` interposer: 101 calls without `set_sm_version` and 1 call with it (the program's one-time query, standing in for the EP's cache). Time per graph went from 2.15-2.28 ms to 1.29-1.67 ms over 3 runs each. - The patch applies cleanly to `v1.28.2`. - Not yet measured: an ORT build with this patch on the real model. The interposer result above is the end-to-end evidence. 🤖 Generated with [Claude Code](https://claude.com/claude-code) Co-authored-by: Claude Opus 5.5 <noreply@anthropic.com>

### Description  Registers `float16` for the WebGPU EP `GRU` and `LSTM` kernels (opset 7–13 and 14). Before this change they only accepted `T = float`. - The GRU/LSTM shaders already load inputs as f32, accumulate in f32, and convert when they write. Registering `WebGpuSupportedFloatTypes()` is enough for the gate math. - The recurrent state stays in f32 for f16 models: the ping-pong `H`/`C` buffers and the GRU `z`/reset scratch. Otherwise the state would be rounded to f16 at every timestep and the error would compound over the sequence. - The state-copy shaders (initial state → state and state → `Y_h`/`Y_c`/`Y`) now convert with `dst_element_t(...)`. The programs that write a graph output from f32 state use `ProgramTensorMetadataDependency::Type` on that output, so the shader cache key includes the output type. **Tests:** `onnxruntime/test/providers/webgpu/rnn/rnn_float16_test.cc` adds tests pinned to WebGPU, with CPU EP fallback disabled: - `GRUTest.WebGpu_Float_Reference` - `GRUTest.WebGpu_Float16` - `LSTMTest.WebGpu_Float_Reference` - `LSTMTest.WebGpu_Float16` The CPU GRU/LSTM kernels are float32-only, so the f16 tests compare against a host reference implementation. The `*_Float_Reference` tests check that reference against the f32 WebGPU path with a 1e-4 tolerance. Coverage includes forward, reverse and bidirectional runs, `linear_before_reset`, Relu activations, LSTM peepholes, and initial states. **Validation:** Windows ARM64, NVIDIA RTX Spark N1X (D3D12), Release build of `onnxruntime_provider_test` with `--use_webgpu`, on this branch (main c026a50 + this commit): - `onnxruntime_provider_test --gtest_filter="*WebGpu*:*GRU*:*LSTM*:*Rnn*"`: 767 tests; 759 passed, 0 failed, 8 skipped. The skips are hardware or build-feature skips unrelated to this change: `WebGpuMatMulAlgorithmTest.ForcedSubgroupMatrix*`/`ForcedPackedSplitK*`, `WebGpuMatMulNBitsTest.DP4A*`, and `WebGpuContextTest.AdapterIndexRejectsUnsupportedBuild`. - All four new tests ran and passed: `GRUTest.WebGpu_Float_Reference`, `GRUTest.WebGpu_Float16`, `LSTMTest.WebGpu_Float_Reference`, `LSTMTest.WebGpu_Float16`. ### Motivation and Context  The W3C WebNN conformance suite (WPT) found that `gru`, `gruCell`, `lstm` and `lstmCell` float16 tests failed on the WebGPU EP, because the kernels were registered for float32 only. **WPT evidence:** - Suite and setup: webgfx/webnn-test full WPT suite, Edge Canary 158 arm64, NVIDIA RTX Spark N1X (D3D12), ORT 1.28.3 runtime, WebGPU plugin EP. - Builds compared: "before" is built from main at 8120434. "After" has this change plus related WebGPU EP fixes in separate PRs; results are attributed by op. 36 subtests fixed, all required: | Op | Before | After | |---|---|---| | gru | 13/25 | 25/25 | | gru_cell | 5/9 | 9/9 | | lstm | 14/29 | 28/29 | | lstm_cell | 7/13 | 13/13 | One `lstm` subtest still fails: a float32 bidirectional case that is 7 ULP off a 3-ULP tolerance. This PR does not address it. --------- Co-authored-by: Aditya Rastogi <adityar@ntdev.microsoft.com> Co-authored-by: Copilot App <223556219+Copilot@users.noreply.github.com>

…arge-M dequantize + cuBLAS GEMM (microsoft#33689) ### Description This PR adds tensor-core CUDA kernels for the `QMoE` and `MatMulNBits` hot spots of Qwen3.8-Flash-Next (B = 1, `S = 8` DFlash2 speculative decoding): `QMoE` INT4 and INT3 grouped GEMV for BF16 activations with packed bf16x2 INT3 dequantization, a warp-coalesced weight/scale layout for constant INT3 weights that the INT3 decode kernel reads with one 16-byte load per lane, a compact 80-route variant of the INT4 FC2 kernel, a dense symmetric-INT8 `MatMulNBits` GEMV for `M <= 8` (SM 9.0), and a large-`M` route for the same dense INT8 layers (`M >= 512`, default on SM 9.0) that dequantizes the weight once and runs one cuBLAS GEMM instead of the row-chunked CUTLASS mixed-input GEMM. It also triggers the programmatic-dependent-launch (PDL) successor early in the two fpA_intB GEMV kernels. The PR was split by operator. The `DynamicSparseAttention` kernels, the gate-mul early PDL trigger and the hyper-connection sigmoid, which this PR carried before, are now in #33707 (see "Related PRs and review guide"). #33740, which was stacked on this PR, is merged into this branch, so its scope (the coalesced INT3 prepack, the prefill repack that consumes it and the compact INT4 FC2 kernel) is part of this PR and of this description. The last commit also addresses the open review feedback of both PRs: four open threads of this PR and two threads of #33740 whose concerns were still present when it merged (see "7. Review follow-up"). On H200 (`B = 1`, DFlash2 `S = 8`, 2 rotated GPU rounds x median of 5 runs), comparing the **final head of this PR** (`d6813ee3ab`) against a baseline built from the exact same companion PRs (`main` `0a12bc0480` + #33714 `ebde3f13e5` + #33707 `204d6593c0` + #33650 / #32126 `78d59a59c7`; GenAI #2808 `5d561a2f5`; package options `prefill.chunk = 8192` and `qmoe_row_tile_size = 8192` held constant): - **Decode (`ms/cycle` at 2K / 8K / 32K context)** drops from `13.82 / 13.96 / 14.02` ms to `10.85 / 10.99 / 11.24` ms (**-21.4% / -21.3% / -19.9%**) for `all_int4`, `14.29 / 14.66 / 14.59` ms to `10.70 / 10.91 / 11.16` ms (**-25.1% / -25.6% / -23.5%**) for `top16`, and `14.42 / 14.51 / 14.92` ms to `10.71 / 10.87 / 11.23` ms (**-25.7% / -25.1% / -24.7%**) for `top08`. - **Prefill TTFT (`8192 / 32768` prompt tokens)** drops from `648.4 / 2760.9` ms to `579.2 / 2444.2` ms (**-10.7% / -11.5%**) for `all_int4`, `668.2 / 2838.6` ms to `593.1 / 2522.0` ms (**-11.2% / -11.2%**) for `top16`, and `672.3 / 2850.2` ms to `599.2 / 2544.0` ms (**-10.9% / -10.7%**) for `top08` (and **-11.0% / -11.8% / -11.2%** at 8192, **-11.2% / -11.2% / -10.8%** at 32768 when toggling `ORT_FPA_INTB_DEQUANT_GEMM_MIN_M` from `0` to `512` on the final head). Every new tensor-core kernel and the new route have a kill switch, available as an environment variable and as a CUDA session option (`ep.cuda.*`); the bit-identical INT3 bf16x2 dequantization and the early PDL trigger have none. The tensor-core kernels fall back to the existing kernels outside their gates or when the loaded PTX target is older than SM80. #### Summary of Changes ##### 1. QMoE decode (BF16): INT3 bf16x2 dequantization, INT4 and INT3 tensor-core GEMV | File | Change | |------|--------| | `onnxruntime/contrib_ops/cuda/moe/qmoe_kernels.cu` | `DequantizeInt3x8Bf16`: two INT3 codes (`j` and `j + 4`) per 32-bit register from one `lop3.b32`, one `HSUB2` and one `HMUL2`. It is bit-identical to the scalar chain (about 6 instructions per weight); the static instruction count of `FusedFc1FastVecKernel<bf16, 3, 64>` drops from 4096 to 3448. `int3_tc::TryLaunchBf16` is tried first on the vectorized decode path. | | `onnxruntime/contrib_ops/cuda/moe/qmoe_int3_tc.cuh` (new) | Tensor-core grouped FC1 (with the interleaved SwiGLU) and FC2 GEMV for INT3 block-64 weights: `mma.sync.m16n8k16` BF16 with FP32 accumulation, up to 8 routes of one expert per CTA. Every warp owns whole-K m-blocks (FC1 `Int3TcKernel<true, 4, 2, 3>`: 4 warps x 2 m-blocks = 128 rows per CTA, 3 CTAs per SM; FC2 `<false, 4, 1, 5>`: 64 rows, 5 CTAs per SM), so there is no cross-warp K split, no shared-memory reduction and no barrier, and the dynamic shared memory holds only the staged activations (`16 * hidden` or `16 * intermediate` B, at most 48 KB). The kernel is one rolled loop over (m-block, 512-wide super-step) items that requests the weights and scales of the next item (two register sets) before it computes the current one. A compile-time unrolled window does not work, because `ptxas` hoists every independent load of a basic block to its top and spills; only a real loop back edge keeps the prefetch distance (FC1: 116 registers instead of 167). The five K partitions (super-step `s % 5`) keep separate FP32 sums that are added in partition order, which is bit-identical to a cross-warp K split. Dequantization (`bd53cf38e2`) uses the range-free packed `bf16x2` sequence (`lop3.b32` into normal BF16 `{8 + c_lo, 128 + c_hi}`, exact centering by `{12, 132}` via `HSUB2`, and `HMUL2` by the block scale) with unscaled shared-memory activations, so zero effective weights (`code = 4`) center to `+0.0` for every finite BF16 scale and all finite BF16 scales and activations are supported without an exponent-window restriction. | | `onnxruntime/contrib_ops/cuda/moe/qmoe_int3_tc.cuh` (#33740) | `PrePackInt3WeightsKernel` and `PrePackInt3ScalesKernel` rewrite constant INT3 weights and BF16 scales once, at session creation, into a warp-coalesced `mma.sync.m16n8k16` layout (512 K per super-step, 16 N per m-block, one `uint4` per lane per 128-K chunk). `Int3TcCoalescedKernel` replaces the scattered row-local 3-bit reads of `Int3TcKernel` with those coalesced loads. It is bit-identical to `Int3TcKernel` (same dequantization, accumulation order and finalize) and has the same numerics contract. | | `onnxruntime/contrib_ops/cuda/llm/moe_gemm/moe_int2_prefill.{h,cu}` (#33740) | `RepackPrepackedInt3ToSm80Int4ActiveExpertsKernel` and `TransposeScalesActiveExpertsVec8Kernel<T, true>` build the SM80 INT4 operands of the active experts for the INT3 prefill GEMM from the prepacked weights and scales, so the prefill path needs no raw-layout fallback. | | `onnxruntime/contrib_ops/cuda/llm/moe_gemm/moe_gemv_tc.cuh` (new) | Tensor-core INT4 (group 32) GEMV for FC1, interleaved-SwiGLU FC1 and FC2. The weights are the A operand of `mma.sync.m16n8k16`; up to 8 consecutive rows of one expert are its N side and share every dequantized weight tile. The K loop is software pipelined: the packed weights and scales of step `s + 1` are in registers while step `s` is computed, and lanes without a route read the activations of the last route so that every load is unconditional. The loop is deliberately not peeled: `__launch_bounds__` carries the occupancy target (5, 6 or 9 CTAs per SM for 4, 2 or 1 m-blocks per warp), and peeling or carrying the activation registers over the loop makes `ptxas` allocate more registers and lose a resident CTA. FC1 is `MoeGemvTcInt4Kernel<true, 4, 4>` (the four warps split K); FC2 is `<false, 1, 1>` (a warp owns 16 columns over the whole K, 56 registers). The summation order is unchanged by the pipelining. #33740 adds a compact FC2 variant for `expanded_num_rows == 80` (`<false, 1, 2, true>`, `kCompact80`): the leaders of the expert row groups are found with warp ballots, which handles experts with more than 8 routes in several chunks without reading `expert_first_token_offset`; `LoadWeightsAfterMb` lowers the register pressure; and the kernel issues `griddepcontrol.launch_dependents` before its global output store. | | `onnxruntime/contrib_ops/cuda/llm/moe_gemm/moe_gemv.cu` | Try `moe_tc::TryLaunchInt4Bf16` before the scalar kernels in `launch_moe_gemv_int_symmetric` and `launch_moe_gemv_int_symmetric_interleaved_swiglu`. | | `onnxruntime/contrib_ops/cuda/moe/moe_quantization.{h,cc}`, `onnxruntime/contrib_ops/cuda/moe/qmoe_kernels.{h,cu}`, `onnxruntime/contrib_ops/cuda/moe/qmoe_fused_route.cuh`, `onnxruntime/contrib_ops/cuda/llm/moe_gemm/{moe_gemv.h,moe_kernels.h,moe_kernels.cu}` | `QMoE` resolves the two disable switches and the two row thresholds once at kernel construction (session option, else environment variable, else default) and passes them to the launchers as plain parameters (a disable switch sets its threshold to 0). From #33740: `QMoE::PrePack` (with a target-device guard) runs the two prepack kernels (`LaunchPrePackInt3Weights`, `LaunchPrePackInt3Scales`) for the constant INT3 weights and scales that `CanUseInt3TcPrePack` and the constructor's shape checks admit, and the decode path then runs the tensor-core kernel on that layout (`Int3TcCoalescedKernel`) from the first route; `moe_kernels.cu` passes `early_routing_loads = fc2_did_gemv` to `finalizeMoeRoutingKernelLauncher`, and `qmoe_fused_route.cuh` skips the route-group packing when `route_groups == nullptr`. | | `onnxruntime/test/contrib_ops/moe_test.cc` | `QMoETest_Bf16Int4TensorCoreGemvParity` (block 32) and `QMoETest_Bf16Int3TensorCoreGemvParity` (block 64): 12 tokens x top-4 over 8 experts, so that expert row groups are full, partial, single or empty, covering the FC1 / FC2 K loops and K tails, compared with a host reference (per-weight BF16 rounding, FP32 accumulation, abs 0.05 / rel 0.01). No existing QMoE test reached these kernels. `QMoETest_Bf16TensorCoreGemvExtremeScales` runs INT3 and INT4 with FC1 block scales near 2^33 and FC2 block scales near 2^-47, and the mirror image, against the host reference (the activations and the tolerance follow the scale factors), which covers the exact scale window of the INT3 kernel. `QMoETensorCoreSessionOptions` runs the parity case with `ep.cuda.disable_qmoe_int3_tc` / `ep.cuda.disable_moe_gemv_tc` = 0 and 1 while the environment variables hold an invalid value, which shows that explicit options override the environment. #33740 extends both parity tests with `num_rows = 20` (80 routes, which reaches the compact INT4 FC2 kernel and the INT3 prefill path). The three tests of the review follow-up are listed under 7. | ##### 2. Early PDL trigger in the fpA_intB GEMV kernels | File | Change | |------|--------| | `onnxruntime/contrib_ops/cuda/llm/fpA_intB_gemv/dispatcher.h` | The scalar fpA_intB GEMV `kernel` issues `griddepcontrol.launch_dependents` first (`__CUDA_ARCH__ >= 900`). | | `onnxruntime/contrib_ops/cuda/llm/fpA_intB_gemv/int8_tc_gemv.cuh` | `Int8TcGemvKernel` (section 3) does the same. | Both kernels feed the hyper-connection down kernel (`HcDownKernel`), which is launched with PDL and reads their output only after `griddepcontrol.wait`; its constant-weight prefetch now overlaps with them. There is no effect for plain launches, no change of any result and no switch. The same trigger in `ProjectedSigmoidGateMulKernel` is part of #33707; together they shorten an `S = 8` decode cycle by 1.4-1.7%. ##### 3. Dense INT8 tensor-core GEMV for `MatMulNBits` (SM 9.0, `M <= 8`) | File | Change | |------|--------| | `onnxruntime/contrib_ops/cuda/llm/fpA_intB_gemv/int8_tc_gemv.cuh` (new) | The weights are the A operand of `mma.sync.m16n8k16` BF16 (16 output columns x the 1-8 activation rows, FP32 accumulation). The runtime's SM80 interleaved weight layout is consumed as it is: the K permutation is absorbed by choosing which activation words feed each `mma`. Dequantization uses one `fma.rn.bf16x2` (`bytes * 2^126 - 1 = (q - 128) / 128`, exact), one `mul.rn.bf16x2` by `128` to form the exact BF16 integer `q - 128`, and one `mul.rn.bf16x2` by the BF16 group scale, with no post-accumulation unscaling. Every weight therefore equals the scalar kernel's `bf16_rn((q - 128) * scale)` bit-for-bit across the entire finite BF16 scale range (including subnormal scales below `2^-119` down to `2^-133`), with no fixed exponent shift that could overflow or flush a valid tiny weight. Register-ring prefetch, a fixed-order cross-warp reduction (a row's result does not depend on `M`), and `LaunchInt8TcGemvAuto` picks the K split (4 or 8 warps), the prefetch depth and the CTAs per SM from `N / 16` relative to the SM count. The launcher returns false when the loaded PTX target is older than SM80, because the kernel body is empty there. | | `onnxruntime/contrib_ops/cuda/llm/fpA_intB_gemv/dispatcher.h` | `CanUseInt8TensorCore` and `exec_kernel_tc`; the gate sits ahead of the scalar kernel in `dispatcher`, which runs the scalar kernel when `exec_kernel_tc` declines. The gate also requires `alpha == 1`: the tensor-core kernel does not apply `alpha`, so a launch with another value keeps the scalar kernel, which does. | | `onnxruntime/contrib_ops/cuda/llm/fpA_intB_gemv/fpA_intB_gemv.h` | `Params::tensor_core`. | | `onnxruntime/contrib_ops/cuda/quantization/matmul_nbits.{h,cc}`, `onnxruntime/contrib_ops/cuda/llm/fpA_intB_gemm_profiler.{h,cc}`, `onnxruntime/contrib_ops/cuda/llm/gemm_profiler.h`, `onnxruntime/contrib_ops/cuda/llm/gemm_tactic_cache.{h,cc}` | `MatMulNBits` resolves the switch at kernel construction and sets `Params::tensor_core` for both the run and the tactic profiler, so the profiler times the kernel that runs. The setting is part of the in-memory tactic-cache key and of the persisted cache (a new `tensor_core_gemv` column; `kTacticSelectionVersion` 4 -> 5), so sessions with different settings do not reuse each other's profiling results. | | `onnxruntime/test/contrib_ops/cuda_kernels/fpA_intB_gemv_int8_tensor_core_test.cc` (new) | `Int8TensorCoreGemvTest.*` (8 tests, SM 9.0 only): one-hot activations isolate the dequantization and must match the scalar kernel as values (a zero may have either sign), for ordinary and wide scale ranges; uniform codes at extreme scales (+-2^100 to +-2^-119) against the BF16 rounding of the exact product, which a fixed exponent shift would fail; random activations against a double-precision reference for the Qwen3.8 projection shapes, all group sizes and row counts; a row's result must not depend on the row count (bitwise); the launch configuration follows the grid size; launches outside the gate, including one with `alpha != 1`, use the scalar kernel, and `NonUnitAlphaScalesTheOutput` checks that an `alpha = 2` launch runs the scalar kernel and returns the scaled reference. | | `onnxruntime/test/contrib_ops/cuda_kernels/gemm_tactic_cache_test.cc`, `onnxruntime/test/contrib_ops/matmul_nbits_fpa_intb_layout_test.cc` | `GemmTacticCacheTest.TensorCoreSettingsHaveSeparateProfiles` (in-memory key and persisted round trip) and `MatMulNBitsFpAIntBLayout.TensorCoreSessionOptions` (`ep.cuda.disable_fpa_intb_gemv_tc` = 0 and 1 on SM 9.0 while the environment variable holds an invalid value). | ##### 4. Large-`M` dense INT8 for `MatMulNBits`: dequantize once + cuBLAS (SM 9.0 default) | File | Change | |------|--------| | `onnxruntime/contrib_ops/cuda/llm/fpA_intB_gemv/int8_dequant.cuh` (new) | `DequantizeInt8InterleavedKernel` expands the SM80-interleaved INT8 weight (64-wide K tiles, two real columns per interleaved row) into a row-major `[N, K]` BF16 / FP16 matrix. `(q - 128) * scale` is exact in FP32 (an 8-bit integer times a 16-bit float) and rounded once to the activation type, so every weight has the value the GEMV kernels and the CUTLASS converter use. | | `onnxruntime/contrib_ops/cuda/llm/fpA_intB_gemv/fpA_intB_gemv.{h,cu}` | `dequantize_int8_interleaved` launches it and returns false (the caller keeps CUTLASS) for `N % 16 != 0`, `K % 64 != 0`, an unsupported group size, unaligned pointers, or `N * K / 32 >= 2^31`. | | `onnxruntime/contrib_ops/cuda/quantization/matmul_nbits.{h,cc}` | `ComputeFpAIntBDequantGemmMinM` resolves the threshold once per node (session option, else `ORT_FPA_INTB_DEQUANT_GEMM_MIN_M`, else 512 on SM 9.0 and 0 elsewhere) and decides eligibility in the constructor. `ComputeInternal` takes the scratch matrix from the EP allocator, dequantizes into it and runs one cuBLAS GEMM instead of the row-chunked CUTLASS call; the route profiles only in a session with CUDA graphs enabled, where it first tunes the CUTLASS `M` buckets that the later capture will use (see 7). When it declines, the CUTLASS path runs as before. | | `onnxruntime/contrib_ops/cuda/quantization/matmul_nbits_workspace_estimate.h`, `include/onnxruntime/core/framework/resource_accountant.h`, `onnxruntime/core/framework/resource_accountant.cc`, `onnxruntime/core/providers/cuda/cuda_execution_provider.cc` | The scratch is charged by both memory estimates. Level 1 becomes `max(CUTLASS workspace, N * K * sizeof(T))` for every admitted `M` at or above the threshold (`WorkspaceEstimatorConfig::cuda_fpa_intb_dequant_gemm_min_m` carries the option from the execution provider to the estimator; an invalid value leaves the node without an estimate, like an invalid chunk size), and Level 2 (`DeclareWorkspaceRequirements`) declares the same maximum. The Level-2 figure is exact when the route runs and an upper bound when CUDA graph capture or deterministic compute keep CUTLASS. | | `onnxruntime/test/contrib_ops/matmul_8bits_test.cc`, `onnxruntime/test/contrib_ops/matmul_nbits_fpa_intb_layout_test.cc` | `MatMulNBits.Fp16_Int8_LargeM_DequantGemm_Cuda` and `MatMulNBits.BFloat16_Int8_LargeM_DequantGemm_Cuda` run the route against the reference. `MatMulNBitsFpAIntBLayout.Int8LargeMDequantGemmRoundTripsExactly` compares the dequantized matrix with the original weights exactly, and `Int8LargeMDequantGemmSessionOption` checks the option, its precedence over an invalid environment variable and the rejection of invalid values (it skips in a build without the fpA_intB kernels, which do not parse the option). `Int8LargeMDequantGemmCudaGraphCapture` is described under 7. | | `onnxruntime/test/providers/cuda/test_cases/matmul_nbits_e2e_workspace_test.cc`, `onnxruntime/test/framework/resource_accountant_test.cc` | `MatMulNBitsWorkspace.Level1DequantGemmScratch`, `DequantGemmScratchAgreementAcrossLevels` (Level 1 and Level 2 charge the same scratch) and `GetCapabilityBudgetChargesDequantGemmScratch`; the factory test of the accountant also checks `cuda_fpa_intb_dequant_gemm_min_m`. | ##### 5. Session options (shared by 1, 3 and 4) | File | Change | |------|--------| | `include/onnxruntime/core/session/onnxruntime_session_options_config_keys.h` | Six `ep.cuda.*` keys: `disable_fpa_intb_gemv_tc`, `disable_moe_gemv_tc`, `disable_qmoe_int3_tc`, `moe_gemv_tc_min_rows`, `qmoe_int3_tc_min_routes` and `fpa_intb_dequant_gemm_min_m`. The comment of `qmoe_int3_tc_min_routes` states that a value above 1 keeps constant INT3 weights in their raw layout (see 7). | | `onnxruntime/contrib_ops/cuda/cuda_session_options.h` (new) | `ResolveCudaSessionOption`: an explicit session option wins, otherwise the environment variable (an empty variable counts as unset), otherwise the default; the whole string is parsed and range-checked, and an invalid value fails. `ReadCudaSessionOption` applies it to the configuration of a kernel at construction, and the memory estimator calls `ResolveCudaSessionOption` directly so that it and the kernel agree. #33707 adds the same file (same blob) for its own switches. | ##### 6. Documentation | File | Change | |------|--------| | `docs/contrib_ops/cuda/moe_qmoe.md` | BF16 INT4 tensor-core GEMV (including the pipelined K loop), packed bf16x2 INT3 dequantization and the INT3 tensor-core GEMV (gates, switches with their `ep.cuda.*` session options, numerics, the exact scale window and activation range, the SM80 PTX requirement, measured gains), and the prepacked INT3 layout of #33740 (conditions, effect on the route threshold, staging of misaligned bindings, measured gain). Corrects the GEMV row gate from `(0, 8]` to `(0, 80]` (`kMaxProfiledExpandedRows`). | | `docs/contrib_ops/cuda/matmul_nbits.md` | Section 6.5 (SM 9.0 tensor-core INT8 GEMV: gate including `alpha == 1`, dequantization and its scale range, the SM80 PTX requirement, the tactic-cache key), the new section 6.6 (large-`M` INT8: gate, threshold, scratch accounting, CUDA-graph warmup, numerics), the `ORT_DISABLE_FPA_INTB_GEMV_TC` and `ORT_FPA_INTB_DEQUANT_GEMM_MIN_M` rows with their `ep.cuda.*` options in section 8 and the tests in section 9. | | `docs/contrib_ops/cuda/gemm_profiler_cache.md` | Tactic-selection version 5 and the `tensor_core_gemv` key column of the persisted cache. | ##### 7. Review follow-up (this PR and #33740) `97dcea1380` fixes the four open review threads of this PR and the two threads of #33740 whose concerns were still present when it merged. Each fix has a regression test that fails on the code before it (checked by reverting the source changes and rerunning the new tests). | Thread | Problem | Fix | Regression test | |--------|---------|-----|-----------------| | `dispatcher.h` | The dense INT8 tensor-core GEMV has no `alpha` multiplier (the scalar epilogue computes `alpha * acc + bias`), so a launch with `alpha != 1` returned the unscaled product. | `CanUseInt8TensorCore` requires `alpha == 1`; other launches use the scalar kernel. | `Int8TensorCoreGemvTest.NonUnitAlphaScalesTheOutput`, and an `alpha = 2` launch in `UnsupportedLaunchesUseScalarKernel`. | | `qmoe_int3_tc.cuh` | The INT3 tensor-core kernel rescaled its FP32 accumulators by 2^60 at the end, but the shifted accumulators underflow in FP32 first, so every result below about 2^-89 became 0. | The activations are multiplied by 2^60 when they are staged in shared memory (exact) and the accumulators are not rescaled. The kernel now also requires activations of magnitude below 2^68 (`ep.cuda.disable_qmoe_int3_tc=1` otherwise). | `QMoETest_Bf16Int3TensorCoreTinyActivations` (expert output 2^-96, constant and non-constant weights). | | `matmul_nbits.{h,cc}` | In a CUDA-graph session the large-`M` route served the warmup runs, so the CUTLASS tactic of the `M` bucket that the capture falls back to was never tuned, and the capture failed with "No valid fpA_intB MatMulNBits tactic". | A graph-enabled node tunes the CUTLASS buckets of the row chunk and of the remainder before it allocates the scratch copy (profiling cannot run during capture). | `MatMulNBitsFpAIntBLayout.Int8LargeMDequantGemmCudaGraphCapture` (`M = 4096`, capture on a user stream). | | `matmul_nbits_fpa_intb_layout_test.cc` | `Int8LargeMDequantGemmSessionOption` failed in a build without the fpA_intB kernels, which do not parse the option. | The test skips there, as `matmul_4bits_test.cc` does under `USE_FPA_INTB_GEMM`. | the test itself | | `moe_quantization.cc` (#33740) | An externally bound input or output that is not 16-byte aligned made the dispatcher choose a byte-safe kernel, which read the prepacked INT3 weights as raw ones and returned wrong values. | With prepacked weights, a misaligned input or output is staged through aligned scratch, so only the tensor-core kernel runs (an `ORT_RETURN_IF_NOT` checks it). | `QMoETest_Bf16Int3PrepackedMisalignedExternalBuffers` (12-70 rows, bindings at misaligned element offsets, compared with the aligned run and with a raw-layout session). | | `moe_quantization.cc` (#33740) | The prepack ran the tensor-core GEMV from the first route and so ignored an explicit `ep.cuda.qmoe_int3_tc_min_routes` or `ORT_QMOE_INT3_TC_MIN_ROUTES`. | An explicitly configured threshold above 1 disables the prepack, so the raw layout honors it for every launch; an explicit 1 or no setting keeps the prepack. | `QMoETest_Bf16Int3ExplicitTensorCoreMinRoutes` (threshold 256 with FC1 scales of 2^-70, which only the vectorized kernels accept). | | `qmoe_int3_tc.cuh` (`bd53cf38e2`) | Shifted subnormal FMA dequantization (`kDequantShift = 60`) overflowed BF16 scale multipliers at `scale = 2^60` (producing `NaN` on zero effective weights `code = 4`) and scaled activations by `2^60`. | Replaced with range-free packed `bf16x2` dequantization (`lop3.b32` + `HSUB2` + `HMUL2`) and unscaled activations in both `Int3TcCoalescedKernel` and `Int3TcKernel` (FC1 and FC2). | `QMoETest_Bf16Int3TensorCoreZeroWeightsExtremeScalesAndActivations` and `QMoETest_Bf16TensorCoreGemvExtremeScales` (`±2^60`). | | `int8_tc_gemv.cuh` (`d6813ee3ab`) | Dividing weights by `128` before rounding to BF16 flushed or distorted tiny weights below `2^-119` (e.g., `code = 129, scale = 2^-130` flushed to `0`; `code = 131, scale = 2^-127` rounded to `0.03125` instead of `0.0234375` with activation `2^120`). | Multiply `(q - 128) / 128` by `128` in BF16 before multiplying by `scale` (`Bf16x2Mul(Bf16x2Mul(Bf16x2Fma(bytes, 2^126, -1), 128), scale)`), removing `kUnscale` and matching `bf16_rn((q - 128) * scale)` bit-for-bit across all 65,024 finite BF16 scales (down to `2^-133`). | `Int8TensorCoreGemvTest.ExtremeScalesMatchReferenceWeights` and `WideScaleRangeMatchesScalarKernel` (scales down to `2^-133` with `1.0` and `2^120` activations across `M = 1, 8` and all 4 launch configurations). | `moe_qmoe.md`, `matmul_nbits.md` and the comment of the `qmoe_int3_tc_min_routes` key describe the new behavior. ##### Gates and switches Every switch is available as an environment variable and as a CUDA session option (`ep.cuda.*`, declared in `onnxruntime_session_options_config_keys.h`). It is resolved once when the kernel is created, so a change needs a new session, and an explicit session option takes precedence over the environment variable. The 0/1 switches accept only `0` and `1`, the row and route thresholds a positive decimal integer and `fpa_intb_dequant_gemm_min_m` an integer in `[0, 2147483647]`; any other value (including an empty session-option string) fails session initialization. | Kernel | Used when | Environment variable / session option | |--------|-----------|--------| | INT3 bf16x2 dequantization | BF16 activations, 3-bit weights, vectorized decode kernels | none (bit-identical) | | INT3 tensor-core GEMV | BF16, FC1 and FC2 3-bit, block 64, no zero points, no `router_weights`, `33 <= routes <= 256`, `hidden % 512 == 0`, `intermediate % 128 == 0`, `hidden <= 3072`, `intermediate <= 3072`, 16-byte-aligned weights, 4-byte-aligned scales, SM80+. Weights prepacked at session creation (constant weights and BF16 scales, see 1) run it from the first route. Exact for block scales in [2^-66, 2^55) and zero and for activations of magnitude below 2^68; a model outside these ranges (real models use scales of about 2^-13 to 2^0) needs the switch. | `ORT_DISABLE_QMOE_INT3_TC=1` / `ep.cuda.disable_qmoe_int3_tc`; `ORT_QMOE_INT3_TC_MIN_ROUTES` / `ep.cuda.qmoe_int3_tc_min_routes` (default 33; a value above 1 also keeps the weights in the raw layout, so the threshold holds for every launch) | | INT4 tensor-core GEMV | BF16, 4-bit, group 32, `33 <= expanded_num_rows <= 80`, `hidden % 512 == 0`, `inter % 128 == 0`, SM80+ | `ORT_DISABLE_MOE_GEMV_TC=1` / `ep.cuda.disable_moe_gemv_tc`; `ORT_MOE_GEMV_TC_MIN_ROWS` / `ep.cuda.moe_gemv_tc_min_rows` (default 33) | | Early PDL trigger | SM 9.0+, PDL successor | none (no effect for plain launches) | | Dense INT8 GEMV | SM 9.0, BF16 x symmetric INT8 (no zero point, bias or activation scale) in the SM80 interleaved layout, `alpha == 1`, `1 <= M <= 8`, `N % 16 == 0`, `K % 64 == 0`, `K % block_size == 0`, 16-byte-aligned activations and weights | `ORT_DISABLE_FPA_INTB_GEMV_TC=1` / `ep.cuda.disable_fpa_intb_gemv_tc` (part of the tactic-cache key) | | Large-`M` dequantize + cuBLAS | FP16 or BF16 activations, symmetric INT8 without zero points or bias in the SM80 layout (weights prepacked for SM 9.0 keep CUTLASS), `N % 16 == 0`, `K % 64 == 0`, a power-of-two `block_size >= 16` that divides `K`, `N * K * sizeof(T) <= 256 MiB` (the 1.3 GB `lm_head` keeps CUTLASS), `M` at or above the threshold, not under CUDA graph capture (a graph-enabled session first tunes the CUTLASS `M` buckets that the capture uses) and not with deterministic compute | `ORT_FPA_INTB_DEQUANT_GEMM_MIN_M` / `ep.cuda.fpa_intb_dequant_gemm_min_m` (default 512 on SM 9.0 and 0 elsewhere; 0 disables) | The tensor-core kernels above (not the large-`M` route, which needs no PTX) also require a loaded PTX target of at least SM80. If only older PTX is available (for example a `compute_75` image that the driver JIT-compiles for a newer GPU), their `__CUDA_ARCH__ >= 800` bodies are empty stubs that would silently leave the output unwritten, so each launcher checks `cudaFuncGetAttributes(...).ptxVersion >= 80` and otherwise uses the scalar or vectorized kernel. ##### Numerics - INT3 bf16x2 dequantization is bit-identical to the scalar chain (generated-token hashes identical in 5 / 5 runs). - The QMoE tensor-core kernels keep the per-weight BF16 rounding of the scalar kernels and differ only in the FP32 accumulation order (about one BF16 ulp). The INT4 pipelining and the INT3 rewrite are bit-identical to the tensor-core kernels they replace (byte-compared outputs in standalone harnesses; generated-token hashes identical in 18 + 12 decode cells x 5 runs). - The INT3 tensor-core kernel decodes each weight to the once-rounded `(code - 4) * scale * 2^-60` and multiplies the activations by 2^60 when it stages them in shared memory (exact), so the FP32 accumulators hold the true dot products. It is exact for block scales in [2^-66, 2^55) and zero and for activations of magnitude below 2^68. The first version of the kernel accepted every scale, so the rewrite keeps a window far wider than the scales of real models; the harness found no mismatch against the previous kernel for scales multiplied by 2^0, 2^+-40 and 2^+-56, and almost every output differs outside the window (2^60, 2^-70), which shows that the check is sensitive. Before the review follow-up (7) the kernel rescaled the accumulators at the end instead, which returned 0 for every result below about 2^-89 (`QMoETest_Bf16Int3TensorCoreTinyActivations`). - The coalesced INT3 prepack of #33740 changes only where the kernel reads weights and scales: `Int3TcCoalescedKernel` is bit-exact with `Int3TcKernel` (measured in #33740), so it adds no numerical difference. End to end, `engine_behaviours.py` passed 27 / 27 on each of the RC5 `all_int4`, `top16` and `top08` packages (#33740). - The dense INT8 GEMV dequantizes to the same weight value as the scalar kernel for every weight that is finite in BF16 when the scale has magnitude >= 2^-119. A standalone harness compared all 256 code bytes with 210 scales of magnitude 2^-119 to 2^110 (three code-to-column assignments, all four launch configurations) against the BF16 rounding of the exact product, with no mismatch, and `OneHotActivationsMatchScalarKernel` and `ExtremeScalesMatchReferenceWeights` check it in the tests. It accumulates in FP32 in a fixed order, where the scalar kernel accumulates in BF16 chains. Compared with a version that pre-scales the weights by a fixed power of two, the two packed FMAs cost +0.6-2.4% kernel time (M = 1-8, N = 2560-16480, K = 2560-6144, group 32-128) and gave bit-identical outputs on realistic scales. - The large-`M` route produces the same per-weight values as the GEMV kernels and the CUTLASS converter (one rounding of the exact `(q - 128) * scale`). Only the FP32 accumulation order of the GEMM differs from CUTLASS, so a near-tie in the first generated token can flip; of the prompts of the TTFT benchmark this happened only at 2048 tokens, deterministically. - The early PDL trigger changes no result (generated tokens identical in 45 / 45 decode runs, measured with both triggers). #### Testing In-tree CUDA 13.0 Release build, `sm_90a`, one H200 (SM 9.0), on this branch at `97dcea1380` (`onnxruntime_provider_test`; the build has `onnxruntime_USE_FPA_INTB_GEMM=ON` and the flash, memory-efficient and lean attention kernels and the INT4 and FP8 KV caches off): | Command (`--gtest_filter=`) | Result | |---|---| | `*Bf16Int4TensorCoreGemvParity*:*Bf16Int3TensorCoreGemvParity*:*QMoETensorCoreSessionOptions*:*Bf16TensorCoreGemvExtremeScales*` | 4 run, **4 passed** | | `*Bf16Int3TensorCoreTinyActivations*:*Bf16Int3ExplicitTensorCoreMinRoutes*:*Bf16Int3PrepackedMisalignedExternalBuffers*` (new, see 7) | 3 run, **3 passed** | | `MoETest.*:*QMoE*:*Moe*:*MoE*` | 134 run, **121 passed**, 13 skipped (CPU-EP, Mixtral and plugin-only cases), 0 failed; includes the seven tests above | | `MatMulNBits*:MatMulNBitsFpAIntBLayout*:*8bits*:*Int8LargeM*` | 112 run, **101 passed**, 11 skipped (CPU-only and SM 12.1-only cases), 0 failed; includes `MatMulNBitsFpAIntBLayout.TensorCoreSessionOptions`, the two `LargeM_DequantGemm` tests and the three `Int8LargeMDequantGemm*` layout tests (`Int8LargeMDequantGemmCudaGraphCapture` is new) | | `CUDA_EP_Unittest.All`, run as `GTEST_FILTER=-GroupQueryAttentionWorkspaceEstimateTest.ValidatesQuantizedCacheMetadata ./onnxruntime_provider_test --gtest_filter=CUDA_EP_Unittest.All` (see the note below) | 516 run, **505 passed**, 11 skipped (FP8 native GEMV, external-data loader and the flash, lean and packed-attention estimator tests that need kernels this build omits), 0 failed; includes the 8 `Int8TensorCoreGemvTest.*` tests (one new), `GemmTacticCacheTest.TensorCoreSettingsHaveSeparateProfiles` and the 45 `MatMulNBitsWorkspace.*` tests | | `RealAccountantTest.*:ResourceAccountantTest.*` (`onnxruntime_test_all`) | 28 run, **28 passed** on the tree before the follow-up; not rerun, because the follow-up changes only a comment of `onnxruntime_session_options_config_keys.h` that these tests include | - `GroupQueryAttentionWorkspaceEstimateTest.ValidatesQuantizedCacheMetadata` is excluded from the `CUDA_EP_Unittest.All` run because it fails in this build, whose flash attention is off: `group_query_attention_workspace_estimate.cc` offers a backend for an INT4 quantized KV cache only under `USE_FLASH_ATTENTION` (the XQA backend takes 8-bit caches only), and only the INT4 iteration of the test fails. This PR changes neither that file nor its test. The 22 other tests of the suite pass. - Each new test fails on the code before the follow-up. I reverted the source changes, rebuilt, reran the new tests and then restored the sources: `Int8TensorCoreGemvTest.NonUnitAlphaScalesTheOutput` and the `alpha = 2` check of `UnsupportedLaunchesUseScalarKernel` (the tensor-core kernel returned 1088 where 2192 is expected); `Int8LargeMDequantGemmCudaGraphCapture` ("No valid fpA_intB MatMulNBits tactic for M=4096, N=256, K=512. The M bucket was not profiled before CUDA graph capture"); `QMoETest_Bf16Int3TensorCoreTinyActivations` (0 where 2^-96 is expected); `QMoETest_Bf16Int3ExplicitTensorCoreMinRoutes` (the explicit threshold of 256 had no effect, so the tensor-core kernel ran outside its scale window and returned wrong values); `QMoETest_Bf16Int3PrepackedMisalignedExternalBuffers` (first output element -0.042 where 0.326 is expected). The skip of `Int8LargeMDequantGemmSessionOption` has no failing state in this build, because it only matters in a build without `USE_FPA_INTB_GEMM`. - Under `nsys` the session-option tests launch the intended kernels. QMoE (`QMoETensorCoreSessionOptions`): `MoeGemvTcInt4Kernel` and `Int3TcKernel` (FC1 and FC2) with the switches at 0, and the scalar `moe_gemv_kernel`, `moe_gemv_interleaved_swiglu_kernel`, `FusedFc1FastVecKernel` and `Fc2RouteFastVecKernel` at 1. Dense INT8 (`MatMulNBitsFpAIntBLayout.TensorCoreSessionOptions`): 49 `Int8TcGemvKernel` launches in the session with the switch at 0, and the scalar `fpA_intB_gemv::kernel` with no tensor-core launch at 1 (measured again on the final tree, after the `alpha == 1` gate of 7: 49 tensor-core and 49 scalar launches in the test process). Large `M` (`Int8LargeMDequantGemmSessionOption`): with the threshold at 512, set as a session option or as an environment variable, `DequantizeInt8InterleavedKernel` followed by a cuBLAS GEMM (`nvjet_*`); at 0 the CUTLASS `GemmFpAIntB`. - `lintrunner` is clean on all 43 changed files. The CUDA translation units built from the changed files (`moe_gemv.cu`, `moe_kernels.cu`, `qmoe_kernels.cu`, `fpA_intB_gemv.cu` and the six `dispatcher_{bf16,fp16}_{int2,int4,int8}.cu`) also compile for `sm_75`, `sm_86` and `sm_120` (30 / 30 compilations before the follow-up; `moe_kernels.cu` for `sm_120a`, which its existing FP4 conversions require). After the follow-up I compiled again the eight translation units that include the headers it changes (`dispatcher_{bf16,fp16}_{int2,int4,int8}.cu`, `moe_gemv.cu` and `qmoe_kernels.cu`) for the same three architectures: 24 / 24, with no warning. The Hopper dispatchers are built for `90a` only, as in the normal build. - The PTX guards (`cudaFuncGetAttributes(...).ptxVersion >= 80`) have no unit test, because they need an executable that carries only `compute_75` PTX. On the H200 (the driver JIT-compiles the PTX for SM 9.0) such an executable shows that the dense INT8 (four configurations), QMoE INT4 (both variants) and INT3 (FC1 and FC2) launchers decline without launching. - Offline tuning tool (`ORT_RUN_CUDA_TACTIC_CACHE_TESTS=1 python onnxruntime/test/python/onnxruntime_test_python_fpa_intb_tune.py` against the Python package of the 13-commit head `c7e138fe27`, one H200): 14 run, **14 passed**, including `test_fresh_process_reuse_repair_and_lock_contention`, which checks the cache header version and key columns end to end. The later commits do not touch the tool, its test, the profiler or the cache format. - CI: the SM 9.0-only `Int8TensorCoreGemvTest.*` tests and `MatMulNBitsFpAIntBLayout.TensorCoreSessionOptions` skip on `sm_86` runners. The QMoE parity and session-option tests need SM80+, so they run there. The quick-build legs (`onnxruntime_QUICK_BUILD=ON`: Windows CUDA and both CUDA Plugin CIs) do not build the native BF16 INT4/INT8 QMoE runner (`moe_quantization.cc`), so there `QMoETest_Bf16Int4TensorCoreGemvParity` is compiled out, and `QMoETensorCoreSessionOptions` and `QMoETest_Bf16TensorCoreGemvExtremeScales` cover INT3 only; the INT3 parity test still runs. The three new INT3 QMoE tests of 7 need SM80+ and are also built in the quick-build legs, because the INT3 runner is. `Int8LargeMDequantGemmCudaGraphCapture` needs SM80+ and `USE_FPA_INTB_GEMM`, and sets `ep.cuda.fpa_intb_dequant_gemm_min_m` to 512 itself, so it also runs on `sm_86`. `Int8LargeMDequantGemmSessionOption` skips in a build without `USE_FPA_INTB_GEMM`. #### Checklist - [x] Tests added/updated and passing - [x] No breaking API changes. Behavior changes are described under Gates and switches and Numerics: on SM 9.0 the dequantize + cuBLAS route replaces CUTLASS for `M >= 512` (`ORT_FPA_INTB_DEQUANT_GEMM_MIN_M=0` restores it), the INT3 tensor-core kernel is exact only for block scales in [2^-66, 2^55) and activations of magnitude below 2^68 (`ep.cuda.disable_qmoe_int3_tc=1` for others), an explicitly configured `qmoe_int3_tc_min_routes` above 1 keeps constant INT3 weights in the raw layout instead of the prepacked one, and persisted tactic caches of version 4 are no longer reused (`kTacticSelectionVersion` 5) - [x] Documentation updated (if applicable) ### Motivation and Context At `S = 8` with top-10 routing a Qwen3.8-Flash-Next decode step routes about 80 rows over 16-25 experts per layer. The scalar QMoE GEMV dequantized every weight once per routed row and was instruction-issue bound. The 97 dense INT8 projections of a decode cycle (block 32, symmetric) ran on the fpA_intB GEMV at about 11 thread instructions per weight, 62% FMA-pipe utilization and only 28-37% of the H200 DRAM bandwidth. The first versions of the tensor-core MoE kernels were latency bound: `ncu` showed the INT4 kernel waiting on `long_scoreboard` for 60% of its warp time (40% of the DRAM peak), and the INT3 FC1 kernel at 167 registers (2 CTAs per SM) with issue slots 38% busy, `long_scoreboard` 27%, `wait` 20% and the cross-warp K-split barrier 9%. In prefill, the CUTLASS mixed-input GEMM converts every weight again for each row chunk and reaches about 300 TFLOPS on H200, where a BF16 cuBLAS GEMM reaches about 730 TFLOPS: 826 `GemmFpAIntB` launches took 138.6 ms (22%) of an 8192-token chunk. #### Performance on NVIDIA H200 (SM 9.0a, CUDA 13.0) Measured on the Qwen3.8-Flash-Next packages (`hidden = 2560`, `intermediate = 640`, 512 experts, top-10; `all_int4` is the all-INT4 package, `top16` and `top08` are the two INT3-mixed packages), B = 1, DFlash2 `S = 8` decode, on a stack that also contains #32126 and #33650. Decode is in ms per cycle, an A/B of consecutive providers on the same packages (single rounds unless noted). | Change | `all_int4` | `top16` | `top08` | |---|---:|---:|---:| | INT3 bf16x2 dequantization (2K / 8K) | unchanged | -5.1% / -4.8% | -6.3% / -6.2% | | INT4 tensor-core MoE GEMV (2K) | -15% | -5% | -2.3% | | INT3 tensor-core MoE GEMV vs. the bf16x2 step (2K / 8K) | n/a | -5.8% / -5.1% | -6.8% / -6.7% | | Dense INT8 tensor-core GEMV (2K / 8K / 32K; mean of 2 rounds) | -3.8% / -4.4% / -3.8% | -5.6% / -5.5% / -4.6% | -5.1% / -4.9% / -6.9% | | Software-pipelined INT4 MoE GEMV (2K / 8K / 32K; mean of 2 rounds) | -1.8% / -2.0% / -1.9% | -0.5% / -0.9% / -1.1% | -0.6% / +0.3% / -0.4% | | Rolled-pipeline INT3 MoE GEMV (2K / 8K / 32K; mean of 2 rounds) | n/a | -4.3% / -4.4% / -4.7% | -5.6% / -5.4% / -5.8% | The gain of the pipelined INT4 kernel scales with the number of INT4 MoE launches per cycle (96 / 32 / 16), and the `+0.3%` cell is a noisy round (-0.5% and +1.1%). The dense INT8 numbers were measured with an earlier dequantization that folded a fixed power-of-two shift into the scale. The range-safe one in this PR (see Numerics) takes 0.6-2.4% longer in a standalone kernel benchmark (cold weights, `M = 1-8`, `N = 2560-16480`, `K = 2560-6144`, group size 32-128) and gave bit-identical outputs on realistic scales; the end-to-end numbers were not re-measured with it. The early PDL triggers (here and in #33707) shorten a cycle by 1.4-1.7%. Cumulative 2K decode of the first 13 commits, ms per cycle (`all_int4` / `top16` / `top08`): | Stack | 2K decode | |---|---| | Without this PR | 14.69 / 15.13 / 15.31 | | + bf16x2 INT3 dequantization, INT4 / INT3 tensor-core MoE GEMV, early PDL | 12.31 / 12.72 / 12.90 | | + dense INT8 tensor-core GEMV = **the first 13 commits of this PR** | **11.83 / 12.04 / 12.21** | The second row also changes the packages' prefill chunk size and QMoE row tile (package configuration, no code), which does not affect decode. It includes the early PDL trigger of the gate-mul kernel, which is in #33707 now. The INT8-KV DSA prefill step that was between these two rows (prefill only, decode 12.30 / 12.76 / 12.86) is not part of this PR. The three later commits were measured as separate steps (rows above), so their gains are not added to this cumulative figure. Kernel level (H200): | Kernel | Before | After | |---|---|---| | INT3 FC1 per layer, `S = 8`, first tensor-core version | 53.9 us (`FusedFc1FastVecKernel`) | 43.1 us | | INT3 FC2 per layer, `S = 8`, first tensor-core version | 33.6 us (`Fc2RouteFastVecKernel`) | 19.1 us | | INT3 FC1 (`n = 1280`, `k = 2560`), rolled-pipeline rewrite (standalone, 80 routes over 56 experts, L2 flushed, including the ~5.5 us launch floor) | 49.06 us | 30.37 us (-38.1%) | | INT3 FC2 (`n = 2560`, `k = 640`), rewrite (same harness) | 24.16 us | 22.37 us (-7.4%) | | INT4 FC1, pipelined K loop (standalone, 80 rows over 29 experts, L2 flushed, including the ~5.5 us launch floor) | 29.2 us | 23.2 us (-20%) | | INT4 FC2, pipelined K loop (same harness) | 19.0 us | 16.1 us (-15%) | | Dense INT8 `in_proj` (`N = 16480`, `K = 2560`, `M = 8`, cold weights) | 24.9 us | 13.0 us | | Dense INT8 `out_proj` (`N = 2560`, `K = 6144`) | 12.7 us | 9.7 us | | Dense INT8 `N = 13952`, `K = 2560` | about 20 us | 11.8 us | | All dense INT8 projections of a decode cycle | 1.84 ms | 1.20 ms | | Dense INT8 `in_proj`, `M = 512` / `M = 8192`, CUTLASS vs. dequantize + cuBLAS (standalone, BF16, group 32) | 128 us / 2050 us | 110 us / 944 us (the dequantization is 50 us of it) | | Dense INT8 `out_proj`, `M = 8192` (same setup) | 723 us | 350 us (dequantization 22 us) | TTFT of the large-`M` route in the model (A/B of this change alone, `all_int4` / `top16` / `top08`, median of 5 runs; decode is unchanged). The route is taken from `M = 512`, so the 128- and 256-token columns are run-to-run noise: | Package | 128 | 256 | 512 | 1024 | 2048 | 8192 | 32768 | |---|---:|---:|---:|---:|---:|---:|---:| | `all_int4` | -1.4% | 0.0% | -4.0% | -5.8% | -8.5% | **-11.8%** | **-11.3%** | | `top16` | -0.7% | 0.0% | -1.9% | -4.9% | -7.4% | **-11.1%** | **-10.7%** | | `top08` | -0.2% | -0.5% | -2.6% | -5.0% | -7.3% | **-11.5%** | **-10.6%** | `all_int4` absolute: 2048 tokens 188.6 -> 172.6 ms, 8192 tokens 650.9 -> 574.0 ms, 32768 tokens 2754.8 -> 2442.5 ms. In an 8192-token chunk (`nsys`, `all_int4`) the critical path drops from 628.8 to 551.6 ms: the 826 CUTLASS `GemmFpAIntB` launches (138.6 ms) are replaced by 97 `DequantizeInt8InterleavedKernel` launches (3.5 ms), cuBLAS GEMMs (31.5 ms for the dense layers) and 2.2 ms of remaining `GemmFpAIntB` (the `lm_head` and small layers). The threshold was swept over 128 / 256 / 512 on the 128-512 token prompts: 128 was 3-5% slower, 256 0-3% slower and 512 2.5-3.5% faster than CUTLASS, so 512 is the default (the extra pass reads 1 B and writes 2 B per weight and the GEMM reads 2 B, which only pays off once the GEMM is large). For the whole stack that contains this PR, #33707, #32126, #33650 and the package configuration, the median over 4 rounds against the previous release candidate is -32.8 / -36.0 / -36.9% TTFT at 8192 tokens, -32.3 / -35.2 / -36.1% at 32768 tokens, and -27.6 / -24.6 / -24.2% decode ms per cycle at 2K (`all_int4` / `top16` / `top08`); these numbers cannot be attributed to a single PR. ##### Final-head isolated benchmarks (`d6813ee3ab`, other PRs held constant) To isolate this PR's final diff from #33707, #33714, #33650, #32126, and the package configuration, two CUDA plugin providers were built from the same companion stack (`main` `0a12bc0480` + #33714 `ebde3f13e5` + #33707 `204d6593c0` + #33650 / #32126 `78d59a59c7`, run with GenAI #2808 `5d561a2f5` and `prefill.chunk = 8192` / `qmoe_row_tile_size = 8192`): `base` (without #33689, md5 `4dc3654493ea4cf8e7b847162d5ab968`) and `final` (with #33689 at `d6813ee3ab`, md5 `d8f2bdb16e4db21966f10a7822da054e`), and benchmarked over 2 rotated GPU rounds on H200 (5 runs per cell per round): | Package | Stack | TTFT 128 | TTFT 256 | TTFT 512 | TTFT 1024 | TTFT 2048 | TTFT 8192 | TTFT 32768 | Decode 2K / 8K / 32K (ms/cycle) | |---|---|---:|---:|---:|---:|---:|---:|---:|---| | `all_int4` | `base` (no #33689) | 0.0463 s | 0.0465 s | 0.0703 s | 0.1120 s | 0.1890 s | 0.6484 s | 2.7609 s | 13.82 / 13.96 / 14.02 | | | `final` (#33689 `d6813ee3ab`) | 0.0460 s | 0.0468 s | 0.0687 s | 0.1068 s | 0.1733 s | 0.5792 s | 2.4442 s | 10.85 / 10.99 / 11.24 | | | **delta (`final` vs `base`)** | -0.7% | +0.7% | **-2.3%** | **-4.6%** | **-8.3%** | **-10.7%** | **-11.5%** | **-21.4% / -21.3% / -19.9%** | | `top16` | `base` (no #33689) | 0.0509 s | 0.0582 s | 0.0840 s | 0.1277 s | 0.2063 s | 0.6682 s | 2.8386 s | 14.29 / 14.66 / 14.59 | | | `final` (#33689 `d6813ee3ab`) | 0.0510 s | 0.0584 s | 0.0817 s | 0.1212 s | 0.1889 s | 0.5931 s | 2.5220 s | 10.70 / 10.91 / 11.16 | | | **delta (`final` vs `base`)** | +0.3% | +0.4% | **-2.8%** | **-5.0%** | **-8.4%** | **-11.2%** | **-11.2%** | **-25.1% / -25.6% / -23.5%** | | `top08` | `base` (no #33689) | 0.0524 s | 0.0610 s | 0.0872 s | 0.1312 s | 0.2086 s | 0.6723 s | 2.8502 s | 14.42 / 14.51 / 14.92 | | | `final` (#33689 `d6813ee3ab`) | 0.0527 s | 0.0616 s | 0.0865 s | 0.1267 s | 0.1956 s | 0.5992 s | 2.5440 s | 10.71 / 10.87 / 11.23 | | | **delta (`final` vs `base`)** | +0.6% | +1.0% | **-0.7%** | **-3.4%** | **-6.2%** | **-10.9%** | **-10.7%** | **-25.7% / -25.1% / -24.7%** | Per-feature ablations on `final` (`d6813ee3ab`, 2 rotated GPU rounds): - **Dense INT8 tensor-core GEMV (`ORT_DISABLE_FPA_INTB_GEMV_TC=1` -> `0`, range-free `Dequant`)**: standalone H200 `M = 8` latencies are `24.9 -> 14.3 us` (`in_proj` `16480 x 2560`), `12.7 -> 10.6 us` (`out_proj` `2560 x 6144`), and `~20.0 -> 12.6 us` (`13952 x 2560`). In-model decode `ms/cycle` at 2K / 8K / 32K drops by **-4.5% / -4.7% / -4.3%** (`11.37 -> 10.85`, `11.54 -> 10.99`, `11.74 -> 11.24` ms) on `all_int4`, **-5.3% / -4.1% / -5.2%** (`11.29 -> 10.70`, `11.38 -> 10.91`, `11.78 -> 11.16` ms) on `top16`, and **-5.1% / -5.3% / -4.9%** (`11.29 -> 10.71`, `11.47 -> 10.87`, `11.81 -> 11.23` ms) on `top08`. - **Large-`M` INT8 dequantize + cuBLAS (`ORT_FPA_INTB_DEQUANT_GEMM_MIN_M=0` -> `512`)**: TTFT at `512 / 1024 / 2048 / 8192 / 32768` tokens drops by **-2.6% / -3.5% / -6.8% / -11.0% / -11.2%** (`all_int4`), **-3.8% / -5.4% / -7.4% / -11.8% / -11.2%** (`top16`), and **-0.6% / -3.7% / -6.5% / -11.2% / -10.8%** (`top08`). - **INT3 tensor-core decode (`bd53cf38e2` range-free `bf16x2` dequantization, H200 standalone, 80 routes over 56 experts, L2 flushed)**: FC1 (`1280 x 2560`) takes `33.02 us` raw (`Int3TcKernel`) and `30.96 us` prepacked (`Int3TcCoalescedKernel`); FC2 (`2560 x 640`) takes `23.76 us` raw and `21.78 us` prepacked. #### Accuracy MMLU-Pro, 800 questions x 3 packages = 2400 items, greedy decoding through the full engine, `max_new = 8192`. The kernels of this PR change only the accumulation precision or order: the dequantizations are bit-exact (INT3 bf16x2, dense INT8 GEMV, large-`M` dequantization), and the INT4 pipelining and the INT3 rewrite are bit-identical to the kernels they replace, and the coalesced INT3 prepack of #33740 is bit-exact with the INT3 kernel it feeds, so no systematic bias is expected. The measurements are of the accumulated stack, not of this PR alone: - The first-round stack (the 13-commit version of this PR, which included the INT8-KV prefill kernel that is now in #33707, plus other changes of the stack) against the previous release candidate: **-0.67 pp** pooled (95% CI -1.36 .. +0.03 pp, exact McNemar p = 0.076); no package differs at the 5% level. The last step alone, the dense INT8 GEMV against the stack before it: 23 / 18 discordant items (p = 0.53). - With the large-`M` route and the split-K DSA decode kernel of #33707 added: 668 / 678 / 667 correct answers against 665 / 676 / 670 for the previous release candidate, pooled **+0.08 pp** (95% CI -0.60 .. +0.77 pp, p = 0.91); no package differs at the 5% level. This run used the stack before the INT4 pipelining, the INT3 rewrite and the sigmoid change, which are bit-exact and were not run separately. - MMLU-Pro covers the large-`M` route only marginally: 27 of its 800 prompts have 512 or more tokens, and 6 of the 61 changed answer letters belong to them. - These MMLU-Pro runs predate the review follow-up of #33707, which keeps more FP16 bits of the INT8-KV attention probabilities below 2^-14 and is therefore not bit-exact. On the final heads of this PR, #33707, #33714, #32126 and #33650, `engine_behaviours.py` passes on the three packages, and TTFT and decode ms per cycle are within the noise of the stack before it (details in the Accuracy section of #33707). The generated tokens differ in near-tie cases (4 of 126 first tokens of the TTFT prompts), and MMLU-Pro was not repeated for it. A drift of up to about 1.4 pp cannot be excluded with 2400 items. ##### Final-head long-prompt accuracy verification for the large-`M` route (`d6813ee3ab`, `min_m = 512` vs `0`) 1. **Kernel execution confirmation (`ORT_FPA_INTB_DEBUG=1`)**: across prompt lengths `256, 512, 1024, 2048, 4096, 8192, 32768` on `final`, all **97** eligible dense INT8 `MatMulNBits` layers per prefill pass launch `kernel=dequant+cuBLAS` at `512, 1024, 2048, 4096, 8192` tokens (`388` launches across four 8192-token chunks at `32768`, and `0` launches at `256` or when `ep.cuda.fpa_intb_dequant_gemm_min_m = 0`, where all 97 run CUTLASS `GemmFpAIntB`). 2. **Node-level accuracy on real Qwen3.8-Flash-Next weights (`M = 512, 1024, 2048, 4096, 8192`)**: probing all `93.6M` weights of the `16480 x 2560`, `2560 x 6144`, and `13952 x 2560` nodes with `M = 512` identity slices gives `0` mismatches between `dequant+cuBLAS` and CUTLASS. Against an FP64 reference of those exact BF16 weights, `dequant+cuBLAS` has `1.11e-4 .. 1.87e-4` relative RMS error across all shapes and `M = 512..8192`: its outputs are **bitwise identical (`0.0` difference)** to CUTLASS whenever CUTLASS selects `split_k = 1` (`16480 x 2560` and `13952 x 2560` at all `M`, and `2560 x 6144` at `M = 4096, 8192`), and are closer to the FP64 reference than CUTLASS on `2560 x 6144` at `M = 512, 1024, 2048` (`1.83e-4` vs `2.15e-3 .. 2.98e-3` relative RMS, cosine similarity `>= 0.9999956`), where CUTLASS reduces split-K partial sums in BF16. 3. **End-to-end long-prompt comparisons (`ep.cuda.fpa_intb_dequant_gemm_min_m = 512` vs `0`, rest of stack unchanged)**: - **55 natural-text prompts per package (`10` each at `512, 1024, 2048, 4096, 8192` tokens and `5` at `32768` tokens = `165` pairs across `all_int4`, `top16`, `top08`)**: first generated token agrees in **`158 / 165` (`95.8%`)** runs (`53/55` `all_int4`, `53/55` `top16`, `52/55` `top08`; `100%` at `4096` and `32768` where all layers use `split_k = 1`), and full 32-token greedy sequences are bitwise identical in `104 / 165` (`63.0%`). - **27 natural MMLU-Pro prompts with `>= 512` tokens (`512..1150` tokens) x 3 packages (`81` evaluations)**: `min_m = 512` scores **`54 / 81` (`66.67%`: `17/27` `all_int4`, `17/27` `top16`, `20/27` `top08`)** vs **`51 / 81` (`62.96%`: `17/27`, `16/27`, `18/27`)** for `min_m = 0`, with **`80 / 81` (`98.8%`)** first-token agreement. - **30 long-context MMLU-Pro prompts per package padded with natural-text history to exact lengths `1024, 2048, 4096` (`10` per length = `90` evaluations)**: `min_m = 512` scores **`57 / 90` (`63.3%`: `21/30`, `18/30`, `18/30`)** vs **`56 / 90` (`62.2%`: `19/30`, `19/30`, `18/30`)** for `min_m = 0`, with **`81 / 90` (`90.0%`)** answer-letter agreement and **`84 / 90` (`93.3%`)** first-token agreement. - **`engine_behaviours.py`**: **27 / 27 PASS** on all three packages (`all_int4`, `top16`, `top08`) with `final` (`d6813ee3ab`). #### Related PRs and review guide - #33707 is the sibling PR (INT8-KV `DynamicSparseAttention` prefill and decode kernels, the gate-mul early PDL trigger and the hyper-connection sigmoid), split from this PR by operator. The changes have no kernel in common. `cuda_session_options.h` is added by both PRs with identical content, and the two edits of `onnxruntime_session_options_config_keys.h` are in different places, so the PRs merge in either order; a trial merge of both branches is clean. - Independent of #32126 (SwiGLU operator and fusion) and #33650 (RouterGateUpProj): no file overlap, and trial merges with both heads and with `main` are clean. The decode graph measured above contains both. - The PR has 43 files and 20 commits, above the usual review size. It is updated only by fast-forward pushes (the first review round saw the 13-commit head `c7e138fe27`), so that the existing review threads and commit hashes are kept. - Commits 1-13 are the previous head: commits 1-3 and 7 (QMoE), commit 4 (early PDL), commit 5 (INT8-KV DSA prefill), commits 6 and 8 (dense INT8 GEMV; commit 8 reads its environment variable with `_dupenv_s` on Windows), commit 9 (docs), commit 10 (review follow-up: range-safe dense GEMV dequantization, the FP32 fallback of the INT8 prefill kernel, the PTX guards and the session options with their tests and docs), commits 11-13 (CI follow-ups: an MSVC narrowing warning in a DSA test, the BF16 INT4 QMoE tests compiled out in quick builds, and `fpa_intb_tune.py`, its test and the cache layout doc for tactic-selection version 5). - Commit 14 (`b32d9eac54`) moves everything `DynamicSparseAttention`, hyper-connection and gate-mul out to #33707. It keeps the early PDL trigger of the fpA_intB GEMV kernels and `cuda_session_options.h`. The hunks of commits 4, 5, 9, 10 and 11 for the moved files are reverted, so the net diff of this PR contains none of them. The earlier review threads on the prefill kernel and its test were answered with the fixes of commit 3 of #33707. - Commits 15-17 are new: the large-`M` dequantize + cuBLAS route, the INT4 K-loop pipelining and the INT3 rewrite, each with its own measurements in its commit message. They can be reviewed separately. - Commits 18-20 come after that: `2b5ceeb045` relaxes the zero-point tolerance of the BF16 large-`M` `MatMulNBits` test (the zero-point variants run CUTLASS), `a8360fcbfe` is #33740 (coalesced INT3 prepack, the prefill repack that consumes it and the compact INT4 FC2 GEMV), and `97dcea1380` is the review follow-up under 7. - The session-option helper and the config keys are shared by the QMoE, dense GEMV and large-`M` groups, and `fpA_intB_gemv/dispatcher.h` is shared by the early-PDL and dense GEMV groups, in different hunks. --------- Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

…tion in f32 (microsoft#33682) ### Description  Fixes WebGPU EP precision and overflow bugs found by the W3C WebNN conformance (WPT) suite. Most of them involve float16. Each commit is self-contained and adds WebGPU-only unit tests with CPU EP fallback disabled. - **Fix f16 reduction overflow and ReduceLogSumExp overflow** (`reduction/reduction_ops.cc`) - Shared-memory reduction: f16 inputs now accumulate in f32 (in both the accumulator and the workgroup array) and convert once at the end. - Naive reduction: ReduceSumSquare and ReduceL2 now square in f32 instead of squaring in f16 and then widening. - ReduceLogSumExp: now computed as `max + log(sum(exp(x - max)))` in f32 on both paths (two-phase on the shared path, online on the naive path). Infinite maxima and `-inf` inputs are handled the same way as in the CPU kernel. - **Keep InstanceNormalization scale and shift in f32** (`nn/instance_norm.cc`): the per-channel scale/shift tensor used to be created with the input type. With float16 subnormal inputs and `epsilon = 0`, `inv_std` is 2^24, which overflows f16. The tensor is now created as f32 and applied in f32 on both the NCHW and NHWC paths, with a single conversion to the output type. - **Evaluate Erf in f32** (`math/unary_elementwise_ops.{h,cc}`, `contrib_ops/webgpu/bert/bias_gelu.cc`, `bias_split_gelu.cc`): `ErfImpl` becomes `ErfImpl(components)` and evaluates `erf_v` in f32 at the caller's vector width. In f16, the final `1 - poly * exp(...)` cancels badly for small inputs. Gelu, BiasGelu and BiasSplitGelu share the helper. - **Evaluate Exp in f32** (`math/unary_elementwise_ops.cc`): f16 `exp()` was up to 2 ULP off; WebNN allows 1 ULP. - **Evaluate Tan with an exact argument reduction** (`math/unary_elementwise_ops.cc`): the built-in `tan()` reduces its argument imprecisely on some GPUs. The new code reduces by pi/2 split into three parts and evaluates a minimax polynomial on [-pi/4, pi/4] in f32. - The reduction is exact while |k| < 2^13 (|x| < ~12867). Beyond that, the built-in `tan()` is used. - The partial differences are OR-ed with a uniform-derived zero the compiler can't prove. This stops the compiler from reassociating the three parts back into one f32 constant; WGSL has no `precise`. New tests (WebGPU EP only, CPU EP fallback disabled): - `ErfTest.WebGpu_Float16` - `ExpTest.WebGpu_Float16` - `TanTest.WebGpu_Float_WebNNConformance`, `TanTest.WebGpu_Float_ReductionRange`, `TanTest.WebGpu_Float16_WebNNConformance` - `InstanceNormalizationOpTest.ZeroEpsilonSubnormal_webgpu` - `ReductionOpTest.*_WebGpu`: f16 ReduceL2/ReduceSum/ReduceMean/ReduceLogSum overflow on the shared and naive paths, and ReduceLogSumExp with large magnitudes, ±Inf rows, f16 and empty input Validation: Windows ARM64, NVIDIA RTX Spark N1X (D3D12), Release, `build.py --use_webgpu ... --target onnxruntime_provider_test`. ``` onnxruntime_provider_test --gtest_filter="*WebGpu*:ReductionOpTest.*:InstanceNormalizationOpTest.*:*Erf*:*Exp*:*Gelu*:TanTest.*:MathOpTest.*:ActivationOpTest.*" ``` Result: 1696 tests, 1680 passed, 0 failed, 16 skipped. All of the new tests pass, as do the existing Gelu, BiasGelu and BiasSplitGelu tests that use the shared `ErfImpl`. The 16 skips are existing and unrelated to this change: - WebGPU MatMul limitations: `MathOpTest.MatMul{Int32,Uint32,ZeroK*}Type`, `MatMulFloat16SharedInitializerWithKernelOwnedPrepack` - CPU-only tests matched by `*Exp*`: `MoETest.QMoETest_CPU_*`, `SparseAttentionIndexerTest.*ExplicitZeroScale*` - Hardware or build-feature skips: `WebGpuMatMulAlgorithmTest.ForcedSubgroupMatrix*` / `ForcedPackedSplitK*`, `WebGpuMatMulNBitsTest.DP4A*`, `WebGpuContextTest.AdapterIndexRejectsUnsupportedBuild` ### Motivation and Context  The WebNN conformance tests in WPT (run via webgfx/webnn-test) fail on the WebGPU EP for these ops because of float16 overflow or precision loss. Setup: Edge Canary 158 arm64, NVIDIA RTX Spark N1X (D3D12), ORT 1.28.3 runtime, WebGPU plugin EP. - Before: plugin built from main at 8120434. - After: plugin built with this PR and four related WebGPU PRs; results below are attributed by op. This PR fixes 37 subtests, all of them required: | Test | Before | After | Cause | |---|---|---|---| | erf | 10/15 | 15/15 | f16 approximation lost ~2% for small inputs | | exp | 9/15 | 15/15 | f16 `exp()` up to 2 ULP off | | tan | 10/15 | 15/15 | built-in `tan(45.5952)` off by 2.4e-3 vs WebNN's 1/1024 ATOL | | reduce_l2 | 31/44 | 44/44 | f16 accumulation overflow | | reduce_log_sum_exp | 44/46 | 46/46 | `exp` overflow; now `max + log(sum(exp(x - max)))` | | reduce-log-sum-small-values | 19/21 | 21/21 | f16 accumulation overflow | | instance-normalization-zero-epsilon | 5/9 | 9/9 | per-channel scale exceeded the f16 range | Across the whole suite with all five PRs applied, failing subtests drop from 343 to 68 versus main, and required failures drop from 275 to 11. --------- Co-authored-by: Aditya Rastogi <adityar@ntdev.microsoft.com> Co-authored-by: Copilot App <223556219+Copilot@users.noreply.github.com>

… branch-free hyper-connection sigmoid (microsoft#33707) ### Description This PR adds tensor-core CUDA kernels for the INT8-KV-cache `DynamicSparseAttention` of Qwen3.8-Flash-Next (12 full-attention layers, 24 query / 2 KV heads, head size 256): a prefill kernel and a split-K decode kernel. It also triggers the programmatic-dependent-launch (PDL) successor early in `ProjectedSigmoidGateMulKernel` and makes the sigmoid of `HyperConnectionLowRankMix` branch free. It is the second half of microsoft#33689, which was split by operator: microsoft#33689 keeps the tensor-core `MatMulNBits` (fpA_intB) and QMoE / MoE GEMM kernels, and this PR carries the DynamicSparseAttention, gate-mul and hyper-connection changes. The two PRs are independent. The only file both add is `cuda_session_options.h`, and this PR's version is a superset of the one in microsoft#33689 (see "Related PRs and review guide"). **Both INT8 tensor-core kernels are opt-in.** `ep.cuda.dsa_prefill_int8_tc` / `ORT_DSA_PREFILL_INT8_TC` and `ep.cuda.dsa_decode_int8_tc` / `ORT_DSA_DECODE_INT8_TC` default to `0`, which runs the previous kernels bit for bit, and the two switches are independent. The end-to-end accuracy and speed results were measured for an earlier revision of the kernels, before the numerics changes of the review, so they are preliminary. The kernels stay off by default until the final revision is qualified end to end (see "Accuracy"). At the operator level on an H200 (`benchmark_dynamic_sparse_attention.py`, added by this PR) the tensor-core decode kernel is 3.3x (1 row) to 5.6x (8 rows) faster than the CUDA-core kernel it replaces, and the prefill kernel is 2.7-3.0x faster than the BF16-dequantizing kernel. Against an FP64 reference the decode error stays within 3% of the previous kernel and the prefill error is 1.6-2.5 times smaller. A CTA whose inputs break the error bounds of the FP16 path is evaluated on a slower FP32 CUDA-core path, which can take longer than the previous kernels (see "Numerics"); the Qwen3.8-Flash-Next layers stay far from its conditions. In earlier-revision end-to-end runs the kernels lowered TTFT by 4.0-9.5% and the `S = 8` decode cycle by 4.6-6.8% (preliminary), and the branch-free sigmoid lowered TTFT by 0.3-1.1%, which is about the size of the repeat noise. #### Summary of Changes ##### 1. DynamicSparseAttention: tensor-core prefill for the INT8 KV cache | File | Change | |------|--------| | `onnxruntime/contrib_ops/cuda/bert/dynamic_sparse_attention_prefill_int8.cuh` (new) | `DynamicSparseAttentionInt8TcKernel<kSplitK>` (`kSplitK = false` for prefill): grid (rows, KV heads), 128 threads, 64-candidate tiles, GQA group <= 16 padded to the `m16n8k16` M dimension. The gathered K/V rows stay INT8 in shared memory (`cp.async`, double buffered, 3-deep index ring) and the `ldmatrix` fragments are converted exactly to FP16 in registers, for FP16 `mma` with FP32 accumulation. `key_scale` is folded into the query once per CTA (`Q' = fp16(Q * key_scale)`) and `value_scale` is applied to the output once, after the accumulator has been divided by 2^15 (so a tiny `value_scale` cannot underflow). Scores and the running maximum stay in the natural-log domain, and only the difference to the new maximum is scaled by log2(e) for `ex2`. The FP16 probability tile holds `p x 2^15` and the row sum stays unscaled FP32. A CTA that cannot bound the FP16 rounding evaluates its query heads in FP32 on the CUDA cores instead (`AttendRowFp32`; the decision is uniform per CTA and the FP32 path starts from the inputs, so nothing is clipped or rounded): an element of `Q * key_scale` is not finite or above 65504; for one head `eps = scale * 128 * sum_d abs(Q'_d - fp16(Q'_d))`, which bounds the shift of a logit by the rounding of `Q'`, exceeds 0.125 nats (`kInt8TcQueryErrorBudget`); a `value_scale` is not finite; or some channel has `abs(value_scale) > 2^16` and a valid candidate is more than 20 nats (`kInt8TcTailDepth`) below the running maximum (a probability below 2^-29 is subnormal in the `p x 2^15` tile). The kernel is selected only for `scale <= 1` (`kInt8TcMaxScale`; the op validates the scale as finite and non-negative, and the default `1 / sqrt(head_size)` is 0.0625 here), because the rounding bound grows with the scale; a larger scale takes the BF16-dequantizing kernel. | | `onnxruntime/contrib_ops/cuda/bert/dynamic_sparse_attention_impl.cu` | Dispatch to the new kernel when its switch is on, the softmax `scale` is at most 1, the buffers are aligned and the loaded PTX target is at least SM80 (`cudaFuncGetAttributes`, otherwise an SM75 PTX image would run an empty kernel stub); `DynamicSparseAttentionPrefillKernel<true>` stays as the fallback. | | `onnxruntime/contrib_ops/cuda/bert/dynamic_sparse_attention.{h,cc}`, `dynamic_sparse_attention_impl.h` | The four `ep.cuda.dsa_*` options are resolved once at kernel construction (session option, else environment variable, else default) and passed to the dispatch as parameters. The two INT8 tensor-core switches default to off. | | `onnxruntime/test/contrib_ops/dynamic_sparse_attention_op_test.cc` | The expected output of the INT8 cases is an FP64 reference computed from the inputs as the kernel receives them (BF16 query, INT8 codes, scales) with max-subtracted weights, rounded to BF16 once. `Int8CachePrefillTensorCoreKernel_CUDA` sweeps query heads {2, 12, 24, 32} x selected candidates {1, 64, 65, 509} x invalid indices with the tensor-core kernel on and off, on random INT8 rows with per-channel scales. A two-candidate fixture (two selected candidates of an otherwise empty cache, so that the output follows from a few numbers) checks the output to one BF16 ulp, each case once with both candidates in one 64-candidate tile and once with them in the first and the last of three tiles, which the decode kernel spreads over up to three splits (two of them active): `Int8TensorCoreQueryFoldingRoundingError_CUDA` (an in-range query whose FP16 rounding cancels the logit gap: 0.5 instead of 1), `Int8TensorCoreLargeCommonLogit_CUDA` (a large common logit, 8290560 at scale 1) and `Int8TensorCoreTinyValueScale_CUDA` (`value_scale = 2^-135`, a BF16 subnormal result). `Int8TensorCoreRandomQueryFolding_CUDA` uses random rows whose folded query products are not FP16 numbers, on both sides of the error budget, and `Int8TensorCoreLargeValueScaleRandom_CUDA` random rows with value scales of 2^14 to 2^17 and logit spreads of 0.7, 11 and 22 nats for the tail-depth check. The FP32 paths are also covered by `Int8PrefillCancellingLargeQuery_CUDA` and `Int8PrefillQueryBeyondHalfRange_CUDA` (four head counts and three candidate counts with invalid entries and the query scaled by 2^28, which sends every CTA to the FP32 path, plus a case where only some CTAs take it). `PrefillSessionOptions_CUDA` checks that options override the environment, that invalid values fail (the failure cases are not compiled with `ORT_NO_EXCEPTIONS`) and that only the strings `0` and `1` are accepted. Misaligned-buffer cases (8-byte K/V, 4-byte output) cover the fallback to the BF16-dequantizing kernel. `Int8TensorCoreLargeSoftmaxScale_CUDA` (`scale = 2^30` with `key_scale = 2^-30`: the `scale <= 1` gate) and `Int8TensorCoreTinyProbabilityLargeValueScale_CUDA` (`value_scale = 2^28` with a candidate of probability about 2^-29: the probability scaling) run prefill and then decode on a two-candidate case with an analytic result. | ##### 2. DynamicSparseAttention: split-K tensor-core decode for the INT8 KV cache | File | Change | |------|--------| | `onnxruntime/contrib_ops/cuda/bert/dynamic_sparse_attention_prefill_int8.cuh` | The kernel body of section 1 with `kSplitK = true`: the grid gets a third dimension, every CTA consumes a contiguous range of the row's 64-candidate tiles and writes an un-normalized partial result (running maximum in the natural-log domain, running sum, value-scaled accumulator) in the layout of the grouped decode kernel. It replaces the CUDA-core `GroupedDecodeDynamicSparseAttentionKernel`, which was issue bound on its per-candidate dot product, warp reduction and exponentials. Decode passes the probability tile as an FP16 high part plus an FP16 residual (about 22 significant bits), because the signed sum of large value codes can cancel to a value far below the rounding of one FP16 probability. A CTA that meets one of the conditions of section 1 evaluates its candidates in FP32 (`AttendRowFp32<true>`) and publishes an FP32 partial in the same layout, so the fallback holds for the splits too. A partial is written as `(acc * 2^-15) * value_scale`, so the merge kernel is unchanged. | | `onnxruntime/contrib_ops/cuda/bert/dynamic_sparse_attention_impl.cu` | `GetInt8TcDecodeSplitCount`: the smallest power-of-two number of 64-candidate tiles per split for which all `ceil(max_selected / 64) * rows * kv_heads` tiles fit one wave of resident CTAs (SM count x `cudaOccupancyMaxActiveBlocksPerMultiprocessor` of the kernel's shared-memory and register footprint: 396 on an H200). The split count is `ceil(tiles / tiles_per_split)`, at most 32 and never more than the workspace was sized for. A positive `ep.cuda.dsa_decode_int8_tc_tiles_per_split` replaces the automatic value; it is a target, because the cap, the workspace size and the even distribution of a row's tiles still apply. The partial vectors come first in the workspace so that they stay 16-byte aligned. `MergeGroupedDecodeDynamicSparseAttentionKernel` combines the partials unchanged. Routing on top of `UseGroupedDecodeAttention` is listed under "Gates and switches"; everything else keeps the CUDA-core kernel. | | `onnxruntime/test/contrib_ops/dynamic_sparse_attention_op_test.cc` | `Int8CacheDecodeTensorCoreKernel_CUDA`: 1, 3 and 8 query rows x 8, 12, 24 and 32 query heads (group sizes 4, 6, 12 and 16, padded to the 16-row MMA tile) x selected candidates {1, 65, 509} around the 64-candidate tile x invalid indices, with the automatic split count, 1 and 3 tiles per split (a split without tiles publishes an empty partial) and the kernel switched off. `Int8TensorCoreDecodeSignedValueCancellation_CUDA` (probabilities 0.5 +- 2^-14 with values +-4064: a single FP16 tile would cancel them to 0) and `Int8TensorCoreDecodeTilesPerSplitCaps_CUDA` (lists of 2048, 2112 and 8192 candidates, that is 32, 33 and 128 tiles, with the automatic split count and 1, 2, 3 and 64 tiles per split, 1 and 8 rows: whatever number of splits results, every candidate reaches the merge; 8192 candidates with 1 tile per split get 32 splits of 4 tiles). `Int8DecodeQueryBeyondHalfRange_CUDA` multiplies every query by 2^28 so that every CTA takes the FP32 path for any split count, plus a case where only some CTAs do. `DecodeSessionOptions_CUDA` checks the switch, the tuning knob, the precedence over the environment and the rejection of invalid values. `Int8CacheDecodeMisalignedBuffers_CUDA`: 16-byte-aligned buffers use the tensor-core kernel, misaligned ones fall back to the CUDA-core kernel. The two scale tests of section 1 run the decode kernel as well. | ##### 3. Early PDL trigger in `ProjectedSigmoidGateMulKernel` | File | Change | |------|--------| | `onnxruntime/contrib_ops/cuda/bert/gated_add_impl.cu` | `ProjectedSigmoidGateMulKernel` issues `griddepcontrol.launch_dependents` first (`__CUDA_ARCH__ >= 900`). | The kernel feeds the hyper-connection down kernel, which is launched with PDL and reads the gate-mul output only after `griddepcontrol.wait`; its constant-weight prefetch now overlaps with the gate-mul kernel. There is no effect for plain launches, no change of any result and no switch. The same trigger in the fpA_intB GEMV kernels is part of microsoft#33689; together they shorten an `S = 8` decode cycle by 1.4-1.7% (preliminary, see "Performance"). ##### 4. HyperConnectionLowRankMix: branch-free sigmoid | File | Change | |------|--------| | `onnxruntime/contrib_ops/cuda/bert/hyper_connection_low_rank_mix_impl.cu` | `SigmoidF32` evaluates `e = expf(-fabsf(v))` once and selects the numerator, `(v >= 0 ? 1 : e) / (1 + e)`, instead of branching between `1 / (1 + expf(-v))` and `expf(v) / (1 + expf(v))`. Both branches already pass `-fabsf(v)` to `expf`, so the result is bit-identical. In the prefill pre-mix neighbouring elements have different signs, so nearly every warp executed both sides of the old branch. | ##### 5. Session options (shared by 1 and 2) | File | Change | |------|--------| | `include/onnxruntime/core/session/onnxruntime_session_options_config_keys.h` | Four `ep.cuda.*` keys: `dsa_prefill_int8_tc`, `dsa_prefill_kernel`, `dsa_decode_int8_tc` and `dsa_decode_int8_tc_tiles_per_split`. The descriptions state the opt-in defaults and the exact `"0"` / `"1"` grammar, and that tiles per split is a target subject to the split cap, the workspace size and the balancing of the row's tiles. | | `onnxruntime/contrib_ops/cuda/cuda_session_options.h` (new) | `ReadCudaSessionOption` (integer threshold) and `ReadCudaSessionSwitch` (exactly `"0"` or `"1"`: `"00"`, `"01"`, `"-0"` and `"+1"` are rejected): an explicit session option wins, otherwise the environment variable (an empty variable counts as unset), otherwise the default. The whole string is parsed and range-checked, and an invalid value fails kernel construction. microsoft#33689 adds the same file for its own switches (see "Related PRs and review guide"). | ##### 6. Documentation and benchmark | File | Change | |------|--------| | `docs/contrib_ops/cuda/dynamic_sparse_attention.md` | The `ep.cuda.dsa_*` options next to their environment variables with the opt-in defaults, the dispatch-table rows and alignment fallbacks of the INT8-cache prefill and split-K decode kernels, the tiles-per-split rule with its sweep and regret statistics, "Numerics and fallbacks" (the conditions of the FP32 path, what the model does, the measured cost of the path), "Rollout" and the operator-level benchmark. | | `docs/contrib_ops/cuda/hyper_connection_low_rank_mix.md` | The early `griddepcontrol.launch_dependents` trigger of the kernel in front of `HcDownKernel` and the branch-free `SigmoidF32`. | | `onnxruntime/test/python/transformers/benchmark_dynamic_sparse_attention.py` (new) | Operator-level benchmark: a single-node `DynamicSparseAttention` model per case (the shapes of the 12 layers; `--rc-model` reads the scales of a model's node), one session per kernel configuration, the GPU time of the kernels of one `Run` from the CUPTI-based torch profiler, alternating rounds, and the RMS error against an FP64 reference. `--query-std`, `--key-scale-multiplier` and `--value-scale-multiplier` reach the FP32 paths, and `--tiles-per-split` sweeps the decode split rule. | ##### Gates and switches Every switch is available as an environment variable and as a CUDA session option (`ep.cuda.*`, declared in `onnxruntime_session_options_config_keys.h`). It is resolved once when the kernel is created, so a change needs a new session, and an explicit session option takes precedence over the environment variable. The 0/1 switches accept exactly the strings `0` and `1`, the tuning knob a non-negative decimal integer; any other value (including an empty session-option string) fails session initialization. `ORT_DSA_PREFILL_KERNEL` already exists on `main` (read on every call; any non-zero integer enabled it) and now follows the same rules. | Kernel | Used when | Environment variable / session option | |--------|-----------|--------| | INT8-KV DSA prefill | BF16 query / output, INT8 main cache, the prefill kernel path (`sequence_length >= 8`, not the grouped decode path), `head_size = 256`, at most 16 query heads per KV head, selected-only from the main cache (no `local_plus_selected`), no smooth softmax, `scale <= 1`, `max_selected > 0`, SM80+, 16-byte-aligned `present_key`, `present_value` and scales, 8-byte-aligned `output`, enough opt-in shared memory (72,960 B) | `ORT_DSA_PREFILL_INT8_TC=1` / `ep.cuda.dsa_prefill_int8_tc` turns it on (default `0`: the BF16-dequantizing kernel, which has no FP16 range limit); `ORT_DSA_PREFILL_KERNEL=0` / `ep.cuda.dsa_prefill_kernel` turns off the whole tensor-core prefill family (default `1`) | | INT8-KV DSA decode | BF16, INT8 main cache, grouped decode path (`1 <= sequence_length <= 8`, 4 to 16 query heads per KV head, `max_selected > 256`), `head_size = 256`, selected-only from the main cache, no smooth softmax, `scale <= 1`, 16-byte-aligned query, `present_key`, `present_value`, scales and attention workspace, SM80+, enough opt-in shared memory (75,264 B) | `ORT_DSA_DECODE_INT8_TC=1` / `ep.cuda.dsa_decode_int8_tc` turns it on (default `0`: the CUDA-core grouped decode kernel, bit for bit as before); `ORT_DSA_DECODE_INT8_TC_TILES_PER_SPLIT` / `ep.cuda.dsa_decode_int8_tc_tiles_per_split` (default 0 = automatic; a requested tiles-per-split target, see section 2) | | Early PDL trigger | SM 9.0+, PDL successor | none (no effect for plain launches) | | Branch-free sigmoid | always | none (bit-identical) | The tensor-core kernels also require a loaded PTX target of at least SM80. If only older PTX is available (for example a `compute_75` image that the driver JIT-compiles for a newer GPU), their `__CUDA_ARCH__ >= 800` bodies are empty stubs that would silently leave the output unwritten, so the dispatch checks `cudaFuncGetAttributes(...).ptxVersion >= 80` and otherwise uses the generic kernels. ##### Numerics - The branch-free sigmoid is bit-identical to the branchy one for all 2^32 `float` bit patterns (exhaustive check in a standalone harness, same nvcc flags, no fast-math). - The early PDL trigger changes no result (generated tokens identical in 45 / 45 decode runs). - The INT8 tensor-core kernels are close to, but not bit-identical with, the previous kernels: they round `Q'` and the probability tile to FP16 and sum in a different FP32 order. The switches' default `0` keeps the previous kernels bit for bit. Against an FP64 reference (the BF16 rounding of the output alone gives about 1.5e-3 relative RMS), decode is at 1.6-1.7e-3 with both kernels (the tensor-core kernel is up to 3% higher), and prefill at 1.7e-3 against 2.7-4.3e-3 for the BF16-dequantizing kernel (tables under "Performance"). - The cases found in review, and how the kernels handle them: | Case | Handling | Test | |---|---|---| | A finite `Q * key_scale` whose FP16 rounding changes the logits (products 32768 and -32766 round to 32768 and -32768, the gap of 16 becomes 0: 0.5 instead of 1) | The per-head bound `eps` above 0.125 nats sends the CTA to the FP32 path | `Int8TensorCoreQueryFoldingRoundingError_CUDA`, `Int8TensorCoreRandomQueryFolding_CUDA` | | Decode: two probabilities 0.5 +- 2^-14 round to the same FP16 value and the signed values cancel (the output 127/256 = 0.496 became 0) | The decode tile is an FP16 high part plus an FP16 residual (prefill keeps one tile, as the BF16 kernel does with a coarser one) | `Int8TensorCoreDecodeSignedValueCancellation_CUDA` | | Scores of 8290560 and 8290560.5 (a large common logit) converted to the log2 domain before the maximum is subtracted lose their gap to float rounding (0.5 instead of 0.62 at scale 1 and 0.508 at scale 1/16) | Natural-log scores and maximum; only the difference is scaled by log2(e) | `Int8TensorCoreLargeCommonLogit_CUDA` | | `value_scale = 2^-135`: `value_scale * 2^-15` underflows before the accumulator is applied (BF16 subnormal result became 0) | Divide the accumulator by 2^15 first, then multiply by `value_scale` | `Int8TensorCoreTinyValueScale_CUDA` | | `abs(value_scale) > 2^16` with a probability below 2^-29 (subnormal in the FP16 tile) | The tail-depth check decides per CTA, instead of sending every CTA to the FP32 path as soon as one channel exceeds 2^16 | `Int8TensorCoreLargeValueScaleRandom_CUDA` | - Cost of the FP32 path (H200, `benchmark_dynamic_sparse_attention.py`; 24 query / 2 KV heads, 2048 selected candidates, context 8192; attention-kernel time, previous kernel against tensor-core kernel): | Data | Rows | Previous | Tensor core | Tensor core vs previous | |---|---|---|---|---| | Queries like the model's | 128 / 8192 | 0.173 / 7.82 ms | 0.063 / 2.83 ms | 2.7x / 2.8x faster | | Queries 8 times larger (`--query-std 11.2`): every CTA is over the error budget | 128 / 8192 / 8 (decode) | 0.172 ms / 7.82 ms / 73.6 us | 3.68 ms / 135.6 ms / 240 us | 21x / 17x / 3.3x slower | | `abs(value_scale)` 24k to 360k (`--value-scale-multiplier 1.2e7`), logit standard deviation about 2.4 nats: 10 to 12% of the CTAs have a candidate deeper than 20 nats | 128 / 512 / 2048 | 0.172 / 0.577 / 2.288 ms | 3.61 / 3.90 / 6.17 ms | 21x / 6.8x / 2.7x slower | | The same value scales, logit standard deviation about 1.3 nats: no CTA is deeper than 20 nats | 8192 / 8 (decode) | 7.78 ms / 73.6 us | 2.86 ms / 13.3 us | 2.7x / 5.6x faster | A CTA on the FP32 path needs about 3.6 ms for 2048 candidates, so the time of a launch with one wave of CTAs is set by its slowest CTA. The fallback keeps the accuracy (relative RMS 1.7e-3 against FP64 with queries 8 times larger; the previous prefill kernel: 8.5e-3). The first group of rows is the Qwen3.8-Flash-Next case: the largest `value_scale` of its 12 layers is 0.285 (the tail-depth check is never armed) and a Monte Carlo over their `q_norm` weights and key scales gives a per-head `eps` of 0.015 on average and 0.022 at most, against the budget of 0.125, which corresponds to queries about 8 times larger than the model produces. Workloads with such queries, or with `abs(value_scale) > 2^16` and wide logit spreads, should leave the switches at `0`. The FP32 path of a CTA is documented in the new "Numerics and fallbacks" section of `docs/contrib_ops/cuda/dynamic_sparse_attention.md`. - Cost of the numerics fixes on the fast path (H200, the same script; separate processes against the previous head of this PR, `ceed60271c`, not paired; attention-kernel time): prefill at 128 / 2048 / 8192 rows 0.062 / 0.815 / 2.760 ms against 0.063 / 0.832 / 2.830 ms (+1.6% to +2.5%); decode with 8 rows 12.23 us against 13.2-13.3 us (+8-9%) and with 1 row 8.84 us against 9.46 us (+7%). The previous head sent every CTA to the FP32 path when one channel had `abs(value_scale) > 2^16` (8192 rows: 2.76 ms became 136 ms). #### Testing In-tree CUDA 13.0 Release build, `sm_90a`, one H200 (SM 9.0), on this branch (`onnxruntime_provider_test`; this checkout needs `-Donnxruntime_ENABLE_CUDA_EP_INTERNAL_TESTS=OFF` because `gemm_tactic_cache_test.cc` does not link otherwise (`undefined reference to onnxruntime::g_host`), which is unrelated to this PR): | Command (`--gtest_filter=`) | Result | |---|---| | `*DynamicSparseAttention*:*HyperConnection*` | 103 run, **103 passed**, 0 failed (73 `DynamicSparseAttentionTest`, 13 `HyperConnectionLowRankMixTest`, 17 `HyperConnectionOpsTest`); about 100 s | | `*GatedAdd*:*ProjectedSigmoid*:*SigmoidGate*:*GateMul*` | 14 run, **14 passed**, 0 failed (12 `ContribOpGatedAddTest`, including `ProjectedGateBFloat16` and `ProjectedGateK1Regression`, and 2 `EngramOpsTest`) | - The review round adds 7 DSA tests (`Int8TensorCoreQueryFoldingRoundingError_CUDA`, `Int8TensorCoreDecodeSignedValueCancellation_CUDA`, `Int8TensorCoreLargeCommonLogit_CUDA`, `Int8TensorCoreTinyValueScale_CUDA`, `Int8TensorCoreDecodeTilesPerSplitCaps_CUDA`, `Int8TensorCoreLargeValueScaleRandom_CUDA`, `Int8TensorCoreRandomQueryFolding_CUDA`), moves the INT8 reference to FP64 with max-subtracted weights, and tightens the two session-option tests (`"00"`, `"01"`, `"-0"`, `"000001"`, `"+1"`, `" 1"`, `"1 "`, `""`, `"2"`, `"-1"`, `"1x"` and `"true"` are rejected for the switches, and `"01"`, `"-0"` and `"+1"` for the environment variables). - Red-green against the previous head of this PR (`ceed60271c`): the final test binary was run against the provider built from that head (`ORT_UNIT_TEST_MAIN_DYNAMIC_PLUGIN_EP_CONFIG_JSON` selects the library) for the 13 tests that concern the review. 7 fail on it: the five numerics tests `Int8TensorCoreQueryFoldingRoundingError_CUDA`, `Int8TensorCoreDecodeSignedValueCancellation_CUDA`, `Int8TensorCoreLargeCommonLogit_CUDA`, `Int8TensorCoreTinyValueScale_CUDA` and `Int8TensorCoreRandomQueryFolding_CUDA`, and `PrefillSessionOptions_CUDA` and `DecodeSessionOptions_CUDA` (the stricter parser). The other 6 pass on it: `Int8TensorCoreDecodeTilesPerSplitCaps_CUDA` and `Int8TensorCoreLargeValueScaleRandom_CUDA` are coverage tests (the split caps; the tail-depth check against FP64 on random data), and `Int8TensorCoreTinyProbabilityLargeValueScale_CUDA`, `Int8TensorCoreLargeSoftmaxScale_CUDA`, `Int8PrefillQueryBeyondHalfRange_CUDA` and `Int8DecodeQueryBeyondHalfRange_CUDA` are the guards of the earlier rounds. All 73 pass with the final provider. - Under `nsys` the session-option tests launch the intended kernels on the final tree. `PrefillSessionOptions_CUDA`: the split / merge kernels for `dsa_prefill_kernel=0`, `DynamicSparseAttentionPrefillKernel<true>` for `dsa_prefill_int8_tc=0` and `DynamicSparseAttentionInt8TcKernel<false>` for `=1` (two launches of each). `DecodeSessionOptions_CUDA`: `GroupedDecodeDynamicSparseAttentionKernel` for `dsa_decode_int8_tc=0` and `DynamicSparseAttentionInt8TcKernel<true>` for `=1` (three launches of each), each followed by `MergeGroupedDecodeDynamicSparseAttentionKernel`. - `lintrunner` (clang-format, ruff, ruff-format) is clean on all 13 changed files. `dynamic_sparse_attention_impl.cu`, `gated_add_impl.cu` and `hyper_connection_low_rank_mix_impl.cu` also compile for `sm_75`, `sm_80`, `sm_86` and `sm_120` with the flags of the build. - The PTX guards (`cudaFuncGetAttributes(...).ptxVersion >= 80`) have no unit test, because they need an executable that carries only `compute_75` PTX. On the H200 (the driver JIT-compiles the PTX for SM 9.0) such an executable reports `ptxVersion = 75` for `DynamicSparseAttentionPrefillKernel<true>` and `<false>`, `DynamicSparseAttentionInt8TcKernel<false>` and `DynamicSparseAttentionInt8TcKernel<true>`, so dispatch rejects them; built for `sm_90a` the same kernels report `ptxVersion = 90`. - CI: the new DSA tests need a BF16-capable device (SM80+) and skip otherwise. They need 72,960 B (prefill) and 75,264 B (decode) of opt-in shared memory for the tensor-core kernels, which `sm_86` has, so they run the tensor-core paths there. #### Checklist - [x] Tests added/updated and passing - [x] No breaking API changes. By default the INT8-KV attention behaves as before (both tensor-core kernels are opt-in; the CUDA-core decode kernel and the BF16-dequantizing prefill kernel run bit for bit as before). Behavior changes: `ORT_DSA_PREFILL_KERNEL` is now resolved once per kernel and strictly 0/1, and the tensor-core kernels, when enabled, are not bit-identical to the previous kernels - [x] Documentation updated (if applicable) ### Motivation and Context At 8192 tokens per prefill chunk the INT8-KV prefill kernel took 12-13% of the chunk with HMMA at only 2.8% of its instructions: it dequantized every gathered K/V tile to BF16 in shared memory before the BF16 MMAs. In decode (`S = 8`, 12 query heads per KV head) the CUDA-core INT8 kernel took 70.3 us plus 5.3 us for the merge per layer at 2K to 32K context, 6.9% of a decode cycle (ncu: issue and shared-load bound on the per-candidate serial chain). In the prefill pre-mix of the hyper-connections, the data-dependent branch of `SigmoidF32` made nearly every warp execute both of its paths. #### Performance on NVIDIA H200 (SM 9.0a, CUDA 13.0) ##### Operator level (final revision) `benchmark_dynamic_sparse_attention.py` builds a single-node `DynamicSparseAttention` model per case (24 query / 2 KV heads, head size 256, INT8 cache, scale 1/16, 2048 selected candidates per row: the shapes of the 12 layers) and one session per kernel configuration, because the `ep.cuda.dsa_*` options are read when the kernel is created. The time is the GPU time of the kernels that one `Run` launches (CUPTI-based torch profiler, so host launch overhead does not enter): "attention" sums the `DynamicSparseAttention` kernels including the decode merge kernel, "node" all GPU activity of the node (cache append, query preparation, ...). Each round measures every configuration once, the order alternates between rounds, the tables show the median of 7 rounds, and the speedup is the median of the paired ratios of the rounds. 3 warmup runs per measurement, 200 profiled runs per round for decode and 5 for prefill. The error is the RMS error of the BF16 output relative to the RMS of an FP64 reference over 64 evenly spaced rows. The data are synthetic: BF16 queries with standard deviation 1.4 (the queries of the model after `q_norm` have a per-head RMS of 1.3 to 1.6), INT8 cache codes with standard deviation 40, and key and value scales drawn from the ranges of that model. "Previous" is the kernel that the switch replaces (the CUDA-core grouped decode, the BF16-dequantizing prefill): ```bash python benchmark_dynamic_sparse_attention.py --rows 1 8 --context 2048 8192 32768 --tiles-per-split 0 1 2 4 8 16 python benchmark_dynamic_sparse_attention.py --rows 128 2048 8192 --context 8192 32768 ``` Decode (`sequence_length <= 8`): | Rows | Context | Attention, previous | Attention, tensor core | Speedup | Node, previous | Node, tensor core | RMS error, previous / tensor core | |---|---|---|---|---|---|---|---| | 1 | 2048 | 31.1 us | 9.5 us | 3.30x | 34.1 us | 12.5 us | 1.59e-3 / 1.62e-3 | | 1 | 8192 | 31.2 us | 9.5 us | 3.30x | 34.2 us | 12.5 us | 1.55e-3 / 1.57e-3 | | 1 | 32768 | 31.2 us | 9.5 us | 3.30x | 34.1 us | 12.5 us | 1.67e-3 / 1.72e-3 | | 8 | 2048 | 73.6 us | 13.2 us | 5.58x | 76.7 us | 16.5 us | 1.64e-3 / 1.69e-3 | | 8 | 8192 | 73.6 us | 13.3 us | 5.57x | 76.8 us | 16.6 us | 1.62e-3 / 1.66e-3 | | 8 | 32768 | 73.7 us | 13.2 us | 5.57x | 77.0 us | 16.6 us | 1.68e-3 / 1.73e-3 | Prefill: | Rows | Context | Attention, previous | Attention, tensor core | Speedup | Node, previous | Node, tensor core | RMS error, previous / tensor core | |---|---|---|---|---|---|---|---| | 128 | 8192 | 0.173 ms | 0.063 ms | 2.72x | 0.179 ms | 0.070 ms | 4.19e-3 / 1.71e-3 | | 128 | 32768 | 0.191 ms | 0.064 ms | 2.99x | 0.199 ms | 0.072 ms | 4.25e-3 / 1.70e-3 | | 2048 | 8192 | 2.287 ms | 0.832 ms | 2.75x | 2.350 ms | 0.894 ms | 4.25e-3 / 1.71e-3 | | 2048 | 32768 | 2.294 ms | 0.832 ms | 2.75x | 2.356 ms | 0.895 ms | 4.21e-3 / 1.70e-3 | | 8192 | 8192 | 7.818 ms | 2.830 ms | 2.76x | 8.057 ms | 3.069 ms | 2.69e-3 / 1.66e-3 | | 8192 | 32768 | 8.877 ms | 3.200 ms | 2.77x | 9.116 ms | 3.439 ms | 4.29e-3 / 1.71e-3 | The decode time does not depend on the context length (the selected candidates are gathered). The error includes the rounding of the BF16 output, whose floor is about 1.5e-3 relative to the RMS: both decode kernels and the tensor-core prefill are near it, and the BF16-dequantizing prefill, which rounds the dequantized keys and values and the probabilities to BF16, has 1.6 to 2.5 times the error of the tensor-core prefill. Decode time against the tiles per split (8 rows, 2048 candidates, kernel plus merge, `--tiles-per-split`): | Tiles per split | 1 | 2 (automatic) | 4 | 8 | 16 | |---|---|---|---|---|---| | Splits, CTAs | 32, 512 | 16, 256 | 8, 128 | 4, 64 | 2, 32 | | Time | 18.8 us | 13.2 us | 12.7 us | 17.7 us | 29.2 us | The automatic rule is a heuristic, not the optimum. Over 20 shapes on the H200 (1, 2, 4 and 8 rows, 512 to 8192 candidates, context 8192, each swept over 1, 2, 3, 4, 6, 8 and 16 tiles per split) it is on average 1.8% and at most 14.8% slower than the best swept value (4 rows and 2048 candidates: 11.9 us against 10.4 us); the previous target of about 384 CTAs is 2.7% and 14.2% slower on the same points. It matters when fewer CTAs are resident: with the kernel's shared memory padded on the H200 so that one or two CTAs fit an SM (132 or 264 resident CTAs; an ad hoc harness that is not part of this change) the rule is on average 0.4% (at most 2.4%) and 3.1% (13.4%) slower than the best swept value, where the 384-CTA target is 23.8% (44.9%) and 14.4% (38.4%) slower. An A100 (108 SMs, two resident CTAs per SM by the 75,264 B of shared memory) is expected to choose 4 tiles per split, 8 splits, for 8 rows and 2048 candidates, the faster setting in the A100 measurement reported in the review (22.1 us against 25.9 us at 16 splits). That choice is derived from the rule, not measured: the A100 has not been run with this script. `ORT_DSA_DECODE_INT8_TC_TILES_PER_SPLIT` tunes other devices. The other kernels of this PR, per launch (H200, `nsys` / CUDA events; unchanged since the earlier revision): | Kernel | Before | After | |---|---|---| | `PreMixVec8Kernel` per launch (hidden 2560, 4 branches, 8192-token chunk; 96 launches per chunk) | 112.1 us | 82.7 us | | `PreMixVec8Kernel`, the same launch in the model (`nsys`, 8192-token chunk, `all_int4`) | 161.2 us | 119.2 us | | `ScaledSiluKernel` | 7.3 us | 6.1 us | ##### End to end (preliminary) These end-to-end numbers were measured for an earlier revision of the kernels, before the numerics changes of the review, and have not been reproduced for the final revision: the stack that they need (the model packages and runtime, and microsoft#33689 together with this PR on top of a `main` that contains microsoft#32126, microsoft#33650 and microsoft#33714) is not available for a final-head run. Each row is an A/B of consecutive providers on the same packages, one round unless noted (not alternated), without raw timings and without confidence intervals. They show the direction of the effect, not a qualified gain. A shorter cycle is not the same as better user-visible throughput, and accepted tokens per cycle and time per accepted output token were not measured for any of these runs. Measured on the Qwen3.8-Flash-Next packages (`hidden = 2560`, `intermediate = 640`, 512 experts, top-10; `all_int4` is the all-INT4 package, `top16` and `top08` are the two INT3-mixed packages), B = 1, DFlash2 `S = 8` decode. Decode is in ms per cycle. TTFT of the INT8-KV prefill kernel in seconds (one round; decode is unchanged): | Package | 2048 | 8192 | 32768 | |---|---|---|---| | `all_int4` | 0.1985 -> 0.1889 (-4.8%) | 0.7122 -> 0.6524 (-8.4%) | 3.0210 -> 2.7782 (-8.0%) | | `top16` | 0.2155 -> 0.2068 (-4.0%) | 0.7339 -> 0.6714 (-8.5%) | 3.0969 -> 2.8608 (-7.6%) | | `top08` | 0.2216 -> 0.2121 (-4.3%) | 0.7478 -> 0.6768 (-9.5%) | 3.1504 -> 2.8540 (-9.4%) | Decode of the split-K INT8-KV kernel, ms per cycle at 2K / 8K / 32K context: | Package | 2K | 8K | 32K | |---|---|---|---| | `all_int4` | 11.672 -> 11.017 (-5.6%) | 11.962 -> 11.143 (-6.8%) | 12.157 -> 11.391 (-6.3%) | | `top16` | 12.148 -> 11.381 (-6.3%) | 12.353 -> 11.608 (-6.0%) | 12.629 -> 11.953 (-5.4%) | | `top08` | 12.216 -> 11.505 (-5.8%) | 12.474 -> 11.675 (-6.4%) | 12.661 -> 12.082 (-4.6%) | Branch-free sigmoid, TTFT in ms (median of 5 repetitions; 9 / 9 cells faster, first tokens identical). Decode is within +-0.3%, because the fused decode kernels call the function for at most 8 rows. The per-cell gains of 0.3-1.1% are about the size of the repeat noise between repeats of one stack (about 1%); the kernel-level `PreMixVec8Kernel` time above is the stronger evidence: | Package | 2048 | 8192 | 32768 | |---|---|---|---| | `all_int4` | 174.1 -> 173.3 (-0.45%) | 574.7 -> 572.9 (-0.31%) | 2450.1 -> 2439.2 (-0.45%) | | `top16` | 191.3 -> 190.5 (-0.42%) | 596.3 -> 592.1 (-0.70%) | 2560.2 -> 2531.7 (-1.11%) | | `top08` | 196.2 -> 194.7 (-0.79%) | 601.1 -> 597.2 (-0.65%) | 2573.5 -> 2554.2 (-0.75%) | The early PDL trigger of the gate-mul kernel and the one of the fpA_intB GEMV kernels (microsoft#33689) together shorten a cycle by 1.4-1.7% (a joint result of two PRs). For the whole stack that contains both PRs (plus microsoft#32126, microsoft#33650 and the package configuration), the median over 4 rounds against the previous release candidate is -32.8 / -36.0 / -36.9% TTFT at 8192 tokens, -32.3 / -35.2 / -36.1% at 32768 tokens, and -27.6 / -24.6 / -24.2% decode ms per cycle at 2K (`all_int4` / `top16` / `top08`); these combined-stack numbers cannot be attributed to a single PR. #### Accuracy The sigmoid change is bit-exact and the early PDL trigger changes no result. At the operator level (final revision) the error against an FP64 reference is in the tables above: decode 1.6-1.7e-3 with both kernels (the tensor-core kernel is up to 3% higher) and prefill 1.7e-3 against 2.7-4.3e-3 for the BF16-dequantizing kernel, with the BF16 rounding of the output alone at about 1.5e-3. The unit tests compare against an FP64 reference with max-subtracted weights and fix the numerical corners found in review to one BF16 ulp. End to end, the final revision has not been qualified, and that is why both tensor-core kernels are opt-in. The one end-to-end accuracy gate is for an earlier revision. The split-K decode kernel changes the generated tokens (a different FP32 accumulation order; for 263 of 800 `all_int4` MMLU-Pro prompts), so it was gated with MMLU-Pro: 800 questions x 3 packages = 2400 items, greedy decoding through the full engine, `max_new = 8192`. Together with the dequantize + cuBLAS route of microsoft#33689, it gave 668 / 678 / 667 correct answers against 665 / 676 / 670 for the previous stack, pooled **+0.08 pp** (95% CI -0.60 .. +0.77 pp, exact McNemar p = 0.91); no package differs at the 5% level. The first-round stack, which contains the INT8-KV prefill kernel, was -0.67 pp pooled against the previous stack (95% CI -1.36 .. +0.03 pp, p = 0.076). These runs measure the accumulated effect of all kernels of the stack, not of one PR, and a drift of up to about 1.4 pp cannot be excluded with 2400 items; a non-significant difference does not establish non-inferiority. They predate the numerics of the final revision (the natural-log domain, the `Q'` rounding bound, the decode probability residual, the value-scale ordering and the tail-depth check), and so does a later check of commit 8 (`engine_behaviours.py` 27 / 27 on the three packages; TTFT and decode within the noise against the stack before commit 8). They do not cover the final revision. Making a switch the default needs the final head qualified with the rest of the stack held fixed and prefill and decode toggled separately: per-package paired accuracy with confidence intervals against a stated degradation bound, including long contexts, exact build and model identities, and for speculative decode the accepted tokens per cycle and the time per accepted token. Packages that qualify the kernels enable them in their session options (`ep.cuda.dsa_prefill_int8_tc` and `ep.cuda.dsa_decode_int8_tc` set to `1`). #### Related PRs and review guide - microsoft#33689 is the sibling PR (tensor-core `MatMulNBits` / fpA_intB and QMoE / MoE GEMM kernels). The changes are divided by operator and have no kernel in common. Both PRs add `cuda_session_options.h`, and the two versions now differ: this PR's is a superset. It adds `ResolveCudaSessionOptionText`, `ResolveCudaSessionSwitch` (exactly `"0"` or `"1"`) and `ReadCudaSessionSwitch`, and `ReadCudaSessionOption` (the integer parser, with unchanged behavior) now calls the text helper. A trial merge of the two heads (`git merge-tree`) has this add/add conflict in that one file and no other, and the two edits of `onnxruntime_session_options_config_keys.h` are in different places. Whichever PR lands second takes this PR's version of the file; the switches of microsoft#33689 can then adopt `ReadCudaSessionSwitch`, which is a change in microsoft#33689. - microsoft#32126 (SwiGLU operator and fusion), microsoft#33650 (RouterGateUpProj) and microsoft#33714 (QMoE INT3 prefill tile) have merged into `main` since this PR was opened. This branch does not touch their files, and a trial merge of this head with the current `main` (`git merge-tree`) is clean. The decode graph measured above contains all three. - The PR has 13 files and 14 commits: commit 1 (early PDL trigger of the gate-mul kernel), commit 2 (INT8-KV prefill kernel), commit 3 (review follow-up: the FP32 fallback for out-of-range queries, the `ep.cuda.dsa_*` session options with their tests), commit 4 (fixes MSVC narrowing warnings in the DSA tests), commit 5 (split-K decode kernel and its options and tests), commit 6 (branch-free sigmoid), commit 7 (docs), commit 8 (review follow-up: the `scale <= 1` condition, the 2^15 probability scaling, `ORT_NO_EXCEPTIONS` guards in the option tests and two regression tests), commit 9 (`ceed60271c`, value-scale gate with an FP32 fallback for all CTAs, replaced by the tail-depth check of commit 10), commit 10 (`1bf2eb17cb`, review fixes of the kernels, the split count, the options and the opt-in defaults), commit 11 (`bd24bb8503`, tests), commit 12 (`2dedd0c4ff`, docs and the benchmark script), commit 13 (`e8b3ad2232`, the end-to-end timings in the hyper-connection doc are labeled preliminary) and commit 14 (`d60b0f4b63`, the DSA doc states that the end-to-end timings have no acceptance measurements). `dynamic_sparse_attention_prefill_int8.cuh` is the main file to review (812 lines). - The review threads of microsoft#33689 on the prefill kernel and its test (clamping of the scaled query, identical K/V rows in the sweep) were answered there with the fixes of commit 3; its other threads concern the QMoE and dense INT8 kernels that stay in microsoft#33689. The review threads of this PR on the numerics (the rounding of the folded query, the decode probability rounding, the log2 domain, the tiny value scale and the value-scale cliff), the tests (FP64 reference, active splits, exact checks), the split target, the switch grammar, the tiles-per-split wording and the rollout and end-to-end claims are answered by commits 10 to 14. --------- Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

…, adaptive promotion, and DMA prefill (microsoft#33452) ### Description Stacked on microsoft#33451, rebased onto the updated `main` + microsoft#33450 + microsoft#33451 stack and structured into 2 commits: 1. **Commit 1 (`18928e7e46`) — QMoE Offload Kernels, SpinPool, Callback Timing, and `QMoE` v4 Admission Hook**: - Adds `LaunchQMoEOffloadClassify`, `LaunchQMoEOffloadGather`, `LaunchQMoEOffloadScatter`, and `LaunchQMoEOffloadMergeFinalizeCombine` in `qmoe_kernels.{h,cu}` and raw NVFP4 GEMV decode support in `moe_gemv_fp4.cu`. - Adds `SpinPool` (`qmoe_spin_pool.h`), callback timing (`qmoe_callback_timing.h`), and v4 `QMoE` offload admission hooks in `moe_quantization.{h,cc}`. 2. **Commit 2 (`9f2abe75e9`) — Portable `QMoEOffload` Backend, Adaptive Promotion, and v4 Offload Tests**: - Adds `QMoEOffload` (`qmoe_offload.{h,cc}`), `test_qmoe_offload_v4.py`, and `docs/contrib_ops/cuda/moe_qmoe.md`. --------- Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

tianleiwu and others added 30 commits October 29, 2026 13:39

Merge remote-tracking branch 'origin/master' into sync_msft_02112026

a75d4aa

ai-fw-intg requested review from Jaswanth51, ankitm3k, jatinwadhwa921 and vthaniel November 1, 2026 20:39

Navigation Menu

Repository navigation

Sync with Microsoft ONNX Runtime - 02112026 - #1379

Sync with Microsoft ONNX Runtime - 02112026#1379
ai-fw-intg wants to merge 31 commits into
ovep-developfrom
sync_msft_02112026

ai-fw-intg commented Nov 1, 2026

Uh oh!

This branch has not been deployed

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

15 participants

Conversation

ai-fw-intg commented Nov 1, 2026

Uh oh!

This branch has not been deployed

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

15 participants