Sync with Microsoft ONNX Runtime - 02092026 - #1277
Merged
Merged
Conversation
## Description Refine fpA-intB GEMV support checks and test coverage for the compact kernel configuration introduced by PR microsoft#32324. The support query now evaluates the physical device architecture separately from the selected kernel/layout architecture, while compact mode remains limited to FP16 groupwise kernels using the non-SM90 layout. ## Summary of Changes ### GEMV support and dispatch | File | Change | |------|--------| | `onnxruntime/contrib_ops/cuda/llm/fpA_intB_gemv/fpA_intB_gemv.h` | Update the support-query declaration to accept both device and kernel architectures. | | `onnxruntime/contrib_ops/cuda/llm/fpA_intB_gemv/fpA_intB_gemv.cu` | Apply compact/full build checks using separate device and kernel architecture inputs, including SM90-layout and BF16 constraints. | | `onnxruntime/contrib_ops/cuda/quantization/matmul_nbits.h` | Pass the physical device architecture and selected packing architecture to the support query. | ### Tests | File | Change | |------|--------| | `onnxruntime/test/contrib_ops/cuda_kernels/fpA_intB_gemm_kernel_test.cc` | Add support-matrix coverage for combinations of device architecture, kernel/layout architecture, and kernel type. | | `onnxruntime/test/python/quantization/test_op_matmulnbits_prepacked_cuda.py` | Use compact-compatible `block_size=32` in the configuration-key fixture so the test exercises fpA-intB instead of falling back. | ## Testing - CUDA internal tests passed: 123 tests, including `FpAIntBGemvTest.SupportUsesDeviceAndKernelArchitectures`. - Four Python configuration-key tests passed against a fresh native build. - Both modified translation units compiled successfully with compact mode disabled/full-mode settings. - `git show --check` passed with no whitespace errors. - `lintrunner -a` was attempted; the Ruff, Ruff-format, and ClangFormat adapters failed without reporting violations in the changed files. ## Motivation and Context A Hopper device can use either the SM80-compatible packing/layout or the native SM90 layout. Treating the selected layout architecture as the device architecture incorrectly rejected valid compatibility-layout dispatch, while ignoring the layout allowed compact mode to select unsupported kernels. Keeping these inputs separate makes the support matrix reflect both hardware capability and the compiled kernel path. ## Checklist - [x] Tests added/updated - [x] No breaking changes - [x] Documentation updated (not applicable)
### Description Validates `blocksize` and uses overflow-safe shape arithmetic shared by `SpaceToDepth` and `DepthToSpace`. Removes the obsolete warning suppression and adds boundary tests. ### Motivation and Context Invalid or extreme block sizes could cause divide-by-zero or signed overflow during output-shape calculation. Co-authored-by: Daniel Song <danielsong@microsoft.com>
…icrosoft#32251) ### Description `SplitProgram` copies one element per invocation, and it works out where that element goes from scratch every time: flat offset to indices, read the split axis, loop over the cumulative split sizes to find the output, subtract the segment base, convert back to an offset. If the input is float32 and every output segment is a whole number of vec4s, none of that is necessary, because the split is just a contiguous partition of the buffer. This adds `SplitContiguousVec4Program` for that case, picked when the input is float32 and every `split_size * inner_size` divides by 4. It treats the input and every output as vec4, so each invocation moves 16 bytes instead of 4 and the dispatch is a quarter the size, and it finds the destination with an `if` chain over segment bounds baked into the shader instead of a runtime loop. Anything that does not qualify, meaning non-float32 inputs or unaligned segments, still goes down the existing path untouched. `segment_vector_sizes` is the only thing the new shader bakes in and it is also the cache hint, so the pipeline key covers exactly what the WGSL depends on. The test is in `js/web/test/data/ops/split.jsonc`: a `[2,4,3]` float32 split on axis 0 gives `inner_size` 12, so both segments are three whole vec4s and the case takes the new path. ### Motivation and Context The index arithmetic in the generic path costs more than the copy it is wrapped around, and it gets worse as tensor rank and output count go up. Detection heads are where you notice it. YOLO26n splits repeatedly in its head, and those splits are float32 with channel counts that are already multiples of 4, so they qualify. Classifiers barely split at all, so I would not expect this to do much for them. --------- Co-authored-by: Ananya Anand <t-anaanand@microsoft.com> Co-authored-by: Copilot App <223556219+Copilot@users.noreply.github.com> Co-authored-by: Ananya Anand <4n4ny4@users.noreply.github.com>
…crosoft#32317) ### Description `WebGpuDataTransferImpl::CanCopyImpl`, the plugin EP path, rejects GPU devices whose vendor id is not `VendorIds::NONE`, on the grounds that they belong to another EP. The built-in `DataTransfer::CanCopy` has no such check and returns true for any GPU device, so in a build with another GPU EP registered it can claim a handle it would then reinterpret as a WGPUBuffer. This applies the same rule in both places, and splits the device-compatibility half of the predicate into a static so it can be covered without a live device. ### Motivation and Context Found while testing the WebGPU data transfer directly. The two paths encode the same policy and had drifted; the plugin one already carries a comment explaining why the vendor check is needed. Co-authored-by: Ananya Anand <t-anaanand@microsoft.com>
### Description - Register native CPU FP16 `Gemm` and `MatMul` kernels only when an accelerated HalfGemm backend is available (ARM64 FP16 vector intrinsics, ARM64 KleidiAI HalfGemm override, or RISC-V configured with RVV/Zvfh). ARM64EC is excluded because MLAS does not compile the Neon HalfGemm dispatch there. - On other targets the FP16 nodes are promoted to FP32 by `InsertCastTransformer` as before. - Centralize the compile-time gate in `mlas.h` as `MLAS_HALF_GEMM_ACCELERATION_POSSIBLE` and the runtime gate in `MlasHalfGemmAccelerationSupported()`, used by registration and by both `Gemm` / `MatMul` `Compute()` paths. - Update kernel documentation and add registry and end-to-end coverage. ### Motivation and Context Unaccelerated x64 CPUs selected the portable FP16 kernels instead of the FP32 cast fallback, causing an approximately 3,500× Windows regression. Removing those registrations on unaccelerated targets restores `InsertCastTransformer` promotion to optimized FP32 execution. ### Notes - Kernel registration is process-wide with the default backend selector config. A session that sets `mlas.disable_kleidiai=1` on a KleidiAI-only build keeps the kernel registered but may fall back to the non-accelerated fp16 path in `Compute()`. Tracked as a follow-up. <!-- START COPILOT CODING AGENT SUFFIX --> - Fixes microsoft#32255 --------- 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: hariharans29 <9969784+hariharans29@users.noreply.github.com>
…ft#32330) first_batch and stride/blockIdx.x are int32; their product can exceed INT32_MAX for large batch counts, silently wrapping before being added to src/dst. This causes out-of-bounds reads/writes, observed as an illegal memory access, silently unwritten output, or occasionally a hang, depending on allocator layout. Fixes both softmax_warp_forward and softmax_warp_forward_resource_efficient, which have the identical pattern. Widened the offset computation to int64_t at the multiply site rather than changing first_batch/stride's types, since those are used elsewhere in signed-subtraction bounds checks that assume int. Verified with a standalone repro: an isolated arithmetic test showing the exact expression wraps to a negative offset, and a full GPU kernel test (~4GB fp16 tensor, ~2.1M rows) showing the unfixed kernel faults with an illegal memory access at the exact predicted boundary row, and the fixed kernel produces correct output across the boundary. Addresses microsoft#32299 Good — this is close, but it looks like the template's section headers (`### Description`, `### Motivation and Context`) got left as empty placeholders below your actual content instead of your content going *into* them. Let's restructure it properly and fold in the test output as evidence. Here's the full corrected PR body: --- ### Description Widens the offset computation in `softmax_warp_forward` and `softmax_warp_forward_resource_efficient` (`onnxruntime/core/providers/cuda/math/softmax_warpwise_impl.cuh`) from `int32` to `int64_t` at the multiply site. `first_batch` and `stride` (and `blockIdx.x`/`stride` in the resource-efficient variant) are both `int32`. Their product can exceed `INT32_MAX` for large batch counts, silently wrapping before being added to `src`/`dst`. This causes out-of-bounds reads/writes — observed as an illegal memory access, silently unwritten output, or occasionally a hang, depending on allocator layout. Kept `first_batch`/`stride`/`batch_size` as `int` rather than widening their declared types, since `local_batches = batch_size - first_batch` relies on signed subtraction elsewhere in the function and widening those types risked an unrelated regression for no benefit — only the multiplication result needed widening. ### Motivation and Context Addresses microsoft#32299. A tensor large enough to trigger this (`batch_size × stride > 2^31`) requires several GB of host memory to construct through the standard `OpTester` float-vector interface, which isn't practical to add as a CI unit test. Instead, I verified the fix two ways: **1. Isolated arithmetic test (no GPU):** confirms `first_batch * stride` computed in `int32` produces a wrapped/negative result for representative large inputs, and that the `int64_t`-cast version produces the correct value. **2. Standalone GPU repro:** a minimal extraction of the kernel's offset logic and warp-reduction structure, run against a ~4GB fp16 tensor (~2.1M rows × 1024 elements) straddling the exact overflow boundary (predicted boundary row: 2,097,152, where `first_batch × stride` first exceeds `INT32_MAX`). Before the fix, the kernel faults deterministically at the predicted boundary: ``` ❯ ./gpu_repro buggy Mode: BUGGY (int32 offset) batch_size=2101248 stride=1024 total_elements=2151677952 (4.01 GiB) overflow boundary row (int32 math): 2097152 Fill complete. *** Kernel execution FAILED: an illegal memory access was encountered *** ``` After the fix, the same tensor — including the rows immediately straddling the boundary — produces correct output: ``` ❯ ./gpu_repro fixed Mode: FIXED (int64_t offset) batch_size=2101248 stride=1024 total_elements=2151677952 (4.01 GiB) overflow boundary row (int32 math): 2097152 Fill complete. Kernel completed without a fault -- checking output correctness... -- Sanity check (well below overflow boundary) -- row 0 | expected peak col 1023 | actual peak col 1023 | PASS row 5 | expected peak col 0 | actual peak col 0 | PASS row 1000 | expected peak col 1023 | actual peak col 1023 | PASS row 500000 | expected peak col 1023 | actual peak col 1023 | PASS -- Boundary check (around row 2097152) -- row 2097150 | expected peak col 1023 | actual peak col 1023 | PASS row 2097151 | expected peak col 0 | actual peak col 0 | PASS row 2097152 | expected peak col 1023 | actual peak col 1023 | PASS row 2097153 | expected peak col 0 | actual peak col 0 | PASS row 2097154 | expected peak col 1023 | actual peak col 1023 | PASS Sanity rows: ALL PASS Boundary rows: ALL PASS ``` Existing `Softmax` op tests (`softmax_test.cc`) are unaffected by this change — it only touches the offset computation, not the reduction/arithmetic logic. **Note on Erf** The linked issue also reports Erf as silently producing wrong output above the same 2³¹-element boundary. I traced Erf's standalone CUDA path end-to-end, unary_elementwise_ops.cc::ComputeInternal (passes Tensor::Shape().Size(), which is int64_t) → Impl_Erf → UnaryElementWiseImpl → the templated _UnaryElementWise kernel in unary_elementwise_impl.cuh — and didn't find the same bug. The kernel's per-thread index is computed as static_cast<int64_t>(NumElementsPerThread) * NumThreadsPerBlock * blockIdx.x + threadIdx.x, which widens to int64_t before any multiplication happens, and the host-side grid-size calculation is already guarded with ORT_ENFORCE(blocksPerGridSize <= INT32_MAX, ...). This path appears correct as-is on main. This PR only fixes the Softmax case (softmax_warp_forward / softmax_warp_forward_resource_efficient). I don't have a repro for the Erf case described in the issue, it's possible it goes through a fused path (e.g. GELU) or is specific to the ROCm backend (the issue's hardware trace is from ROCm/MI300A) rather than the standalone CUDA unary-elementwise dispatch I checked. Flagging this rather than guessing at a fix for code I couldn't confirm is broken, happy to dig further if a maintainer can point me at the right path, or if the issue author can confirm which path they hit.
onnxruntime struggle to load qwen3.6-27B-int4 quantized with ModelBuilder. This PR improves the parallelization of the prepacking. It improves the creation of the session by 25%. Processor is Intel(R) Xeon(R) Platinum 8480C. ``onnxruntime.InferenceSession(<model>, provides=["CPUExecutionProvider"])`` Before: ``` -- model qwen36-27-cpu-int4/model.onnx -- start: 2026-08-06 11:16:48.388195 ---- end: 2026-08-06 11:17:32.072674 loading time: 43.68443649681285 -- start: 2026-08-06 11:17:32.072778 ---- end: 2026-08-06 11:18:16.123237 loading time: 44.0504179620184 -- start: 2026-08-06 11:18:16.123356 ---- end: 2026-08-06 11:19:00.205266 loading time: 44.08187505789101 ``` After: ``` -- model: qwen/qwen36-27-cpu-int4/model.onnx -- start: 2026-08-06 11:33:41.232455 ---- end: 2026-08-06 11:34:14.474138 loading time: 33.24157801596448 -- start: 2026-08-06 11:34:14.474373 ---- end: 2026-08-06 11:34:47.198557 loading time: 32.72414276842028 -- start: 2026-08-06 11:34:47.198684 ---- end: 2026-08-06 11:35:19.824560 loading time: 32.625834794249386 ``` --------- Co-authored-by: copilot-swe-agent[bot] <198982749+Copilot@users.noreply.github.com> Co-authored-by: xadupre <22452781+xadupre@users.noreply.github.com> Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
This pull request improves validation in the `ScatterND` operator to ensure that indices are empty when any indexed data dimension has size zero, and adds a corresponding unit test to verify this behavior. Validation logic improvements: * Updated `ValidateShapes` in `scatter_nd.h` to return an error if indices are non-empty and any indexed input data dimension is zero, preventing invalid scatter operations. Testing: * Added a new test (`ScatterND_nonempty_indices_reject_indexed_zero_dimension`) in `scatter_nd_op_test.cc` to assert that the validation correctly rejects non-empty indices when an indexed data dimension is zero. * Included the necessary header for `scatter_nd.h` in the test file to support the new test.
This pull request introduces a validation check in the `GatherND` operator to ensure that the `indices` tensor is not scalar (i.e., has rank greater than 0), and adds a corresponding unit test to verify the new error handling. This prevents invalid inputs from proceeding and improves the robustness of the operator. Input validation improvements: * Added a check in `gather_nd.cc` to return an error if the `indices` tensor has zero dimensions, with a clear error message. Testing enhancements: * Added a unit test `GatherND_runtime_scalar_indices_error` in `gather_nd_op_test.cc` to verify that passing a scalar `indices` tensor triggers the expected error.
…oft#29879) ### Description This PR fixes incorrect CUDA `ScatterElements` results for arithmetic reductions on element types that share the same byte width but have different numeric semantics. The current CUDA implementation selects representative template types by element size: * 1 byte → `int8_t` * 2 bytes → `half` * 4 bytes → `float` * 8 bytes → `double` This is valid for `reduction="none"`, where values are copied without numeric interpretation. It is incorrect for `add`, `mul`, `min`, and `max`, which must operate using the tensor's actual element type. This change preserves byte-width dispatch for `reduction="none"` and uses actual element-type dispatch for arithmetic reductions. ### Motivation and Context The existing dispatch can silently evaluate one numeric type using another type's semantics. For example: ```text data type: int32 data: [2] indices: [0] updates: [3] reduction: mul expected: [6] previous CUDA result: [0] ``` Because `int32` is four bytes wide, the previous dispatcher selected the `float` specialization. The input and update bit patterns were therefore interpreted and multiplied as floating-point values rather than integers. The byte-width dispatch was originally used to reduce template specializations for the assignment path. When CUDA arithmetic reductions were later added, the same representative-type dispatch remained in place. Byte width is sufficient when only the stored bits matter. It is not sufficient when the operation depends on the numeric meaning of those bits. The CUDA kernel is registered for all fixed-size tensor types and accepts these reductions, so returning incorrect values for the affected combinations is a correctness issue in the existing supported path. ### Implementation * Keep the existing representative-type dispatch for `reduction="none"`. * Dispatch `add`, `mul`, `min`, and `max` using the input tensor's actual element type. * Add the concrete `ScatterElements` template instantiations required by the corrected dispatch. * Extend the existing CAS-based atomic helper pattern with overloads for the additional registered types. * Add regression tests for the affected dtype and reduction combinations. `GatherElements` and `ScatterElements` previously shared one explicit-instantiation macro because both used the same representative types. Arithmetic `ScatterElements` now requires semantic element types, while `GatherElements` still only copies values and can retain the existing representative instantiations. The macros are therefore separated so that the additional instantiations are limited to `ScatterElements` rather than unnecessarily expanding the Gather path. Multiple updates may target the same output element, so arithmetic reductions require atomic updates. The new overloads reuse the existing generic CAS implementations instead of introducing separate atomic algorithms for each type. For `bool`, the corrected dispatch matches the ONNX reference semantics: * `add` and `max` use logical OR; * `mul` and `min` use logical AND. Boolean results are stored canonically as `0` or `1`, correcting the previous one-byte representative path's possible non-canonical result for `add(true, true)`. This change does not modify the ONNX schema, public API, registered element types, CPU provider, ROCm provider, or the existing `reduction="none"` behavior. ### Validation The regression coverage includes: * `add`, `mul`, `min`, and `max`; * signed and unsigned integer types; * floating-point types and `BFloat16`; * `bool` reductions and canonical `0`/`1` storage; * both `int32` and `int64` indices; * duplicate-index contention paths; * representative types that were already correct; * preservation of `reduction="none"` behavior. Local validation results: * Release CUDA `onnxruntime_provider_test` build: passed * Focused `ScatterElements` regression tests: 19/19 passed * Focused dtype-dispatch regression cases: 5/5 passed * `Scatter.*:ScatterElements.*`: 46/46 passed * Repeated contention stress cases: 600/600 passed * Lintrunner on the four modified files: passed * `git diff --check`: passed Validation environment: ```text WSL2 Ubuntu 24.04 NVIDIA RTX 4060 Laptop GPU (SM89) CUDA 12.8 cuDNN 9.7.1 Release configuration ```
Address two issues in Win Arm64 Cuda Plugin Packaging: (1) Authenticate Windows ARM64 CUDA downloads for new machine pool (2) Fix cudnn-frontend build error for a missing attribute in Windows ARM64 cuDNN 9.24.0.17 headers --------- Co-authored-by: Baiju Meswani <bmeswani@microsoft.com> Co-authored-by: copilot-swe-agent[bot] <198982749+Copilot@users.noreply.github.com> Copilot-Session: 28bad07c-098e-4917-a893-b7a772d5f036
ai-fw-intg
requested review from
Jaswanth51,
ankitm3k,
jatinwadhwa921 and
vthaniel
September 1, 2026 20:37
hdharpure9922
self-requested a review
September 2, 2026 04:58
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Automated daily backmerge from ORT main to ovep-develop. No conflicts detected. Do NOT squash or rebase - use merge commit only.