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

ai-fw-intg · 2026-10-28T20:39:33Z

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

…de (microsoft#32951) # Honor explicit rotary positions in CUDA GroupQueryAttention fast decode ### Description Exclude explicit rotary positions from the CUDA GroupQueryAttention fused FlashDecoding path when rotary embedding is enabled. Those cases use the existing position-aware PrepareQKV path followed by FlashAttention. The guard is applied before fast-decode scratch sizing and sequence-length preparation. It preserves existing eligibility for implicit positions and non-rotary inputs, and adds no device-to-host position comparison. Add deterministic operator regression tests with an independent scalar reference, shared-cache alias checks, and bit-exact K/V cache assertions. Coverage includes shifted, unshifted and implicit positions, batch sizes one and two, partial/full rotary dimensions, split/interleaved pairing, and three single-batch multi-token controls. ### Motivation and Context The fused rotary path derives rotary-table indices from cache offsets and does not forward supplied explicit position IDs. Explicit positions can differ from cache append offsets, producing incorrect attention and appended keys. The existing PrepareQKV path honors those positions. ### Reproduction The commands below reconstruct the tested source configuration, not current main. They are portable adaptations of saved successful commands, checked against source and records but **not rerun during PR publication**. No model, saved prompts or token vectors are needed. Prerequisites: Linux, an assigned A100/SM80 GPU, GCC 13.2, CMake 4.1.2, GNU Make, CUDA 13.0.48, cuDNN 9.12, and an existing build-capable Python environment (the saved CMake cache reports Python 3.12.9). Provision dependencies separately: the commands do not fetch, install or download. Supply the 17 extracted source trees listed below, using the URLs/checksums in [`cmake/deps.txt` at the validated base](https://github.com/microsoft/onnxruntime/blob/8633745c6844d76cef5b811e305a2cb8ce5c24a1/cmake/deps.txt). Those pins match the saved build. Neither the local dependency bundle nor compiled binaries are included in this PR. <details> <summary>Prepare identical regression tests and build both variants (Bash)</summary> Set the variables checked below to absolute paths, except the job count. `ORT_REPO` must already contain the pinned candidate commit below and its parent. `REPRO_ROOT` must be a new, nonexistent directory under an existing parent. `DEPS_ROOT` contains `<name>-src` directories. Use the same dependencies and toolchain for both builds; the historical build used eight parallel jobs. ```bash set -euo pipefail : "${ORT_REPO:?existing local ONNX Runtime repository}" CANDIDATE_COMMIT=c7557c9a1061eb5ab04301e65d2f356d018469b4 : "${REPRO_ROOT:?new absolute reproduction directory}" : "${DEPS_ROOT:?absolute directory of pre-extracted pinned sources}" : "${CUDA_HOME:?CUDA toolkit root}" : "${CUDNN_HOME:?cuDNN root}" : "${CC:?absolute GCC C compiler path}" : "${CXX:?absolute GCC C++ compiler path}" : "${BUILD_JOBS:?positive bounded parallelism appropriate for your host}" [[ "$BUILD_JOBS" =~ ^[1-9][0-9]*$ ]] || { echo "BUILD_JOBS must be positive" >&2; exit 1; } BASE=8633745c6844d76cef5b811e305a2cb8ce5c24a1 test "$(git -C "$ORT_REPO" rev-parse "$CANDIDATE_COMMIT^")" = "$BASE" mkdir -- "$REPRO_ROOT" git -C "$ORT_REPO" worktree add --detach "$REPRO_ROOT/candidate-src" "$CANDIDATE_COMMIT" git -C "$ORT_REPO" worktree add --detach "$REPRO_ROOT/baseline-src" "$BASE" TEST_FILE=onnxruntime/test/contrib_ops/group_query_attention_position_test.cc test ! -e "$REPRO_ROOT/baseline-src/$TEST_FILE" cp -- "$REPRO_ROOT/candidate-src/$TEST_FILE" "$REPRO_ROOT/baseline-src/$TEST_FILE" cmp -- "$REPRO_ROOT/candidate-src/$TEST_FILE" "$REPRO_ROOT/baseline-src/$TEST_FILE" deps=(abseil_cpp cudnn_frontend cutlass date eigen3 flatbuffers googletest gsl mp11 nlohmann_json onnx onnxruntime_curl onnxruntime_mbedtls protobuf pytorch_cpuinfo re2 safeint) dep_args=() for dep in "${deps[@]}"; do test -d "$DEPS_ROOT/$dep-src" || { echo "Missing source: $dep-src" >&2; exit 1; } dep_args+=("-DFETCHCONTENT_SOURCE_DIR_${dep^^}=$DEPS_ROOT/$dep-src") done for variant in baseline candidate; do CUDA_VISIBLE_DEVICES='' cmake \ -S "$REPRO_ROOT/$variant-src/cmake" -B "$REPRO_ROOT/$variant-build" \ -G "Unix Makefiles" \ -DCMAKE_BUILD_TYPE=RelWithDebInfo -DCMAKE_CUDA_ARCHITECTURES=80 \ -DCMAKE_C_COMPILER="$CC" -DCMAKE_CXX_COMPILER="$CXX" \ -DCMAKE_CUDA_COMPILER="$CUDA_HOME/bin/nvcc" \ -Donnxruntime_USE_CUDA=ON -Donnxruntime_CUDA_HOME="$CUDA_HOME" \ -Donnxruntime_CUDNN_HOME="$CUDNN_HOME" \ -Donnxruntime_BUILD_SHARED_LIB=ON -Donnxruntime_BUILD_UNIT_TESTS=ON \ -Donnxruntime_ENABLE_PYTHON=OFF -Donnxruntime_USE_TELEMETRY=OFF \ -Donnxruntime_USE_FLASH_ATTENTION=ON -Donnxruntime_CUDA_MINIMAL=OFF \ -Donnxruntime_BUILD_CUDA_EP_AS_PLUGIN=OFF -Donnxruntime_USE_WEBGPU=OFF \ -Donnxruntime_QUICK_BUILD=OFF -Donnxruntime_RUN_ONNX_TESTS=OFF \ -Donnxruntime_MATERIALIZE_ONNX_NODE_TESTS=OFF \ -Donnxruntime_FLASH_NVCC_THREADS=1 \ -DFETCHCONTENT_FULLY_DISCONNECTED=ON -DFETCHCONTENT_UPDATES_DISCONNECTED=ON \ "${dep_args[@]}" CUDA_VISIBLE_DEVICES='' cmake --build "$REPRO_ROOT/$variant-build" \ --target onnxruntime_provider_test --parallel "$BUILD_JOBS" done ``` The base commit alone has no new regression test. Copying only the test into the baseline leaves its production implementation unfixed. Fresh CMake configuration discovers that test in both source trees; no CMake source edit is required. Do not substitute a quick/minimal build or a different GPU architecture and call it the recorded configuration. </details> <details> <summary>Run ordinary numerical regressions, then nearby tests (Bash)</summary> Continue in the same shell after successful builds. Define `CUDA_DEVICE` as the explicitly assigned GPU UUID/index. Define `CUDA_LIB_DIR` and `CUDNN_LIB_DIR` as the absolute shared-library directories for the installations above (historically `CUDA_HOME/lib64` and `CUDNN_HOME/lib`). Each variant loads its own freshly built provider first. ```bash : "${CUDA_DEVICE:?assigned GPU UUID or index}" : "${CUDA_LIB_DIR:?absolute CUDA shared-library directory}" : "${CUDNN_LIB_DIR:?absolute cuDNN shared-library directory}" common=(--gtest_color=no --gtest_repeat=1 --gtest_shuffle=0 --gtest_stack_trace_depth=0) flash_env=(ORT_DISABLE_FLASH_ATTENTION=0 ORT_DISABLE_FLASH_DECODE=0 ORT_ENABLE_CUDNN_FLASH_ATTENTION=0 ORT_ENABLE_XQA=0) FOCUSED_FILTER='GqaSharedCachePositions/GqaSharedCachePositionTest.MatchesReference/*' run_test() ( build="$REPRO_ROOT/$1-build" report="$2" filter="$3" cd -- "$build" || exit 2 env CUDA_VISIBLE_DEVICES="$CUDA_DEVICE" "${flash_env[@]}" \ LD_LIBRARY_PATH="$build:$CUDA_LIB_DIR:$CUDNN_LIB_DIR" \ ./onnxruntime_provider_test "${common[@]}" \ --gtest_filter="$filter" --gtest_output="xml:$build/$report.xml" ) for variant in baseline candidate; do build="$REPRO_ROOT/$variant-build" (cd -- "$build" env CUDA_VISIBLE_DEVICES='' LD_LIBRARY_PATH="$build:$CUDA_LIB_DIR:$CUDNN_LIB_DIR" \ ./onnxruntime_provider_test "${common[@]}" \ --gtest_filter=GroupQueryAttentionPositionTest.ReferenceSanityCpu \ --gtest_output="xml:$build/reference.xml") done baseline_status=0 run_test baseline positions "$FOCUSED_FILTER" || baseline_status=$? test "$baseline_status" -eq 1 run_test candidate positions "$FOCUSED_FILTER" GQA_FILTER='GroupQueryAttentionTest.CudaCacheAliasingFlash:GroupQueryAttentionTest.CudaCacheAliasingUnfused:GroupQueryAttentionTest.CudaFlashFastDecodeClampsNegativeSeqlensK:GroupQueryAttentionTest.CudaFlashFastDecodeClampsOversizedSeqlensK' ROTARY_FILTER='ContribOpRotaryEmbeddingTest.RotaryEmbedding_Interleaved_SmallData_LlamaMSFT:ContribOpRotaryEmbeddingTest.RotaryEmbedding_NotInterleaved_SmallData_LlamaMSFT:ContribOpRotaryEmbeddingTest.RotaryEmbedding_CustomRotaryDim_SmallData_Phi' run_test candidate nearby-gqa "$GQA_FILTER" run_test candidate nearby-rotary "$ROTARY_FILTER" ``` At this pinned source revision, the focused wildcard selects exactly the same 27 parameterized cases as the saved explicit-name filter. Require **27 tests executed with no skips**, baseline **eight shifted single-token failures and 19 passing controls**, and candidate **27 passes**. Inspect the XML and console results: an exit code alone is not sufficient (including baseline exit 1). The CPU reference check must pass for both variants; nearby invocations select four GQA tests and three rotary tests, respectively. Outputs are `reference.xml` and `positions.xml` in each build directory, plus `nearby-gqa.xml` and `nearby-rotary.xml` in the candidate build directory. These commands test numerical/cache assertions only. They do **not** generate the historical debugger dispatch observations or reproduce its provenance and resource-monitoring controller. The ordinary nearby commands are proposed reproduction commands, not a claim that the stopped historical invocation completed successfully. </details> ### Validation Validated against `8633745c6844d76cef5b811e305a2cb8ce5c24a1` on Linux, NVIDIA A100/SM80, CUDA 13.0.48 and cuDNN 9.12. Baseline and candidate used identical regression tests and matched build configuration. These are saved September 25 results; publication adds no fresh build or GPU validation. | Check | Result | | --- | --- | | Baseline focused matrix | Eight expected shifted-position numerical failures; 19 passing controls | | Candidate focused matrix | 27/27 passed; no skips | | Nearby GQA | Four tests completed successfully | | Nearby standalone RotaryEmbedding | Three tests passed in a separate ordinary native invocation; no skips | The baseline result was accepted by a supplemental saved-evidence audit after correcting a controller mapping-classification error. Its original invalid controller record was retained. The earlier nearby invocation completed four GQA tests, then stopped on an unavailable debugger local in a rotary test; it has no complete aggregate XML. The three rotary tests were subsequently completed without the debugger, with CUDA FP32/FP16 participation verified. This was an observer failure, not an ORT numerical failure. Attention tolerance remains 0.002 and cache assertions remain bit-exact. Actual dispatch observations under forced Flash recorded: - Baseline: 24 FlashDecoding and three FlashAttention cases. - Candidate: 16 explicit-position single-token cases use PrepareQKV/ FlashAttention; eight implicit-position cases retain FlashDecoding; three multi-token controls use FlashAttention. ### Limitations and Follow-ups - This validates the forced-Flash configuration, not default XQA routing. - The affected-path matrix covers FP16, separate Q/K/V, BNSH shared caches and SM80. Broader dtypes, layouts and architectures remain untested. - Multi-token cases are controls on this revision, not evidence of an affected multi-token fused path. - Current-target performance has not been measured. Explicit positions conservatively take the preparation path even when equal to cache offsets. - Results do not establish validation on a newer upstream revision. - LongRoPE cache-history/transition policy remains a separate unresolved issue; this change does not establish full-model parity or assume that re-rotating stored keys alone is sufficient. This is an ORT correctness fix, not a GenAI allocation change; no GenAI memory-saving claim is made. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

### Description  Change CodeQL's Java 17 distribution from Microsoft to Temurin so `actions/setup-java` uses the JDK already provisioned in the runner image. No other workflow behavior is changed. Validated commit `48511836a7c34075624dd9d7d329d4da8c58af79` with a branch workflow dispatch before opening this PR: [CodeQL run 40580423733](https://github.com/microsoft/onnxruntime/actions/runs/40580423733) passed all three language jobs, including Java autobuild and analysis. The [Java job log](https://github.com/microsoft/onnxruntime/actions/runs/40580423733/job/122458014829) reports `Resolved Java 17.0.20+1 from tool-cache`, using `/opt/hostedtoolcache/Java_Temurin-Hotspot_jdk/17.0.20-1/x64` on image `Onnxruntime-Ubuntu-2204-CPU-nvme-full` version `516.0.0`. Java setup did not perform remote version resolution or download a JDK. ### Motivation and Context  Five investigated CodeQL Java runs failed during Microsoft JDK setup and succeeded on retry at the same SHA. Failures included inability to load the Microsoft OpenJDK manifest and connection timeouts during remote resolution. The image already caches Temurin, but the Microsoft distribution request cannot use that cache. Selecting Temurin removes the observed Java manifest/download dependency by using the existing image cache; no runner-image changes are required. Co-authored-by: Copilot App <223556219+Copilot@users.noreply.github.com>

## Summary - Support activation-type GDN gates and fused packed projection handling. - Extend causal convolution state/normalization handling with schema and provider validation. - Keep existing default modes and upstream regression coverage. ## Clean main-based replacement Supersedes microsoft#33461. Reconstructed directly on `main` (`28225b8583474cf6a6e7f97222227ae668f80e8b`), without the old divergent feature-branch history. This PR targets `main` directly and has no dependency on another operator PR. Family documentation is bundled here; the separate old documentation PR microsoft#33465 is retired. ## Review and pipeline fixes - Initialize `mid_state` for empty ragged requests from the incoming state, or zero when absent, in both warp-specialized and generic recurrent kernels. - Add regression coverage for `cu_seqlens={0,0,1}`, head sizes 128 and 32, and both initial-state cases. - Apply the same fused-decay reference and input setup for the `softplus` and `qwen` aliases. The existing first-row-state regression retains the `softplus` alias. - Replace constructor use of `Node::OutputDefs()` with `ComputeContext::OutputCount()` validation. The former API is absent in the WebGPU plugin adapter and caused the Linux static-plugin and Windows plugin compile failures. - Synchronize the operator documentation with the extended schemas and CPU/CUDA/DML registrations. These edits are source-derived; exact generated-document validation remains with CI. - The earlier React Native Android failure was an npm feed download `ETIMEDOUT`, not an operator compilation failure. ## Validation - Used the `qwen-flash` Conda environment. - Ran `lintrunner -a` before committing, followed by a passing check-only run over the full PR delta. - `git diff --check` passed. - Before the request to skip further builds, the changed GDN CUDA and C++ test translation units compiled successfully. An isolated GPU3 harness using the current CUDA kernel bodies passed empty-request state initialization for both warp/generic geometries, with and without initial state. - No full ONNX Runtime or GenAI rebuild was performed for these review fixes. The isolated harness is not an end-to-end operator test run. - No benchmark/profiling/evaluation scripts or local task metadata are included. ### Runtime qualification The earlier accumulated-build failure in `GatedDeltaNetTest.RecurrentOutputsIncludeFirstRowState` compared different recurrences because the reference did not recognize the `softplus` alias. The reference has been corrected. A fresh full native test-suite/model run has not been performed; keeping this PR in draft pending CI and runtime qualification.

…t#33472) ## Summary - Extend DynamicSparseAttention cache and grouped CUDA paths. - Add FP32 key-state support for BF16 QSA, optional mid-state/status outputs and deterministic tiled selection. - Align optional QSA positions across schema, CUDA and WebGPU; explicitly reject CUDA-only modes on WebGPU. - Keep experimental radix/warp-head controls default-off and document their scope. ## Review follow-up and pipeline fixes - Allocate grouped floating-cache decode partial buffers even for a single split; only the generic split path can omit them. - Add FP16/BF16 grouped prefill/decode reference-parity cases with nonzero logits, partial query-head groups, empty/partial selection rows, auxiliary selections/local windows, and sink logits. Include the batch-128, four-head, width-256, max-selected-512 single-split decode regression. Inputs use exactly representable binary values so the reference sees the actual stored values. - Reject connected optional outputs 8-10 in CUDA IndexShare reuse instead of returning with unproduced outputs; document the restriction and test paired mid-state and validation outputs. - Reject CUDA-only committed_length and validation_status on WebGPU DynamicSparseAttention, with model-level rejection tests. - Count optional QSA position_ids only when supplied, retaining strict checks for required bindings. Existing no-position regressions cover the eight-binding case; add explicit-position tests for nine-binding success and eight-binding rejection. - Infer state_update activation type from past_kv_buffer, not FP32 past_key_state. Add BF16-activation/FP32-state capture coverage with raw keys and compression-block transitions. - Fix Windows C4244 warning-as-error failures in the test reference: preserve the int64 KV-head index and explicitly convert INT8 test data. - Regenerate the affected contrib/operator-kernel documentation using official renderers with source-derived CPU/CUDA/DML metadata, including current signatures, T_KEY, and optional-output descriptions. ## Clean main-based replacement Supersedes microsoft#33462. Reconstructed directly on `main` (`28225b8583474cf6a6e7f97222227ae668f80e8b`), without the old divergent feature-branch history. This PR targets `main` directly and has no dependency on another operator PR. Family documentation is bundled here; the separate old documentation PR microsoft#33465 is retired. ## Validation - Used the `qwen-flash` Conda environment. - Ran `lintrunner -a` before committing, followed by a passing check-only run on the full PR delta, C++/CUDA formatting checks, and source whitespace checks. - Verified affected generated sections/rows against official renderer output using source-derived metadata. Exact renderer-generated paragraph/EOF whitespace is retained; full native `--gen_doc validate` remains a CI check. - Per request, no ORT/GenAI rebuild or new native/GPU/WebGPU test execution was performed for this follow-up. Execution is left to CI. - Original-branch schema syntax and affected CUDA/host/test translation-unit compilation (SM90) are historical evidence, not validation of this follow-up. - Inspected Windows CUDA/plugin/TensorRT logs identify the C4244 failures fixed above; macOS WebGPU reports eight QSA failures from the missing optional-position binding. React Native Android npm ETIMEDOUT and CodeQL Java Gradle-download timeout are infrastructure failures, not source errors. - The new push starts fresh CI; all pipelines are not yet claimed green. Full main-based model/GPU requalification remains outstanding; keeping this PR in draft. - No benchmark/profiling/evaluation scripts or local task metadata are included. ### Runtime qualification notes Earlier accumulated-build testing excluded malformed-state expectations `QsaMalformedStateLengthsAreClamped` and `CsaMalformedStateLengthsAreClamped`; their upstream-baseline status is not established. Those tests are not disabled by this follow-up. The current WebGPU QSA log fails before those numerical assertions due to the optional-position binding bug fixed here. The output-count regression expectation is updated for the intentional optional-output extension to 11 slots. WebGPU changes have not been built or executed locally.

### Description Support empty packed weights in CPU and WebGPU `MatMulNBits`. With `weight_prepacked=0`, one or more dimensions of B's packed shape may be zero; the other dimensions must retain their normal values. The output keeps the logical shape determined by A, K and N, and contains zeros or the broadcast bias. Unused quantization parameters may also be empty. CPU skips unused parameter packing and matrix multiplication; WebGPU records a fill dispatch so graph replay rewrites the output correctly. QKV/MLP fusion remains enabled: when a projection has empty weights, the fused operator uses the shared zero/bias fallback while computing the nonempty projections. Includes operator documentation and focused tests for FP16/FP32, 2/4/8-bit weights, constant/runtime inputs, bias, malformed shapes, empty/nonempty overrides, and graph capture. ### Motivation and Context Optional projections, such as disabled LoRA adapters, should not require full-size zero weight tensors. Zero-sized defaults let the same raw `MatMulNBits` graph run on CPU and WebGPU, with nonempty weights supplied when needed. Split from microsoft#33370. That PR retains these changes until this PR merges; QLoRA-specific operators, fusion and shaders remain there. ### Validation Built the standalone branch against main with built-in WebGPU on Windows x64. All 16 empty-weight tests, 28 existing QKV/MLP fusion regressions, and 8 existing nonempty-weight regressions passed; 4 Arm64-specific tests were skipped. Coverage includes fused empty/nonempty overrides, MLP bias, residual outputs, and capture replay. FP16 WebGPU cases require confirmed ShaderF16 support. Lint and `git diff --check` passed.

…crosoft#33531) ### Description  Restore `adapter_info_.nextInChain = &subgroup_matrix_configs_` for WASM builds in `WebGpuContext::Initialize`. microsoft#33055 moved the ` #endif of the !defined(__wasm__)` block past this assignment. As a result, on WASM `subgroup_matrix_configs_` is never chained into `GetAdapterInfo()`. All subgroup matrix kernels will fall back to non subgroup matrix kernels. ### Motivation and Context  see above Signed-off-by: Zhenwei Jin <zhenwei.jin@intel.com>

### Description Add `103-real` to the Linux aarch64 CUDA 13 architecture lists used by: - Python GPU wheel packaging - CUDA plugin packaging - the Linux Python build-script fallback The existing CUDA architecture normalization maps `103-real` to `103a-real`. Other platforms and CUDA architecture lists remain unchanged. ### Motivation and Context NVIDIA GB300 reports compute capability 10.3. While testing on WSL, the current Linux aarch64 CUDA 13 package list omits SM103, so an ONNX Runtime GPU package can import and advertise `CUDAExecutionProvider` but fail its first CUDA kernel on GB300 with: ```text CUDA error cudaErrorNoKernelImageForDevice:no kernel image is available for execution on the device ``` This follows the Linux aarch64 architecture-list approach established by microsoft#33315 for SM100. The existing `onnxruntime_test_python_cuda_kernel.py` test covers the provider-execution failure mode, so this change reuses that coverage rather than adding a duplicate test. ### Testing A local ONNX Runtime 1.30.0 Linux aarch64 wheel built with `CMAKE_CUDA_ARCHITECTURES=103-real;120-real`: - passed wheel archive validation - contained both `sm_103a` and `sm_120a` cubins - loaded the CUDA provider without unresolved internal symbols Signed-off-by: Haoxi Zhang <haoxiz@nvidia.com>

### Description Optimize WebGPU DP4A activation quantization and matrix tiling for wide Qualcomm subgroups. - Reduce each lane's four absolute activation components locally before XOR shuffles, preserving independent 32-lane partitions for each 128-value quantization block. - Use a 128x32 output tile with two explicitly requested 128-lane subgroups on compatible Qualcomm adapters. This route requires 4-bit weights, block size 32, and M >= 128. - Reuse each activation scale across its four 32-value K tiles, retaining both workgroup barriers. - Add regression coverage for independent quantization blocks, scale reuse across K tiles, partial workgroups, matrix tile tails, batching, bias, and zero points. ### Motivation and Context Adreno X1-85 exposes 64-128-lane subgroups. Local component reduction reduces quantization shuffle traffic; the matrix tile assigns one activation row and 16 output columns to each lane. A fresh source-built ONNX Runtime GenAI native benchmark comparison on Windows ARM64 / Adreno X1-85 (driver 31.0.170.0), Aion batch 1, and exactly 1,024 prompt tokens measured: | KV-cache capacity | Baseline prefill | PR prefill | Baseline TPS | PR TPS | Latency reduction | TPS gain | |---|---:|---:|---:|---:|---:|---:| | 8,192 | 3.76782 s | 3.63298 s | 271.78 | 281.86 | 3.58% | 3.71% | | 1,025 | 3.77155 s | 3.61767 s | 271.51 | 283.06 | 4.08% | 4.25% | Both ORT variants were rebuilt with identical settings and used the same source-built GenAI executable, model and prompt. Runs used D3D12, generator reuse, graph capture disabled, one generated token, two warmups and five measured iterations per process. Reordered runs collected 20 measured prefills per version at capacity 8,192 and 10 at capacity 1,025. Loaded ORT DLL paths and hashes were verified. GenAI times prompt processing separately from first-token sampling; timing scopes and loops were unchanged. The local GenAI integration patch only adjusts external-QNN registration and output flushing. Earlier GPU profiling measured activation quantization at 244.4 ms before and 82.9 ms after (approximately 3x faster), and exploratory matrix-stage timing at 5.586 s versus 5.342 s. These stage profiles came from an earlier session and are not a decomposition of the fresh GenAI results. Performance evidence is limited to this device and workload. The performance measurements above precede the activation-scale reuse review follow-up. No additional end-to-end gain is claimed for that follow-up. ### Validation - Fresh native Windows ARM64 Release build with WebGPU/D3D12 succeeded. - Six focused provider tests passed after the review follow-up: `WebGpuMatMulNBitsTest.DP4AActivationScaleReuse`, `MatMul2Bits.Float32_2b_Accuracy4`, `WebGpuMatMulNBitsTest.DP4AQualcommWave128Tiles`, `WebGpuMatMulNBitsTest.DP4AQuantizationLargeSubgroups`, `MatMulNBits.Float32_4b_Accuracy4`, and `MatMulNBits.Float32_4b_Accuracy4_Batch`. - Fresh GenAI prefill logits and sampled KV-cache values were bit-identical between the rebuilt variants at capacity 8,192. - Aion prefill and one-token decode logits were bit-identical to the original build at 128 and 1,024 input tokens. - C++ lintrunner and whitespace checks passed. Three broader FP16 MatMulNBits cases at M=100, N=64, K=32 fail on both the original build and this branch through a different kernel. Browser WASM and other GPU vendors have not been benchmarked.

…osoft#32276) ### Description - Remove NvTensorRTRTXExecutionProvider source files/ references in operator tests/macros from the ORT repo- - For perftest, remove cudart dependency to enable use of CPU build with TRT RTX plugin EP. - AppendExecutionProvider with legacy EP reports deprecation error with the new plugin EP path ### Motivation and Context  https://github.com/NVIDIA/TensorRT-RTX-EP-ABI plugin EP is the preferred path for TRT RTX EP. --------- Co-authored-by: Copilot Autofix powered by AI <175728472+Copilot@users.noreply.github.com> Co-authored-by: Tianlei Wu <tlwu@microsoft.com> Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

### Description - Set the minimum Pillow version to 13.0.0 in the transformer tooling requirements. - Add the same explicit minimum to the transformer test requirements, including Pillow installed through torchvision. ### Motivation and Context Keep the Pillow requirement consistent across transformer tooling and tests. The minimum is compatible with ONNX Runtime's Python 3.11 minimum and torchvision 0.25.0. ### Validation - Verified both requirement specifiers enforce the Pillow 13.0.0 minimum. - Resolved Pillow 13.0.0 and torchvision 0.25.0 wheels on Windows with Python 3.14. - Verified torchvision 0.25.0's declared Pillow requirement accepts 13.0.0. - Passed PNG, JPEG, TIFF, and morphology smoke checks with Pillow 13.0.0. - `git diff --cached --check` passed.

…ore and add GatedRMSNormWarp128Kernel (microsoft#33601) ## Description Optimizes `GatedDeltaNetChunkedKernel` (and shared `gated_delta_net_mma.cuh` helpers) and `GatedRMSNorm` for Qwen3.5 / Qwen3.8 linear-attention prefill and speculative decode on SM80+/SM90 GPUs: 1. **Hardware `ldmatrix` & Direct Shared-Memory Epilogues in `gated_delta_net_mma.cuh`**: - Replaces scalar 16-bit shared-memory loads + `PackHalf2` in `LoadFragA<Transposed>` and `LoadFragB<Transposed>` with `ldmatrix.sync.aligned.m8n8.{x4,x2}{.trans}.shared.b16` on `__CUDA_ARCH__ >= 800`. - Adds 64-bit/32-bit `StoreC2` (`float2` / `PackHalf2`), `SyncBeforeStore`, and `IdentityMinus` to `SmemGemm`, and vectorizes 3-slice `S -> s_h` staging in `SmemGemmFp32State` with `float2` loads and `PackHalf2` stores. - Stages the 16-element forward-substitution column in `BuildTriInverse` in unused `tf` shared memory instead of a thread-local `float col[16]` array (`STACK:64 -> STACK:0`) and writes Horner steps (`I - nb @ ti`) and `ti @ db` directly to `ti`. 2. **Barrier & Shared-Memory Round-Trip Elimination + `uint4` Vectorized Global I/O in `GatedDeltaNetChunkedKernel` (`gated_delta_net_impl.cu`)**: - Vectorizes global `key`/`query`/`value` loads and `output` stores with 16-byte `uint4` transactions. The vector paths are taken only when the bound buffers are 16-byte aligned (checked at run time, so offset IOBinding views use the existing scalar paths), and the dynamic shared-memory base is `__align__(16)`. - Passes `c_scale = __expf(g_total)` directly to `SmemGemm<DK, kDVB, BT, true, false, true>(s.s_f, ...)`, eliminating the 8,192-element `s.s_f *= decay_all` shared-memory loop and barrier. - Writes `s.u_h = s.ti_h @ s.u_h` directly via `SyncBeforeStore=true`, caches `qk_l2_norm` inputs in registers, overlaps `InclusiveScanBT` before the barrier, and vectorizes `s.v_h -> s.u_h` and `s.k_h`/`s.q_h` decay scaling in pairs, eliminating 6 `__syncthreads()` barriers and 5 shared-memory matrix round-trips per chunk. 3. **4-Row/CTA Warp-Shuffle `GatedRMSNormWarp128Kernel` (`linear_attention_gates_impl.cu`)**: - For `sizeof(T) == 2 && norm_size == 128`, replaces the 1-row/CTA scalar `cub::BlockReduce` kernel (which read `input` twice and used `__syncthreads()`) with a 4-warp/CTA `uint2`-vectorized warp-shuffle kernel (`REG:26 STACK:0 SHARED:0`, 0 barriers, 4x fewer CTAs). It is selected only when the output, input, scale and gate pointers are 8-byte aligned; otherwise the generic kernel runs. ## Testing - `GatedDeltaNetTest.ChunkedMisalignedFloat16IoBuffers` / `ChunkedMisalignedBFloat16IoBuffers`: bind query/key/value/output through IOBinding at 2-byte and 8-byte offsets into a larger allocation, with `ORT_GDN_PLAN=chunked` pinned, `chunk_size` 64 and 32, and the BF16 `fp32`/`fp16`/`bf16` precision modes. They fault if either alignment guard is removed. - `ContribOpGatedRMSNormTest.Float16_PartialFinalCta` / `BFloat16_PartialFinalCta`: 1, 3 and 5 rows (a partially filled last CTA) with SiLU and sigmoid gates. - Verified on an H200 (SM90): `GatedDeltaNet*:ContribOpGatedRMSNormTest.*:ContribOpLinearAttentionGateTest.*:*LinearAttention*` passes (142 passed, 19 WebGPU tests skipped). ## Performance (NVIDIA H200, Qwen3.8-Flash-Next 36 GDN layers) - **`GatedDeltaNetChunkedKernel` (=2048$, 36 layers)**: - Executed instructions cut by **`-45.5%`** (`264.53M -> 144.05M` per launch), stack frame reduced from `64 B -> 0 B`. - Latency reduced from **`34.74 ms` (`964.96 us/call`) -> `22.18 ms` (`616.11 us/call`) (`1.57x` faster, `-12.56 ms` saved per 2K chunk)**. - **`GatedRMSNormWarp128Kernel` (=2048$, 36 layers)**: - Latency reduced from **`3.08 ms` (`85.48 us/call`) -> `1.05 ms` (`29.08 us/call`) (`2.94x` faster, `-2.03 ms` saved per 2K chunk)**, achieving `85.36%` Compute SOL and `89.73%` occupancy. --------- Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

### Description  Batch compatible streamless GPU-to-GPU `OrtApi::CopyTensors` vectors into one call-local recording and queue submission by default. - Encode copies in array order, preserving dependencies between entries. - Retain the context through submission and isolate recordings for concurrent calls and error recovery. - Preserve existing paths for single copies, CPU transfers, explicit streams, and incompatible vectors. - Add provider tests for submission counts, stream routing, copy dependencies, error recovery, buffer lifetimes, and concurrent Session teardown. No public API/ABI changes or experimental flags are introduced. ### Motivation and Context  Fixed-state caching can issue many small GPU copies during each execution phase. Submitting every copy separately adds avoidable host and queue-submission overhead. This change allows callers to submit those copies together without adding a host-completion wait. The companion consumer is microsoft/onnxruntime-genai#2775. Together, the changes saved approximately 0.92 seconds for baseline generation and 1.52 seconds for DFlash2 width 3 in isolated A100 measurements with 512 generated tokens. These savings are configuration-dependent; generated tokens, speculative counters, and final fixed-state bytes were unchanged. Validation: Release runtime/plugin builds passed. The expanded WebGPU provider suites passed 66 tests, and serialized compatibility passed 24 tests, with no skips and two pre-existing disabled tests. Formatting/lint checks passed. Local execution used Linux/Vulkan on an A100; other backends remain for platform CI. --------- Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

…t#33551) ### Description  Automatically use `Queue.WriteBuffer` for eligible CPU-to-GPU uploads up to and including 1 MiB. - Flush deferred readers before queue writes to preserve read-before-overwrite ordering. - Preserve mapped-buffer handling and staging for larger, unaligned, and capture/replay uploads. - Retain existing completion, validation, device-loss, and failure handling. - Add focused upload tests and document ordering, source-lifetime guarantees, and the size threshold. ### Motivation and Context  Profiling Qwen/DFlash2 on A100 identified substantial host overhead from allocating, mapping, copying, and submitting a fresh staging buffer for each upload. Queue writes let Dawn reuse staging allocations and batch pending transfers with subsequent queue work, without adding CPU completion waits. The conservative 1 MiB cap is the largest consistently beneficial size measured on A100/Vulkan. Larger transfers can regress, so they retain the existing staging path rather than extrapolating the small-upload improvement. **Validation:** On Linux/A100/Vulkan, 80 native tests passed with one unrelated build-mode skip, and 100 GenAI tests passed. Both 512-token baseline and DFlash2 generation runs preserved exact tokens and speculative counters. Formatting and lint checks passed. ### Update: Rebased onto latest main, including merged PR microsoft#33491, and reran A100/Vulkan benchmarks to measure the incremental benefit of `Queue.WriteBuffer`. **Results**: - Baseline: **3.3% faster** (short), **6.8% faster** (long) - DFlash2 width 3: **2.5% faster** (short), **3.1% faster** (long) All generated tokens and speculative counters matched. Native and plugin tests passed. Existing batched uploads remain unaffected, and larger transfers retain the staging path. This confirms that the optimization provides additional performance benefits beyond microsoft#33491. --------- Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

…ll-N projections (microsoft#33600) ### Description - Extend the `M = 8` narrow-`CtaN` (`CtaNLargeM / 2 = 2`) specialization in `onnxruntime/contrib_ops/cuda/llm/fpA_intB_gemv/dispatcher.h` beyond FP16 INT4. It now also covers `BF16DetailsA` and `Int8DetailsW` (`Details::kStepK == 16 && Details::kInterleave == 2`), so the scale-only SM80-interleaved INT4 and INT8 kernels use it for both FP16 and BF16. - The tile choice is a pure selector, `PickM8GemvCtaN` in `fpA_intB_gemv.h`: - **A100 (SM 8.0, 108 SMs)** uses windows qualified with same-shape narrow-versus-default timings. INT4 uses the narrow tile for `3456 < N <= 4320`, or for `K <= 3072` and `1728 < N <= 2560`. INT8 uses it for `864 < N <= 1304`. The windows apply only to scale-only launches (no activation scale, bias, or alpha-in-advance) with group size 32, 64, or 128; other A100 launches keep the default tile. - **Other devices** keep the original cutoffs: `grid.y = N / (CtaNLargeM * kInterleave) <= 320` for INT4 (`N < 5136`) and `<= 160` for INT8 (`N < 1288`). Larger `N` keeps the default tile. - The caller reports the device: `MatMulNBits` and its tactic profiler set `Params::device_sm` (`major * 10 + minor`) and `Params::device_sm_count` from the execution provider's device properties, and `IsA100Device()` (SM 8.0 with 108 SMs, which excludes the 56-SM A30) selects the A100 rule. The dispatcher makes no CUDA runtime call, and a caller that leaves both fields at 0 gets the device-independent cutoffs. The opt-in wave-aware `M = 8` tile (`ep.cuda.fpa_intb_gemv_wave_aware`, FP16/INT4, sm_12x only) takes its SM count from the same fields through `GemvPickSmCount()`, which used to query the device on every launch, so a launch that reports no device gets no wave-aware tile. The profiler reports the device too, so it times the tile that production launches use. On A100 the dispatched launch's warm CPU enqueue time exceeds a forced launch of the same tile by 0.003-0.022 µs (0.1-0.8 %), with no change in device time. - The native Hopper layout (`kInterleave == 1`), the zero-point kernels and INT2 (`kStepK == 64`) are unchanged. The tile only decides which block covers which columns, so outputs are bit-identical between tiles. - The `[fpA_intB_debug]` launch log now prints `M` and `CtaN`, so tests can assert which tile was launched. - Add regression tests for the narrow and fallback paths and for the A100 windows (see Testing). Document the tile rule in `docs/contrib_ops/cuda/matmul_nbits.md` and the A100 measurements in `docs/contrib_ops/cuda/matmul_nbits_small_m_experiments.md`. ### Motivation and Context - In `M = 8` speculative decoding (`DFlash2`), `fpA_intB_gemv` already halved `CtaNLargeM` from `4` to `2` for `FP16DetailsA` + `Int4DetailsW` (`kStepK == 32, kInterleave == 4`) to avoid wave under-utilization on small-`N` projections. - For `BF16` models with `INT8` dense projections (`kStepK == 16, kInterleave == 2`), an `N = 1280` projection with `CtaNLargeM = 4` (`CtaN * kInterleave = 8`) launched only `160` CTAs across `132` SMs (`1.21` CTAs/SM). - Extending the `CtaNLargeM / 2` dispatch for `grid.y <= 160` doubles the grid from `160` to `320` CTAs (`2.42` CTAs/SM), reducing Drafter `o_proj` latency on H200 from `14.50 µs` to `13.83 µs/call` while remaining bit-identical. - The H200 measurement is for BF16/INT8. FP16/INT8 and BF16/INT4 reuse the tile and cutoff of the layout they share with the measured kernels and were not measured on H200. - The `grid.y` cutoffs do not depend on the SM count, and the narrow tile is not a uniform win on a 108-SM A100. To qualify the newly enabled combinations there, I timed the narrow tile (`CtaN = 2`) against the default tile (`CtaN = 4`) with the production kernels on an A100-SXM4-80GB and gated the heuristic on the result. **A100 measurements** (`M = 8`, `K = 4096`, group size 32; weights rotated over 256-276 MiB, more than 6.4x the 40 MiB L2, to approximate DRAM streaming; per-kernel latency in µs, narrow / default, median of five independent rounds): | Weights | N | FP16 | BF16 | A100 tile | |---|---:|---|---|---| | INT4 | 1024 | 12.91 / 11.75 | 13.07 / 11.83 | default | | INT4 | 2880 | 15.22 / 14.41 | 15.24 / 14.67 | default | | INT4 | 4096 | 17.17 / 18.59 | 17.17 / 18.36 | narrow | | INT4 | 4320 | 17.52 / 18.71 | 17.51 / 18.46 | narrow | | INT4 | 4336 | 23.53 / 18.79 | 23.68 / 18.50 | default | | INT4 | 5120 | 29.87 / 19.15 | 29.81 / 18.83 | default | | INT8 | 640 | 9.80 / 9.36 | 9.98 / 10.09 | default | | INT8 | 1024 | 10.34 / 10.80 | 10.72 / 11.28 | narrow | | INT8 | 1280 | 10.82 / 11.24 | 11.23 / 11.76 | narrow | | INT8 | 1536 | 11.56 / 11.53 | 12.15 / 12.10 | default | - **INT4:** the narrow tile wins for `3456 < N <= 4320` (up to 19 % faster) and loses from `N = 4336`. It fits five resident blocks per SM, so `N = 4320` fills all `108 * 5 = 540` slots and larger `N` starts a second wave. It is 4 % to 58 % slower than the default tile there, and the previous `N <= 5120` cutoff selected it for that whole range. BF16/INT4 behaves like FP16/INT4: the paired narrow/default ratio is `0.95` at `N = 4320` and `1.28` at `N = 4336`. - **INT4, lower window:** the narrow tile also wins for `1728 < N <= 2560` when `K <= 3072` (the default tile is up to 8.4 % slower at `K = 1024`), but not at `K = 4096`, where the result is mixed. - **INT8:** the narrow tile wins for `864 < N <= 1304` (4-5 % at `K = 4096`, up to 9 % at `N = 872`, `K = 1024`). The edges are soft; see Known limitation. FP16 and BF16 show the same windows. - **Audit:** over 1,216 unique dtype x weight bits x `N` x `K` x group size x residency conditions (`K = 1024` to `8192`, group size 32 plus 64 / 128 spots, L2-resident and rotating weights), the windows select the narrow tile in 428 conditions. The worst regret of a selected narrow tile is 1.95 % and none exceeds 3 %. - **Production check:** a standalone harness built from the unmodified production dispatcher headers ran 86 launches in three rounds (258 logged tile choices) over all four dtype/weight combinations, narrow and default boundaries, group size 32 / 64 / 128 and `K` up to 8192. Every choice matches the windows, `M = 7 / 9` and a bias-enabled BF16/INT4 launch keep `CtaN = 4`, and production, forced-`CtaN = 2` and forced-`CtaN = 4` outputs are bit-identical. The dispatcher made 0 `cudaGetDevice` / `cudaDeviceGetAttribute` calls over cold, 1,000 warm, `M = 7 / 9` and bias-enabled launches, and a launch that reported no device took the device-independent tiles. The same harness covered the opt-in wave-aware tile: on the A100 with a reported SM 12.0 / 170-SM device, 52 launches (13 `N` x 4 reported devices) took the tile the wave-aware policy specifies (`CtaN = 2` at `N = 5184 / 8192 / 10240 / 16384`, `4` at `N = 5760 / 6144 / 12288`) with bit-identical outputs. The dispatcher made 0 runtime calls over cold and 1,000 warm launches, where the previous header made one `cudaGetDevice` call per warm launch (warm CPU enqueue time +0.106-0.112 µs, 3.4-3.6 %, versus +0.016-0.021 µs, 0.5-0.7 %, now). No sm_12x GPU was available, so these runs check the dispatch decision and its host cost, not the wave-aware speed-up. - **Method:** standalone SM80 builds of the production `kernel` template (not an ORT session) with `M = 8` and 128 threads. Each case replays a CUDA graph of at least 128 kernels and reports the median of 11 batch-average per-kernel latencies. Rounds are separate processes with shuffled shape and tile order. L2-resident cases reuse one 0.13-34 MiB weight buffer; rotating cases cycle distinct weight buffers over 256-276 MiB (more than 6.4x the 40 MiB L2). These are kernel-level timings with synthetic data, and rotating weights only approximate DRAM pressure. ### Testing - New `{Fp16,Bf16}Int{4,8}SymmetricGroupwiseTest.*M8NarrowTile` tests in `onnxruntime/test/contrib_ops/cuda_kernels/fpA_intB_gemm_kernel_test.cc`. They use scale-only fixtures (`zeros = nullptr`) with the SM80-interleaved layout (also on Hopper), group size 32, and `K = 1024` and `4096`. Each case checks the launched `CtaN` from the debug log, the output against the CUTLASS reference, and for `M = 8` bit equality with two `M = 4` launches. The fixtures report the device through `Params`, as `MatMulNBits` does. The four tests also check, tile only, that a launch that reports no device uses the non-A100 cutoffs, also on an A100. The expected tile follows the device: - Non-A100: INT4 `M = 8` with `N = 2880 / 5120` (narrow) and `N = 5136 / 6144` (fallback); INT8 `M = 8` with `N = 640 / 1280` (narrow) and `N = 1288 / 1536 / 2560` (fallback); `M = 7 / 9` fall back. - A100: INT4 `N = 3472 / 4320` (narrow) and `N = 2880 / 3456 / 4336 / 5120 / 5136 / 6144` (default); INT4 `N = 1744 / 2560` narrow at `K = 1024 / 2048 / 3072` and default at `K = 4096 / 8192`; INT8 `N = 872 / 1280 / 1288 / 1304` (narrow) and `N = 640 / 864 / 1312 / 1536 / 2560` (default); `M = 7 / 9` keep the default tile. - `FpAIntBGemvTest.A100M8DispatchUsesMeasuredWindows` and `FpAIntBGemvTest.OtherDevicesKeepM8DispatchCutoffs` call `PickM8GemvCtaN` directly for both device classes, including the window edges, `M = 7 / 9` and ineligible layouts, so both rules are tested regardless of which GPU runs the suite. `FpAIntBGemvTest.A100DeviceRequiresMatchingSmAndSmCount` covers the predicate: SM 8.0 with 108 SMs only, not the 56-SM A30, other SM versions, or a caller that reports nothing. - `get_m_list()` now includes `M = 8`, so the existing GEMV/CUTLASS comparison sweeps also run it. The tile those shapes take depends on the device: on non-A100 devices the compact build's scale-only INT4 shapes (`N = 2880`, `5120`) take the narrow tile, while its INT8 shapes and the full build's zero-point fixtures keep the default tile. On A100 all of them keep the default tile. - Run with `GTEST_FILTER='*M8NarrowTile*:FpAIntBGemvTest.*' ./onnxruntime_provider_test --gtest_filter=CUDA_EP_Unittest.All`. - Validated on an A100 (SM80, CUDA 13.0): all 28 tests matching `*Groupwise*`, `FpAIntBGemvTest.*` and `Int4DecodeTileTest.*` pass. Not run on H200 or other devices. Only the compact fpA_intB build (`USE_COMPACT_FPA_INTB_GEMM=1`) was run; the non-compact configuration of the test file was only syntax-checked. ### Known limitation - Only an A100-SXM4-80GB was available for this qualification. A100 40 GB and PCIe parts share SM80, 108 SMs and a 40 MiB L2, so they use the same windows, but they were not measured separately. On every other device (H100, L4, A10, RTX, ...), FP16/INT8 and BF16/INT4 still inherit the H200-derived `grid.y` cutoffs without timing, and the SM-count-agnostic INT4 cutoff may spill into a second wave there as it did on A100. Timings from other devices are welcome; `PickM8GemvCtaN` is where a device-specific window would be added. - The A100 windows are conservative, not universally optimal. At very small `N` the default tile gives up measured narrow wins (up to 6.5 % at INT4 `N = 256`, 4.0 % at `N = 512`, 9.1 % at INT8 `N = 128`). The INT8 window ends at `N = 1304`. Narrow wins continue to about `N = 1360` at `K = 1024 / 4096` (excluding `N = 1312` costs up to 5.8 %), but the benefit is marginal or reverses at `K = 8192`, the `N = 1328`-`1360` samples are group size 32 only, and `N = 1376` loses 3.8-5.2 % with L2-resident weights at `K = 8192`, so I did not extend the window. The lower INT4 window was validated at `K = 1024`, `2048` and `3072` only. - A wider `CtaN = 8` tile beats the better of `CtaN = 2` and `4` at some large `N` (for example BF16/INT4 `N = 6144`, `K = 8192`: 15.8 % lower latency). It is outside this PR and a possible follow-up. - The timings are kernel-level with synthetic weights and activations; end-to-end model latency was not measured on A100. --------- Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

…e dispatch (microsoft#33585) ## Description This PR optimizes the CUDA execution paths for 3-bit (`INT3`) and mixed-width integer `com.microsoft::QMoE` layers across both multi-token speculative decoding ($S=8$) and long-context prefill, and refines the multi-row decode dispatch threshold for `is_moe_gemv_fused_finalize_supported`. Together, these optimizations deliver a **2.05x–2.25x** speedup on cold-L2 INT3 decode GEMV kernels, a **16.5x** per-layer speedup on 2048-token INT3 prefill (**3.05x** faster end-to-end 2048-token TTFT on NVIDIA H200), and an **18%** end-to-end $S=8$ decode throughput improvement (`265 -> 312.76 tok/s` median, `322.00 tok/s` peak at context 2048). ## Summary of Changes ### File Changes | File | Change | |------|--------| | `onnxruntime/contrib_ops/cuda/moe/qmoe_kernels.cu` | Add branchless `__shfl_down_sync` + `__funnelshift_r` 3-bit word alignment (`AlignPackedCodeWord<3>`), `lop3.b32` + `FADD` centered dequantization (`Unpack8CenteredCodesFromWord<Bits>`), vectorized `FusedFc1FastVecKernel<T, Bits, BlockSize>`, `Fc2RouteFastVecKernel<T, Bits, BlockSize>`, `FinalizeRoutesVec8Kernel<T>`, and `BuildInt3RouteGroupsPairedFirstKernel` (`LaunchBuildInt3RouteGroups` throws on `num_pairs` outside `(0, 256]`), and dispatch the decode kernel through `SelectInt3DecodeKernel` in `LaunchInt3Typed`. | | `onnxruntime/contrib_ops/cuda/moe/qmoe_fused_route.cuh` | Add `FusedInt3SoftmaxTopKGroupsKernel<T, K>` for single-CTA warp-per-row `SoftmaxTopK` + paired-first `int2 route_groups` construction (`1 <= rows <= 8, num_experts == 512, k <= 16`). | | `onnxruntime/contrib_ops/cuda/moe/qmoe_kernels.h` | Declare `LaunchFusedInt3RouteGroups` and `LaunchBuildInt3RouteGroups`, add `const int2* route_groups` to `Int3MoeParams`, and add `Int3DecodeKernel` / `SelectInt3DecodeKernel`, the single place that chooses the scalar, fast, or vectorized decode kernel. | | `onnxruntime/contrib_ops/cuda/llm/moe_gemm/moe_int2_prefill.cu` | Add `RepackRawIntToSm80Int4ActiveExpertsKernel<SrcBits>` (`SrcBits` in `{2, 3, 4}`) and `TransposeScalesActiveExpertsKernel<T>` to repack only active experts on the fly into a shared scratch buffer before CUTLASS SM80 `MoeGemmRunner<T, cutlass::uint4b_t, T>`. For raw (INT3) weights the FC1/FC2 expert biases are passed to the grouped GEMM epilogue, so they are added to the FP32 accumulators and rounded once, matching the SIMT kernels. The raw-weight repack launcher enforces the word alignment its loads need (`GetInt2MoePrefillRawWeightAlignment`: 8, 4, and 16 bytes for INT2, INT3, and INT4). | | `onnxruntime/contrib_ops/cuda/llm/moe_gemm/moe_int2_prefill.h` | Add `bool weights_are_raw = false` to `Int2MoePrefillParams` and declare `GetInt2MoePrefillRawWeightAlignment`. | | `onnxruntime/contrib_ops/cuda/moe/moe_quantization.h` | Declare the default-off session option key `ep.cuda.qmoe_disable_int3_tensorcore_prefill` and its `disable_int3_tc_prefill_` member. | | `onnxruntime/contrib_ops/cuda/moe/moe_quantization.cc` | Route `QMoE::ComputeInt3` to `RunInt2MoePrefill` when `rows >= 16` (configurable via `ORT_INT3_PREFILL_MIN_ROWS`, opt-out via the session option `ep.cuda.qmoe_disable_int3_tensorcore_prefill=1` or `ORT_DISABLE_INT3_TENSORCORE_PREFILL=1`; the Tensor Core path also requires aligned packed weights for each projection (8/4/16 bytes for INT2/INT3/INT4) and 16-byte-aligned `input`, `output`, and biases, otherwise the SIMT path is used), and wire `LaunchFusedInt3RouteGroups` / `LaunchBuildInt3RouteGroups` into the decode path, building route groups only when `SelectInt3DecodeKernel` picks the vectorized kernels that read them. | | `onnxruntime/contrib_ops/cuda/llm/moe_gemm/moe_gemv.cu` | Gate `is_moe_gemv_fused_finalize_supported` at `num_rows <= 4` when `expanded_num_rows > 32` so $S=8$ speculative decoding (`num_rows = 8, expanded_num_rows = 80`) uses the parallel-route `moe_gemv_kernel` + `finalizeMoeRoutingKernel`. | | `onnxruntime/test/contrib_ops/moe_test.cc` | Add CUDA QMoE INT3 / mixed-width tests: SIMT vs Tensor Core prefill parity, vectorized GEMV per-weight rounding, 512-expert fused routing execution parity, and `QMoETest_Int3CudaPrefillBiasPrecisionParity` (bias added in FP32 before rounding, including FP16 overflow-then-cancel), and `QMoETest_Int3CudaPrefillMisalignedExternalBuffers_*` (IOBinding operands at byte offsets across the prefill threshold), `QMoETest_Int3CudaPrefillPreActivationOverflowExactWhenTensorCoreDisabled` (finite-input FP16 pre-activation overflow, exact with the Tensor Core prefill disabled), and `QMoETest_Int3CudaPrefillDisableTensorCoreOptionRejectsInvalidValue`. | | `onnxruntime/test/contrib_ops/cuda_kernels/softmax_topk_kernel_test.cc` | Add `FusedInt3RouteGroups_512ExpertsParityAndGrouping` for the fused softmax/top-k + paired-first route grouping kernel, and `FusedInt3RouteGroups_DecodeKernelSelection` for the `SelectInt3DecodeKernel` boundaries. | | `docs/contrib_ops/cuda/moe_qmoe.md` | Document the new INT3 decode/prefill kernel pipeline, paired-first route grouping, active-expert SM80 `fpA_intB` Tensor Core prefill repack, bias and precision semantics (including the FP16 pre-activation overflow limitation and the FP32 opt-in), the Tensor Core prefill alignment contract, route-group gating, and `is_moe_gemv_fused_finalize_supported` threshold in Section 6.7. | ### 1. INT3 Decode Kernel Optimizations (`rows < 16`) - **Branchless Non-Divergent 3-Bit Unpacking + IEEE-754 `lop3.b32` + `FADD` Dequantization (`qmoe_kernels.cu`)**: - Eliminates the warp-divergent conditional second global load `(sub == 1 || sub == 2) ? __ldg(&warp_u32[word_idx + 1]) : 0u` when 24-bit 8-weight packets straddle a 32-bit `uint32_t` boundary. Every lane now issues a single unconditional `__ldg` and obtains the next word from `lane + 1` via `__shfl_down_sync(0xffffffffu, w0, 1)` + `__funnelshift_r(w0, w1, bit_shift)` in `AlignPackedCodeWord<3>`. - Replaces 7 bit-shifts + 8 masks + 8 integer subtracts + 8 `I2F` conversions per 8 weights with IEEE-754 mantissa insertion via `lop3.b32` + `FADD` bias subtraction in `Unpack8CenteredCodesFromWord<Bits>` (`Bits` $\in \{2, 3, 4\}$), which is **100% bit-exact** (`0` diffs across all elements). - **Single-CTA Fused Routing & Paired-First Expert Route Grouping (`qmoe_fused_route.cuh`, `qmoe_kernels.cu`)**: - `FusedInt3SoftmaxTopKGroupsKernel<T, K>` (`1 <= rows <= 8, num_experts == 512, k <= 16`) computes warp-per-row `SoftmaxTopK` in registers (bit-identical to `SoftmaxTopKMergeKernel<T, 128, 4>`) and simultaneously packs `int2 route_groups` in a single CTA launch. - Both `FusedInt3SoftmaxTopKGroupsKernel` and `BuildInt3RouteGroupsPairedFirstKernel` (for general `routes <= 256`) place 2-route shared-expert pairs `(r0, r1)` ahead of singletons `(r0, -1)` using warp ballot/popc prefix sums so paired routes execute in the leading wave and load/unpack expert weights once for both tokens without creating tail-wave stragglers. - Route groups are built only when `SelectInt3DecodeKernel` selects the vectorized decode kernels, the only readers of `route_groups`; the scalar fallbacks (for example explicit zero points) and the non-vectorized fast path skip the build and run only `SoftmaxTopK`. - **Vectorized Decode GEMV & Route Reduction (`qmoe_kernels.cu`)**: - `FusedFc1FastVecKernel<T, Bits, BlockSize>`: `__launch_bounds__(128, 7)`, 4 output columns (8 gate/up rows) per warp, upfront 8-row weight & scale `__ldg` pipelining, `#pragma unroll 2`, compile-time `BlockSize` $\in \{32, 64, 128\}$, and compact `T` shared-memory staging (`10 KB` for 2 routes at `hidden = 2560` $\rightarrow$ 4–6 CTAs/SM). - `Fc2RouteFastVecKernel<T, Bits, BlockSize>`: `__launch_bounds__(128, 6)`, 4 output columns per 16-lane half-warp, upfront 4-row weight & scale `__ldg` pipelining, `#pragma unroll 5` (fully unrolling `640 / 128 = 5` iterations), and `2.5 KB` shared memory. - `FinalizeRoutesVec8Kernel<T>`: 128-bit `uint4` loads and stores reducing 8 columns per thread across `top_k` routes. ### 2. INT3 Prefill Active-Expert SM80 Tensor Core Grouped GEMM (`rows >= 16`) - **`RepackRawIntToSm80Int4ActiveExpertsKernel<SrcBits>` & `TransposeScalesActiveExpertsKernel<T>` (`moe_int2_prefill.cu`, `moe_quantization.cc`)**: - For prefill batches (`rows >= 16`, configurable via `ORT_INT3_PREFILL_MIN_ROWS`, opt-out via the session option `ep.cuda.qmoe_disable_int3_tensorcore_prefill=1` or `ORT_DISABLE_INT3_TENSORCORE_PREFILL=1`), inspects `expert_offsets` on device after `threeStepBuildExpertMapsSortFirstToken` and skips inactive experts (`expert_offsets[e + 1] == expert_offsets[e]`) in 0 cycles. - Repacks active experts' canonical row-local `[E, N, ceil(K * bits / 8)]` weights into SM80 `fpA_intB` `uint4b_t` layout (mapping centered 3-bit unsigned codes $u_3 \in [0, 7]$ to unsigned 4-bit nibbles $u_4 = u_3 + 4 \in [4, 11]$ via `0x44444444u` nibble bias and permuting 12 contiguous INT3 bytes into 16 contiguous `uint4` SM80 `fpA_intB` bytes) and transposes active experts' scales (`[E, N, K_blocks] -> [E, K_blocks, N]`) into a single shared scratch buffer reused between FC1 and FC2 before running CUTLASS `MoeGemmRunner<T, cutlass::uint4b_t, T>`. - **Bias and precision semantics (`moe_int2_prefill.cu`, `moe_quantization.cc`)**: - The FC1/FC2 expert biases are passed to the CUTLASS grouped GEMM epilogue, so `acc + bias` is computed in FP32 and rounded once, matching the SIMT decode/prefill kernels. The epilogue loads biases as 16-byte vectors, so a bias pointer that is not 16-byte aligned selects the SIMT path. - **Known limitation (FP16 only):** the grouped GEMM output is still `T`-typed, so an FP16 FC1 pre-activation with `|x| >= 65520` is stored as `+/-inf` before SwiGLU, as in the packed INT2/INT4 prefill and the generic CUTLASS QMoE path, whereas the SIMT kernels evaluate SwiGLU from FP32. This also happens for finite inputs: a gate accumulator of `-65536` (for example an input value of `16384` times a `-4` weight code) is finite in FP32, where `silu(-65536)` is `0`, but becomes `-inf` in FP16, where `silu(-inf)` is NaN and propagates through FC2. BF16 is unaffected. Real LLM FC1 pre-activations are not expected to approach this range, and keeping FP32 through the activation would need an FP32-output grouped GEMM or a fused SwiGLU epilogue (extra scratch and epilogue traffic on the default path), so the default stays on the Tensor Core path. - **FP32 prefill opt-in (default off):** set the session option `ep.cuda.qmoe_disable_int3_tensorcore_prefill` to `1` (it must be `0` or `1` and takes precedence over the environment variable) or `ORT_DISABLE_INT3_TENSORCORE_PREFILL=1` to run INT3 and mixed-width prefill on the FP32 SIMT kernels. SwiGLU is then always evaluated from the FP32 accumulators, whether or not the operands meet the alignment contract below, at the cost of prefill throughput. - **Alignment contract (`moe_quantization.cc`, `moe_int2_prefill.cu`)**: - The Tensor Core path reads raw packed weights as 8-, 4-, and 16-byte words (INT2, INT3, INT4) and the activations and biases as 16-byte vectors. Externally bound tensors, for example IOBinding buffers that start at an offset into a larger allocation, may not be aligned, so `QMoE::ComputeInt3` checks the FC1 and FC2 weights against their own bit width and `input`, `output`, and both biases against 16 bytes, and otherwise falls back to the byte-safe SIMT kernels. Scales are read element by element and need no check. The raw-weight repack launcher enforces the same weight alignment. ### 3. Multi-Row Decode Gate for `is_moe_gemv_fused_finalize_supported` (`moe_gemv.cu`) - Gates `is_moe_gemv_fused_finalize_supported` at `num_rows <= 4` when `expanded_num_rows > 32` so $S=8$ speculative decoding (`num_rows = 8, expanded_num_rows = 80`) dispatches to the parallel-route `moe_gemv_kernel` + `finalizeMoeRoutingKernel` (`46.03 us` vs `52.59 us` for `moe_gemv_fused_finalize_kernel`, which iterated serially over `experts_per_token = 10` inside each CTA). ## Motivation and Context ### Performance Benchmarks on NVIDIA H200 (SM90a, CUDA 13.0) Measured on Qwen3.8-Flash-Next (`hidden = 2560, intermediate = 640, num_experts = 512, top_k = 10, block_size = 64`, 40 INT3 QMoE layers + 8 INT4 QMoE layers, DFlash2 $S=8$ speculative decoding): #### Kernel-Level & Layer-Level Microbenchmarks (Nsight Compute / CUDA Graph Cold-L2) | Kernel / Stage | Baseline (Before) | Optimized (After) | Speedup / Delta | | :--- | :---: | :---: | :---: | | **Cold-L2 INT3 Decode FC1 (`rows=8, routes=80, b=64`)** | `125.41 us` (`FusedFc1FastKernel`) | **`55.80 us`** (`FusedFc1FastVecKernel`) | **2.25x faster** (`-69.61 us/layer`) | | **Cold-L2 INT3 Decode FC2 (`rows=8, routes=80, b=64`)** | `66.63 us` (`Fc2RouteFastKernel`) | **`32.58 us`** (`Fc2RouteFastVecKernel`) | **2.05x faster** (`-34.05 us/layer`) | | **INT3 Decode Routing (`rows=8, E=512, K=10`)** | `6.02 us` (`SoftmaxTopKMergeKernel`) | **`~4.0 us`** (`FusedInt3SoftmaxTopKGroupsKernel`) | **1.50x faster** (includes `route_groups`) | | **INT3 Decode Finalize (`rows=8, hidden=2560, K=10`)** | `3.97 us` (`FinalizeRoutesFastKernel`) | **`~3.0 us`** (`FinalizeRoutesVec8Kernel`) | **1.32x faster** (`uint4` vec8) | | **INT4 Decode FC2 (`rows=8, routes=80, b=64`)** | `52.59 us` (`moe_gemv_fused_finalize_kernel`) | **`46.03 us`** (`moe_gemv_kernel` + `finalize`) | **1.14x faster** (`-6.56 us/layer`) | | **INT3 Prefill per Layer (`rows=2048, routes=20,480`)** | `41.07 ms/layer` (`FusedFc1Fast` + `Fc2RouteFast`) | **`2.49 ms/layer`** (`0.81 ms` repack + `1.68 ms` TC GEMM) | **16.5x faster** (`-38.58 ms/layer`) | #### End-to-End Model Package Benchmarks on H200 (`qmoe_top08_44_33_b64`) | Workload / Metric | Baseline (Before) | Optimized (After) | Improvement | | :--- | :---: | :---: | :---: | | **2048-Token Prefill TTFT (`prompt_len=2048`)** | `2.310 s` (`1.572 s` with fast SIMT) | **`0.757 s` (`757 ms`)** | **3.05x faster** (`-1,553 ms` TTFT) | | **$S=8$ Decode Cycle Latency (`ctx=2048`)** | `18.89 ms/cycle` | **`14.56 ms/cycle`** | **`-4.33 ms/cycle` (`-22.9%`)** | | **$S=8$ Decode Throughput (`ctx=2048, gen=256`)** | `~265 tok/s` | **`312.76 tok/s` median (`322.00 tok/s` peak)** | **`+18.0%` tok/s** | ## Testing - **Unit & Parity Tests** (run from the build directory): ```bash ./onnxruntime_provider_test --gtest_filter='MoETest.*:CUDA_EP_Unittest.SoftmaxTopK*:CUDA_EP_Unittest.FusedInt3*' ``` - `MoETest.QMoETest_Int3AndMixedWidthCudaPrefillParity` and `MoETest.QMoETest_Int3CudaPrefillBiasPrecisionParity` compare the Tensor Core prefill path and the SIMT path with a reference; the latter uses biases that cancel exactly (and, for FP16, overflow) before the single rounding. - `MoETest.QMoETest_Int3Fused512ExpertRoutingExecutionParity` and `CUDA_EP_Unittest.FusedInt3RouteGroups_512ExpertsParityAndGrouping` cover the fused routing and paired-first route grouping. - `MoETest.QMoETest_Int3CudaPrefillMisalignedExternalBuffers_W{3W3,4W3,2W3,3W4,3W2}` bind packed weights (+1 and +4 bytes), scales, biases, and activations at byte offsets into larger device allocations through IOBinding, at 15 rows (SIMT) and 16 rows (Tensor Core threshold), and compare against aligned buffers. They fail with a misaligned-address fault if the alignment gate is removed. - `MoETest.QMoETest_Int3CudaPrefillPreActivationOverflowExactWhenTensorCoreDisabled` covers finite-input FP16 pre-activation overflow (a gate accumulator of `-65536` that is finite in FP32 but `-inf` in FP16) with the Tensor Core prefill disabled through the session option and through the environment variable, and BF16 on both paths. `MoETest.QMoETest_Int3CudaPrefillDisableTensorCoreOptionRejectsInvalidValue` checks that values other than `0` or `1` are rejected. The default FP16 Tensor Core behavior for such inputs is a documented limitation and is intentionally not pinned. - `CUDA_EP_Unittest.FusedInt3RouteGroups_DecodeKernelSelection` pins the `SelectInt3DecodeKernel` boundaries (alignment, zero points, block size, shape, route count, shared-memory limits). Route groups are built only for the vectorized kernels. - On H200 (SM90) this filter runs 114 tests: 97 passed, 0 failed, 17 skipped (CPU-only tests and tests that do not apply to the CUDA plugin build). - **Offline Kernel Checks**: - Nsight Systems on `MoETest.QMoETest_Int3AndMixedWidthCudaFastPathParity` (20 vectorized decode runs and 20 runs with zero points that use the scalar kernels): `BuildInt3RouteGroupsPairedFirstKernel` launches 20 times, once per vectorized run, and never on the scalar runs. - Verified bit-exactness of `Unpack8CenteredCodesFromWord<Bits>` (`lop3.b32` + `FADD`) against integer shift/mask/subtract `I2F` across all `2, 3, 4`-bit code values (`0` diffs across 51,200 FC1 and 204,800 FC2 elements). - Verified byte-for-byte identity (`0` mismatched bytes across `409,600` bytes per expert) between `RepackRawIntToSm80Int4ActiveExpertsKernel<3>` and offline `prepack_qmoe_sm80_torch` with shifted 4-bit codes (`u4 = u3 + 4`). - **End-to-End Model Behaviour Suite (`engine_behaviours.py --prefill fast`)**: - Passed all **27 / 27** deterministic generation and multi-turn behaviour checks (`BEHAVIOURS PASS`) on H200. - **Backward Compatibility**: - Set the session option `ep.cuda.qmoe_disable_int3_tensorcore_prefill=1` (or `ORT_DISABLE_INT3_TENSORCORE_PREFILL=1`) to force the exact FP32 SIMT path for prefill, for debugging, A/B comparison, or models that cannot rule out FP16 FC1 pre-activations beyond the FP16 range. The default behavior is unchanged. - Shapes or alignments outside the vectorized decode or Tensor Core prefill gates transparently fall back to the existing `FusedFc1FastKernel` / `Fc2RouteFastKernel` / scalar kernels. ## Checklist - [x] Tests verified (`MoETest.*` and `CUDA_EP_Unittest` gtests, plus 27/27 end-to-end behaviour suite) - [x] No breaking changes - [x] Documentation updated (`docs/contrib_ops/cuda/moe_qmoe.md`) --------- Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

### Description A legal rank-0 Gemm bias terminates the process when the packed WebGPU path reads C's last dimension. Guard this access in the packed and Intel implementations so scalar bias uses scalar storage and the existing broadcast handling. Add FP32/FP16 regressions for rank-0 and singleton biases, nonzero/zero beta, scalar/vec4 inputs, transposed inputs, and Intel-eligible shapes, with CPU fallback disabled. Validation on Windows x64 with bundled Dawn/D3D12: - Both baseline regression tests terminated with exit code 3; the fixed tests pass all 72 subcases with 72 logged packed Gemm dispatches. - `GemmOpTest.*`: 10 passed; `GemmOpTypedTests/*.*`: 46 passed; no GTest skips. Shared typed tests include CPU controls. - Native provider-test build, lintrunner, and diff checks passed. Intel-specific kernels were compiled but not executed locally; Metal was not tested. ### Motivation and Context Fixes microsoft#32839.

…33668) ### Description Return empty WebGPU pooling outputs after shape validation and allocation, before constructing a GPU program or dispatch. Add regression coverage for AveragePool, MaxPool and global pooling in NCHW/NHWC, including small/large windows and empty batches, with CPU fallback disabled. ### Motivation and Context An unpadded 65x65 AveragePool window over a nonempty 64x64 input produces an empty output. WebGPU currently attempts a zero-sized dispatch and fails instead of returning that output. Fixes microsoft#32841. ### Validation - Native Windows x64, bundled Dawn/D3D12: 85 pooling tests passed; 2 DirectML-only tests skipped. - All 3 new tests (16 scenarios) failed before the fix and passed afterward. - The original issue model passed three repeated runs per layout with full validation and CPU fallback disabled. - Lintrunner and `git diff --check` passed. Metal, browser/Wasm and plugin configurations were not tested.

### Description  Add a read-only GitHub Actions pinning job to the Lint workflow using SHA-pinned pinact-action v3.0.0, and document the equivalent local check. The job checks action and reusable-workflow commit SHAs and version comments without validation API calls or file modifications. Pin the remaining `actions/setup-python@v6` and `actions/setup-node@v6` references in `windows_webgpu.yml` to the verified v6.3.0 and v6.5.0 release commits, respectively. Validated in two separate commits, preserving the failing CI evidence: - `2172317414`: Added the linter without changing the existing unpinned actions. The [GitHub Actions pinning PR check failed](https://github.com/microsoft/onnxruntime/actions/runs/40593895342/job/122504780563), reporting exactly `actions/setup-python@v6` at line 453 and `actions/setup-node@v6` at line 469 of `windows_webgpu.yml`. - `b4a158719d`: Added the missing SHA pins. The [same PR check passed](https://github.com/microsoft/onnxruntime/actions/runs/40594080576/job/122506071131). Local validation: `pinact run --fix=false --no-api` (pinact v4.0.0) failed on the two unpinned references before the fix and passed afterward. `git diff --check` passed. ### Motivation and Context  Prevent mutable action tags from being introduced or overlooked. The new PR check was demonstrated to reject actual existing unpinned references before their SHA pins were restored. --------- Co-authored-by: Copilot App <223556219+Copilot@users.noreply.github.com>

…t#33612) ### Description  Add `tools/ci_build/github/azure-pipelines/templates/component-governance-variables.yml` as the shared definition of Component Governance directory exclusions. All 13 pipelines with explicit CG exclusion lists import the variable template and pass its value to `sdl.componentgovernance.ignoreDirectories` using a compile-time expression. Preserve all 17 existing exclusions and add only the CoreML Tools vendored pybind11 documentation directory: - `$(Build.SourcesDirectory)/build/Release/_deps/coremltools-src/deps/pybind11/docs` The vendored `deps/protobuf/java` directory remains in scan scope. Pipeline-specific alert policies, timeouts, and other settings remain unchanged. Validation: local PyYAML assertions passed for the shared template and all 13 consumers, the exact 18 exclusions, removal of only the proposed protobuf Java exclusion, and preservation of the documentation exclusion and all original exclusions. Earlier baseline semantic comparisons passed for unrelated pipeline settings. `git diff --check` passed. Azure-hosted 1ES template expansion and actual Component Governance scans were not run. ### Motivation and Context  Avoid maintaining identical exclusion lists independently across pipelines. CoreML Tools' vendored pybind11 documentation requirements register dependencies that ORT neither builds nor installs. Exclude only this unused documentation directory rather than all of `coremltools-src`; CoreML model formats, MILBlob, ModelPackage sources, and vendored protobuf Java projects remain in scan scope. --------- Co-authored-by: Copilot App <223556219+Copilot@users.noreply.github.com>

### Description Update four stale CUDA training call sites to match the current stream APIs: - pass the raw CUDA stream from `CudaKernel::Stream(ctx)` to the three `SoftMaxComputeHelper` calls; - pass both the ORT `Stream*` and raw CUDA stream to `GatherNDGrad::PrepareCompute`. The three loss changes are the same fix proposed in microsoft#29899, which was closed without being merged. I added the current `GatherNDGrad` caller as part of the same build repair and kept the original loss-fix author as a co-author. ### Validation Built a CUDA Training wheel from `b2097437af1184d2c0231fca2f2e6468158494b0` on WSL2 with CUDA 12.8.1, cuDNN 9.8, and `sm_120`. The four failures were reached during one full build; after these caller changes, the shared library and `onnxruntime_training-1.31.0+cu128` wheel completed successfully. The touched files are unchanged on current main `02f169292255f115c69dd15923b10d35767b024f`. I also loaded the resulting CUDAExecutionProvider on an RTX 5080 with CPU fallback disabled and ran a minimal CUDA graph successfully. Wheel SHA-256: `d19c746866854f23a7d63c541e36acee3d882b0174710a02b860ebba59441dd3` Co-authored-by: Lukáš Záborský <lukas.zaborsky@firma.seznam.cz>

…osoft#33674) - Hoist uniforms.K out of the K-loop condition so the loop test doesn't reload the uniform buffer each iteration. - Compute the subgroupMatrixStore stride as max(N, kSgMatDimN) once so Tint's range analysis can skip its own stride clamp on every store. - Compute the subgroupMatrixLoad offset from the loop index instead of accumulating it.

…t#33616) ### Description Fix QDQ MatMul fusion that can produce an unsupported `QLinearMatMul` or incorrect results when per-axis quantization parameters are used. The selector now validates that: - A and Y use scalar or single-element quantization parameters. - B uses per-tensor or supported 2D per-column quantization, with a supported axis and shapes. Unsupported cases fall back to the original `DQ → MatMul → Q` path by skipping fusion, preserving correctness without changing the kernels. fix: microsoft#33333

@timothyanderson096-ocdealcheck

…64 (microsoft#33598) ### Description `FusedConvFp16` steps through its packed weights with an `MLFloat16` pointer, but `MlasHalfGemmPackBSize` returns a size in bytes. From the second group on, `PrePack` writes each group's packed weights past the end of the buffer and `Compute` reads them back from there, so anything else allocated in that memory overwrites them. Grouped Conv also passed the first group's bias to every group's GEMM. This keeps the packed-weight offsets in bytes and gives each group its own slice of the bias. ### Motivation and Context Fixes microsoft#33156. The packing path only runs on arm64, where the NEON half GEMM has had a `CopyPackB` routine since microsoft#28786, so x64 builds never reach it. The bias problem predates that, as the issue notes. The fix follows the patch @timothyanderson096-ocdealcheck posted in the issue. ### Testing I cross-built `onnxruntime_provider_test` for aarch64 (GCC 13, Release) at bd6010c and ran it under `qemu-aarch64 -cpu max`. - On main, the new `ConvFp16Test.Conv2D_Group2_Bias_Issue33156` fails. It uses the issue's shapes (8x8 input, 7x7 kernel, stride 2, pad 3, two groups of three output channels). Group 1's channels should be 8, 10 and 12 and come out as 20, 22 and 24, as in the issue, and with a bias group 1 gets group 0's values (13.5 expected, 12.75 returned). - With the fix, all 36 `ConvFp16Test` tests pass. Of the 389 tests matching `*Conv*`, 365 pass and 23 are skipped (CUDA, DML or WebGPU only). `DeformConvTest.OnnxModelTest` fails in this setup because it loads `testdata/deform_conv_test.onnx` by a relative path; it doesn't use this kernel. I haven't run it on an arm64 device. --------- Co-authored-by: TIM ANDERSON <281533835+timothyanderson096-ocdealcheck@users.noreply.github.com>

hdharpure9922

LGTM

BBaditi and others added 25 commits October 27, 2026 16:26

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

d26a2e3

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

hdharpure9922 self-requested a review October 29, 2026 04:39

hdharpure9922 approved these changes Oct 29, 2026

View reviewed changes

hdharpure9922 merged commit 6d7eb62 into ovep-develop Oct 29, 2026
5 of 11 checks passed

Navigation Menu

Repository navigation

Sync with Microsoft ONNX Runtime - 29102026 - #1374

Sync with Microsoft ONNX Runtime - 29102026#1374
hdharpure9922 merged 25 commits into
ovep-developfrom
sync_msft_29102026

ai-fw-intg commented Oct 28, 2026

Uh oh!

hdharpure9922 left a comment

Uh oh!

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

18 participants

Conversation

ai-fw-intg commented Oct 28, 2026

Uh oh!

hdharpure9922 left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

18 participants