Sync with Microsoft ONNX Runtime - 29052026#1111
Merged
Merged
Conversation
### Description Swap the operand order of `vminps` in the FMA3 Erf kernel (both `x86_64/ErfKernelFma3.S` GAS and `amd64/ErfKernelFma3.asm` MASM, 5 instructions each) so the input sits in src2. `VMINPS` returns its src2 operand when either input is NaN, so this is what lets NaN propagate through the kernel. Adds `MathOpTest.ErfNaN` regression test (36 elements — covers both the 4×8 main loop and the 1×8 tail). ### Motivation and Context Fixes microsoft#28462. The FMA3 kernel had the input in src1, so a NaN input became the clamping constant (3.925f) and the polynomial then produced ~1.0. The C++ scalar / SSE2 / NEON / SVE paths already preserve NaN — only the FMA3 asm needed fixing.
… multi-threaded kernel (microsoft#28589) ### Description Replace the naive single-threaded scalar loop for 2-bit dequantization with float/MLFloat16 zero points with a multi-threaded kernel (`DequantizeBlockwise2Bits`) that: - **Parallelizes via `TrySimpleParallelFor`** — distributes work across all intra-op threads (previously single-threaded) - **Processes 16 elements per iteration** — one `uint32_t` = 16 packed 2-bit values, reducing per-element overhead - **Hoists scale/zp lookups** — all 16 elements share a block, so scale and zero_point are loaded once per batch Follows the same threading pattern as the existing 4-bit `DequantizeBlockwise` path for consistency. **Files changed:** - `matmul_nbits_impl.h` — declare `DequantizeBlockwise2Bits` - `matmul_nbits_impl.cc` — implement `Dequantize2BitsKernel` + `DequantizeBlockwise2Bits` with instantiations for `<float,float>` and `<float,MLFloat16>` - `matmul_nbits.cc` — replace naive loops in both `MatMulNBits<float>` and `MatMulNBits<MLFloat16>` `ComputeBUnpacked` ### Motivation and Context The `bits=2` + float zero_point path (added in microsoft#28354) was flagged with `// !!!!!!!!!!!!!! naive implementation, need to be optimized !!!!!!!!!!!!!!`. It ran ~20× slower than the `bits=4` MLAS path because it was a tight scalar `for n × for k` loop with no threading — the entire N×K dequantization ran on a single core before calling `MlasGemmBatch`. With 8 intra-op threads this should recover most of that gap. ### Benchmark Results Tested on a 96-core x86_64 Linux machine, ORT 1.27.0 CPU Release build, using typical LLM matrix shapes with `block_size=128` and float zero points. #### Multi-thread speedup (2-bit dequantization, 1 thread → 8 threads) | Shape (M×K×N) | 1-thread (ms) | 8-thread (ms) | Speedup | |---|---|---|---| | 1×4096×4096 | 41.0 | 8.5 | **4.84×** | | 32×4096×4096 | 47.9 | 8.8 | **5.46×** | | 1×4096×11008 | 120.7 | 24.2 | **4.99×** | | 32×4096×11008 | 146.8 | 28.2 | **5.21×** | | 1×11008×4096 | 119.2 | 24.5 | **4.87×** | | 32×11008×4096 | 154.4 | 28.2 | **5.47×** | | 1×1024×1024 | 1.18 | 0.16 | **7.61×** | #### 2-bit vs 4-bit comparison (ratio = 2-bit / 4-bit; <1.0 means 2-bit is faster) | Shape (M×K×N) | Threads | 4-bit (ms) | 2-bit (ms) | Ratio | |---|---|---|---|---| | 1×4096×4096 | 1 | 52.0 | 41.0 | **0.79×** | | 1×4096×4096 | 8 | 9.4 | 8.5 | **0.90×** | | 1×4096×11008 | 1 | 141.6 | 120.7 | **0.85×** | | 1×4096×11008 | 8 | 26.8 | 24.2 | **0.90×** | | 1×11008×4096 | 1 | 141.2 | 119.2 | **0.84×** | | 1×11008×4096 | 8 | 26.6 | 24.5 | **0.92×** | | 32×4096×4096 | 1 | 56.1 | 47.9 | **0.85×** | | 32×4096×4096 | 8 | 9.6 | 8.8 | **0.92×** | | 1×1024×1024 | 1 | 1.66 | 1.18 | **0.71×** | **Key findings:** - Multi-threading delivers **4.8–7.6× speedup** with 8 threads across all LLM shapes - 2-bit is now **10–30% faster** than 4-bit (ratio 0.71–0.93×), due to fewer bytes read from memory - The original ~20× regression (issue microsoft#28552) is fully resolved --------- Co-authored-by: copilot-swe-agent[bot] <198982749+Copilot@users.noreply.github.com> Co-authored-by: tianleiwu <30328909+tianleiwu@users.noreply.github.com> Co-authored-by: Tianlei Wu <tlwu@microsoft.com>
This pull request improves the accuracy and safety of device matching logic for GPU execution providers, especially when using numeric GPU indices (e.g., `gpu:1`) in layering rules. The changes ensure that matching by index only occurs when a runtime device ordinal is available (from `device_memory_info`), preventing accidental matches with hardware PCI device IDs. The update also enhances logging and expands test coverage for these scenarios. **Device matching logic improvements:** * Added a `has_device_ordinal` flag to `EpDeviceView` and updated matching logic so that index-based GPU matching only occurs when the device ordinal is known to be a runtime ordinal (from `device_memory_info`), not a hardware PCI ID. [[1]](diffhunk://#diff-a8f614056d63b5b3325eea1d855afc96550c977c16d8fdba641012a79194b7b5R169) [[2]](diffhunk://#diff-a8f614056d63b5b3325eea1d855afc96550c977c16d8fdba641012a79194b7b5L193-R198) [[3]](diffhunk://#diff-a8f614056d63b5b3325eea1d855afc96550c977c16d8fdba641012a79194b7b5L288-R299) [[4]](diffhunk://#diff-a8f614056d63b5b3325eea1d855afc96550c977c16d8fdba641012a79194b7b5R324) * Updated log messages to provide clearer error information when a layering rule with a numeric GPU index cannot be mapped, including guidance for troubleshooting. **Test improvements:** * Updated and expanded unit tests to cover correct and incorrect GPU index matching, including cases where only hardware IDs are present and should not match, and added a new test for execution providers with specific GPU ordinals. [[1]](diffhunk://#diff-37d64a2aa66018cc6a40ca2227432eae6c33dd6c1456d19ef539e869ee9d4f72L364-R366) [[2]](diffhunk://#diff-37d64a2aa66018cc6a40ca2227432eae6c33dd6c1456d19ef539e869ee9d4f72L375-R387) [[3]](diffhunk://#diff-37d64a2aa66018cc6a40ca2227432eae6c33dd6c1456d19ef539e869ee9d4f72R566-R584)
## Changes - Update `requirements.txt` to `protobuf>=4.25.8` - Update `requirements-training.txt` to `protobuf>=4.25.8` - Update `requirements-dev.txt` to `protobuf>=4.25.8` - Update `docs/python/requirements.txt` to `protobuf>=4.25.8` ## Notes This change addresses the direct Python manifest surface only. It does not claim to resolve every transitive Component Governance finding. Co-authored-by: arajendra <arajendra@users.noreply.github.com>
### Description Update GatherBlockQuantized to support 2-bits. Updated op schema, implemented the CPU and WebGPU EP. This helps to make the model smaller.
rev some npm packages --------- Co-authored-by: Copilot Autofix powered by AI <175728472+Copilot@users.noreply.github.com>
Co-authored-by: github-actions[bot] <41898282+github-actions[bot]@users.noreply.github.com> Co-authored-by: copilot-swe-agent[bot] <198982749+Copilot@users.noreply.github.com> Co-authored-by: guschmue <22941064+guschmue@users.noreply.github.com> Co-authored-by: Copilot <copilot@github.com>
…osoft#28659) ### Description Fix the CUDA plugin EP package test pipeline failure where the plugin is built with the latest code (which includes `float8e8m0` and other newer data types), but the host ORT 1.26 release does not support these types. When the plugin attempts to register kernel type constraints containing unsupported types, `GetTensorDataType` fails and the plugin load crashes. ### Motivation and Context The plugin EP architecture allows plugins to be built against a newer version of the ONNX Runtime headers while being loaded into an older host ORT. However, the existing `KernelDefBuilder::TypeConstraint` methods call `GetTensorType` (which throws on unsupported types), making it impossible for a forward-compatible plugin to register kernels that include newer data types in their type constraint lists. ### Changes - Add `TryGetTensorType()` — a non-throwing variant of `GetTensorType()` that returns `nullptr` when the host ORT does not recognize a tensor element type. - Add `TryMLDataTypeToOrtDataType()` — a non-throwing variant of `MLDataTypeToOrtDataType()` that returns `nullptr` instead of asserting/throwing. - Update `KernelDefBuilder::TypeConstraint` (both vector and single-type overloads) to use the `Try` variants and gracefully skip unsupported types rather than failing. ### Impact - Plugins built with newer headers can now load into older host ORT releases without crashing on unknown data types. - If all types in a constraint list are unsupported, the constraint is simply not registered (the kernel will not match, which is the correct fallback behavior). - No behavioral change when the host supports all types — the code path is identical to before.
…osoft#28675) ## Description The `windows_x64_asan / build_x64` CI pipeline has been failing with OOM (out-of-memory) because ASan-instrumented test binaries consume significantly more memory than normal builds, and CTest was running them at full CPU-count parallelism. This PR adds a `--test_parallel` argument to `build.py` that allows CTest concurrency to be configured independently from the build parallelism (`--parallel`). It then uses `--test_parallel 4` in the Windows x64 ASan workflow to cap test execution to 4 parallel jobs, preventing OOM while keeping build parallelism at full speed. ## Motivation and Context - ASan instrumentation inflates per-process memory usage by ~2-3x. - The existing `--parallel` flag controls both MSBuild and CTest concurrency together; there was no way to keep fast parallel builds while limiting test concurrency. - The CI runner has limited memory, and running all tests in parallel under ASan exceeded available RAM. ## Changes | File | Change | |------|--------| | `tools/ci_build/build_args.py` | Add `--test_parallel` argument (default: `None`, falls back to `--parallel`) | | `tools/ci_build/build.py` | Add `number_of_test_parallel_jobs()` helper; use it for CTest `--parallel`; validate negative values | | `.github/workflows/windows_build_x64_asan.yml` | Pass `--test_parallel 4` to cap ASan test concurrency | ## Testing - `python -m py_compile tools/ci_build/build.py tools/ci_build/build_args.py` — passes - `python tools/ci_build/build.py --help` shows the new `--test_parallel` option - `git diff --check` — no whitespace issues - When `--test_parallel` is omitted, behavior is unchanged (falls back to `--parallel` value)
…6345 (microsoft#28524) This pull request introduces comprehensive validation and error handling improvements for the ConvTranspose operator across CPU, CUDA, WebGPU, and XNNPACK backends, as well as in shape inference and unit tests. The main focus is to ensure that invalid input shapes (especially rank-0 or rank-1 tensors) are properly detected and reported, preventing undefined behavior and improving robustness. Additionally, error messages are clarified, and several helper functions now return `Status` for better error propagation. **Validation and Error Handling Improvements:** * All ConvTranspose implementations (CPU, CUDA, WebGPU) now explicitly check that input `X` and filter `W` tensors have at least 3 dimensions, returning clear error messages if not. (`[[1]](diffhunk://#diff-72fa27d94d5d92dd1e78ff510ef9a84d1ad74426c19af9722cf6511f8d38a5a8R65-R79)`, `[[2]](diffhunk://#diff-d1bbcb0542b5acea587ac929cd6362cfd11172c522505c6db8b457a9d470c63dR273-R289)`, `[[3]](diffhunk://#diff-b615243d0702e9613bd815173108306495b0f690294001e606823b77322f6fafR22-L28)`) * The shape inference function for `ConvTransposeWithDynamicPads` now fails gracefully with descriptive errors if input or weight tensors have fewer than 2 dimensions. (`[onnxruntime/core/graph/contrib_ops/contrib_defs.ccL62-R67](diffhunk://#diff-81f57d9adc2cce94f85a2949a895b7ff82efcc13d05e23ee6567661f0fecb7c0L62-R67)`) * Additional validation ensures that `output_padding` and dynamic pads have correct sizes, and that `output_padding` values are within ONNX-specified limits. (`[[1]](diffhunk://#diff-72fa27d94d5d92dd1e78ff510ef9a84d1ad74426c19af9722cf6511f8d38a5a8R138-R153)`, `[[2]](diffhunk://#diff-72fa27d94d5d92dd1e78ff510ef9a84d1ad74426c19af9722cf6511f8d38a5a8R171-R187)`) **Refactoring for Robustness:** * Helper functions such as `ComputePadsAndOutputShape` and `ComputeTransposePadAndOutputShape` now return `Status`, allowing errors to propagate and be handled appropriately rather than causing crashes or silent failures. (`[[1]](diffhunk://#diff-72fa27d94d5d92dd1e78ff510ef9a84d1ad74426c19af9722cf6511f8d38a5a8L165-R234)`, `[[2]](diffhunk://#diff-72fa27d94d5d92dd1e78ff510ef9a84d1ad74426c19af9722cf6511f8d38a5a8L194-R262)`, `[[3]](diffhunk://#diff-72fa27d94d5d92dd1e78ff510ef9a84d1ad74426c19af9722cf6511f8d38a5a8L220-R282)`, `[[4]](diffhunk://#diff-72fa27d94d5d92dd1e78ff510ef9a84d1ad74426c19af9722cf6511f8d38a5a8R291-R302)`) * All call sites (CPU, CUDA, WebGPU, XNNPACK) are updated to handle and propagate these errors using `ORT_RETURN_IF_ERROR` or `ORT_THROW_IF_ERROR`. (`[[1]](diffhunk://#diff-72fa27d94d5d92dd1e78ff510ef9a84d1ad74426c19af9722cf6511f8d38a5a8R171-R187)`, `[[2]](diffhunk://#diff-d1bbcb0542b5acea587ac929cd6362cfd11172c522505c6db8b457a9d470c63dL362-R379)`, `[[3]](diffhunk://#diff-b615243d0702e9613bd815173108306495b0f690294001e606823b77322f6fafL48-R60)`, `[[4]](diffhunk://#diff-6a2f8672090f25850b90b266aff3c7212552fc81b14bb7b539e9e5161c9fd526L494-R497)`) **Unit Test Enhancements:** * New negative tests are added to verify that rank-0 and rank-1 weight tensors are properly rejected and produce the expected error messages, increasing test coverage and reliability. (`[onnxruntime/test/contrib_ops/conv_transpose_with_dynamic_pads_test.ccR22-R56](diffhunk://#diff-cb5bfc51d0c8096922eb674d142f0e970d5becd140b47bdfd7729a06a818b598R22-R56)`) **Minor Code Quality Improvements:** * Improved memory management in the CPU implementation by wrapping the allocated buffer in `BufferUniquePtr` immediately to prevent leaks if exceptions are thrown. (`[onnxruntime/core/providers/cpu/nn/conv_transpose.ccR79-R89](diffhunk://#diff-0dcb5a9c8ba0c4e67940e9d77f77cb706bbf82d67bf6757967099b0a69c797b5R79-R89)`) * Minor includes and type safety improvements (e.g., use of `SafeInt` for overflow protection). (`[[1]](diffhunk://#diff-72fa27d94d5d92dd1e78ff510ef9a84d1ad74426c19af9722cf6511f8d38a5a8R22)`, `[[2]](diffhunk://#diff-72fa27d94d5d92dd1e78ff510ef9a84d1ad74426c19af9722cf6511f8d38a5a8R291-R302)`) **Summary of Most Important Changes:** **1. Validation and Error Handling** - All ConvTranspose implementations now check that input and filter tensors have at least 3 dimensions, returning clear errors if not. (`[[1]](diffhunk://#diff-72fa27d94d5d92dd1e78ff510ef9a84d1ad74426c19af9722cf6511f8d38a5a8R65-R79)`, `[[2]](diffhunk://#diff-d1bbcb0542b5acea587ac929cd6362cfd11172c522505c6db8b457a9d470c63dR273-R289)`, `[[3]](diffhunk://#diff-b615243d0702e9613bd815173108306495b0f690294001e606823b77322f6fafR22-L28)`) - Shape inference for `ConvTransposeWithDynamicPads` fails with descriptive errors for invalid input or weight tensor ranks. (`[onnxruntime/core/graph/contrib_ops/contrib_defs.ccL62-R67](diffhunk://#diff-81f57d9adc2cce94f85a2949a895b7ff82efcc13d05e23ee6567661f0fecb7c0L62-R67)`) - Additional checks for `output_padding` and dynamic pads sizes and values, with ONNX spec compliance. (`[[1]](diffhunk://#diff-72fa27d94d5d92dd1e78ff510ef9a84d1ad74426c19af9722cf6511f8d38a5a8R138-R153)`, `[[2]](diffhunk://#diff-72fa27d94d5d92dd1e78ff510ef9a84d1ad74426c19af9722cf6511f8d38a5a8R171-R187)`) **2. Error Propagation and Refactoring** - Helper functions now return `Status` and propagate errors; all call sites updated to handle these errors. (`[[1]](diffhunk://#diff-72fa27d94d5d92dd1e78ff510ef9a84d1ad74426c19af9722cf6511f8d38a5a8L165-R234)`, `[[2]](diffhunk://#diff-72fa27d94d5d92dd1e78ff510ef9a84d1ad74426c19af9722cf6511f8d38a5a8L194-R262)`, `[[3]](diffhunk://#diff-72fa27d94d5d92dd1e78ff510ef9a84d1ad74426c19af9722cf6511f8d38a5a8L220-R282)`, `[[4]](diffhunk://#diff-72fa27d94d5d92dd1e78ff510ef9a84d1ad74426c19af9722cf6511f8d38a5a8R291-R302)`, `[[5]](diffhunk://#diff-d1bbcb0542b5acea587ac929cd6362cfd11172c522505c6db8b457a9d470c63dL362-R379)`, `[[6]](diffhunk://#diff-b615243d0702e9613bd815173108306495b0f690294001e606823b77322f6fafL48-R60)`, `[[7]](diffhunk://#diff-6a2f8672090f25850b90b266aff3c7212552fc81b14bb7b539e9e5161c9fd526L494-R497)`) **3. Unit Testing** - Added tests to ensure invalid weight tensor ranks are rejected with proper error messages. (`[onnxruntime/test/contrib_ops/conv_transpose_with_dynamic_pads_test.ccR22-R56](diffhunk://#diff-cb5bfc51d0c8096922eb674d142f0e970d5becd140b47bdfd7729a06a818b598R22-R56)`) **4. Code Quality** - Improved buffer management and type safety in CPU backend. (`[[1]](diffhunk://#diff-0dcb5a9c8ba0c4e67940e9d77f77cb706bbf82d67bf6757967099b0a69c797b5R79-R89)`, `[[2]](diffhunk://#diff-72fa27d94d5d92dd1e78ff510ef9a84d1ad74426c19af9722cf6511f8d38a5a8R22)`, `[[3]](diffhunk://#diff-72fa27d94d5d92dd1e78ff510ef9a84d1ad74426c19af9722cf6511f8d38a5a8R291-R302)`)
- disable dynamic wgsl template tesst - disable shader cache key checks
…osoft#28369) BiasLoader hardcoded 128-bit vectorized loads (`ElementsPerAccess = 128/sizeof_bits = 8` for fp16) regardless of the `isAligned` template flag. When bias stride was not a multiple of 8, the unaligned kernel was selected but BiasLoader still used 128-bit loads → `cudaErrorMisalignedAddress`. **Fix**: Use `kAlignmentA` (4 for unaligned, 8 for aligned) instead of hardcoded 8. Tested with Gemma4 Attention + mask at all seq lengths 1–32. --------- Signed-off-by: Justin Chu <justinchu@microsoft.com> Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> Co-authored-by: Tianlei Wu <tlwu@microsoft.com> Co-authored-by: copilot-swe-agent[bot] <198982749+Copilot@users.noreply.github.com>
…rosoft#27707) ## Summary - Fix ORT raising "does not have type information set by parent node" when a subgraph references an initializer declared in the outer (parent) graph without explicit `value_info` in the subgraph - Propagate type info from implicit input defs to subgraph NodeArgs before subgraph verification in `VerifyNodeAndOpMatch` - Add regression test with an `If` node whose subgraph references an outer scope initializer without `value_info` ## Motivation Fixes microsoft#24880 When a node's op schema type inference function does not invoke subgraph inferencing (e.g., contrib ops like `BeamSearch`, `GreedySearch`, `WhisperBeamSearch`, `Sampling`), `InferAndVerifySubgraphTypes` is never called. This means type info from outer scope values — such as initializers declared in the parent graph — is never propagated to the subgraph's NodeArgs. When the subgraph is later verified in the second pass of `VerifyNodeAndOpMatch`, nodes referencing these outer scope values fail with a null type error. The existing workaround in `convert_generation.py` (manually adding `value_info` entries for moved initializers) confirms this gap in the type propagation path. ## Changes **`onnxruntime/core/graph/graph.cc`**: In `VerifyNodeAndOpMatch`'s subgraph verification loop, propagate type info from the containing node's `implicit_input_defs` to the subgraph's NodeArgs before calling `VerifyNodeAndOpMatch` on the subgraph. The propagation is guarded by `subgraph_nodearg->Type() == nullptr`, making it a safe no-op for standard ONNX ops (If/Loop/Scan) where `InferAndVerifySubgraphTypes` already set the types. For nested subgraphs, the recursive call to `VerifyNodeAndOpMatch` handles propagation at each level. **`onnxruntime/test/ir/graph_test.cc`**: Add `OuterScopeInitializerTypeInfoPropagatedToSubgraph` test that constructs a model proto with an `If` node whose subgraphs reference an outer graph initializer without `value_info`, and verifies `Model::Load` (which calls `Graph::Resolve`) succeeds. ## Test Plan - [ ] New C++ unit test `OuterScopeInitializerTypeInfoPropagatedToSubgraph` verifies model resolution succeeds - [ ] Existing `graph_test.cc` tests continue to pass (no regression in type inference for standard ONNX ops) - [ ] Existing control flow tests (If/Loop/Scan) continue to pass - [ ] CI lint checks pass (verified locally with `lintrunner`)
…lash_nvcc_threads, and enable quick build mode (microsoft#28645) ## Description Speed up CUDA CI build times by splitting the monolithic CUDA provider into architecture-specific OBJECT libraries with independent `nvcc --threads` control, and introducing a quick build mode (`onnxruntime_QUICK_BUILD`) that reduces kernel instantiations for CI validation. ## Motivation and Context CUDA builds were bottlenecked by `--nvcc_threads 1` across all targets because flash attention (48 .cu files, SM80+) requires ~4GB per nvcc thread and caused OOM when compiled with higher thread counts. The old heuristic in `build.py` used `psutil` to auto-detect memory but was unreliable and always conservative. By splitting flash attention into its own OBJECT library, the rest of the build can safely use `--threads 4` while flash attention stays at `--threads 2`. Combined with quick build mode (fewer kernel variants), this significantly reduces CI wall-clock time. ## CI Time Saving * N1F1: `--nvcc_threads 1`. CI time is from checks of [PR 28607](https://github.com/microsoft/onnxruntime/pull/28607/checks). * N4F2: `--nvcc_threads 4 --flash_nvcc_threads 2`: CI time is from this PR. * N8F4: `--nvcc_threads 8 --flash_nvcc_threads 4`: CI time is from this PR. * N4F4: `--nvcc_threads 4 --flash_nvcc_threads 4`: CI time is from this PR. This is the final candidate. * Saving = N1F1 - N4F4 * Saving Ratio = (N1F1 - N4F4) / N1F1 Here is CI time (Build + Test time in minutes) saving: CI | N1F1 | N4F2 | N8F4 | N4F4 | Saved Minutes | Saving Ratio -- | -- | -- | -- | -- | -- | -- Linux CI | 35 + 38 | 35 + 32 | 35 + 32 | 36 + 27 | 10 | 14% Windows CI | 58 + 36 | 53 + 38 | 54 + 38 | 48 + 36 | 10 | 11% Plugin Linux CI | 53 + 26 | 38 + 17 | 39 + 39 | 39 + 15 | 25 | 32% Plugin Windows CI | 77 + 16 | 57 + 14 | 54 + 14 | 53 + 12 | 28 | 30% Windows TRT CI | 54 + 43 | 38 + 38 | 42 + 43 | 41 + 37 | 19 | 20% Note that this is only one time comparison. Cache might take effect with more runs, and might change the statistics. The CI time is reduced in the range of 11% to 32%. Total CI time saving is more than 90 minutes. ## Key Changes ### 1. CMake: Architecture-specific OBJECT Libraries | File | Change | |------|--------| | `cmake/onnxruntime_cuda_source_filters.cmake` | New macros: `onnxruntime_extract_flash_attention_sources()`, `onnxruntime_extract_llm_sources()`, `onnxruntime_extract_sm_specific_cuda_sources()` to partition sources by SM arch | | `cmake/onnxruntime_providers_cuda.cmake` | Create `flash_attention` (SM80+), `llm` (SM75+), `sm90_tma`, and `sm120_tma` OBJECT libraries with per-target `--threads`; merge fpA_intB SM90 launchers into SM90 TMA lib | | `cmake/onnxruntime_providers_cuda_plugin.cmake` | Mirror OBJECT library pattern for plugin EP build; consolidate shared compile options into a variable; fix `-Xcudafe --diag_suppress=550,2810` and `--std c++20` for CUDA 12.8 compatibility | | `cmake/onnxruntime_unittests.cmake` | Link new OBJECT libraries into test target | ### 2. Build Script: `--flash_nvcc_threads` and Default 4 | File | Change | |------|--------| | `tools/ci_build/build.py` | Remove `psutil`-based memory heuristic; add `--flash_nvcc_threads` forwarding; default `nvcc_threads` to 4 | | `tools/ci_build/build_args.py` | Add `--flash_nvcc_threads` CLI argument (default: same as `--nvcc_threads`) | ### 3. Quick Build Mode (`onnxruntime_QUICK_BUILD`) - Reduces flash attention kernels to hdim128 fp16 only (skips hdim32/64/96/192/256) - Guards some MoE SM90 generated launchers with `#ifndef ORT_QUICK_BUILD` - Restricts CUTLASS SM80 tile configs to 3 instantiations - Skips test cases that depend on excluded kernel variants (e.g., `test_gqa_fp8_fallback_unsupported_head_size` needs hdim64) - Applied to all CI pipelines **except** Linux CUDA CI (full build) and packaging pipelines ### 4. CI and Packaging Pipeline Updates All CUDA CI pipelines updated from `--nvcc_threads 1` to `--nvcc_threads 4 --flash_nvcc_threads 4`: - `.github/workflows/linux_cuda_ci.yml` - `.github/workflows/linux_cuda_plugin_ci.yml` (+ `QUICK_BUILD=ON`) - `.github/workflows/linux_tensorrt_ci.yml` (+ `QUICK_BUILD=ON`) - `.github/workflows/windows_cuda.yml` (+ `QUICK_BUILD=ON`) - `.github/workflows/windows_cuda_plugin.yml` (+ `QUICK_BUILD=ON`) - `.github/workflows/windows_tensorrt.yml` (+ `QUICK_BUILD=ON`) Packaging pipeline updated to use `--nvcc_threads 4 --flash_nvcc_threads 2`, except `--nvcc_threads 2 --flash_nvcc_threads 1` for cuda plugin: - Azure Pipelines: `custom-nuget-packaging-pipeline.yml`, `nuget-win-cuda-packaging-stage.yml`, `plugin-win-cuda-stage.yml`, `py-win-gpu-stage.yml` - Linux scripts: `build_cuda_plugin_package.sh`, `build_linux_python_package.sh` ### 5. Bug Fix: CUTLASS Heuristic for SIMT Kernels - `onnxruntime/contrib_ops/cuda/llm/cutlass_heuristic.cc`: Fixed `ORT_QUICK_BUILD` path to return proper tile config for SIMT (float) gemm type instead of discarding the type info ## Architecture Mapping | OBJECT Library | Min SM | Sources | Threads | |---|---|---|---| | `*_flash_attention` | SM80+ | `bert/flash_attention/*.cu` (48 files) | `onnxruntime_FLASH_NVCC_THREADS` (default: same as nvcc_threads) | | `*_llm` | SM75+ | `contrib_ops/cuda/llm/*.cu` (excl. SM90/SM120 launchers) | `onnxruntime_NVCC_THREADS` (default 4) | | `*_sm90_tma` | 90a-real | MoE TMA + fpA_intB SM90 launchers | `onnxruntime_NVCC_THREADS` | | `*_sm120_tma` | SM120+ | MoE SM120 TMA generated files | `onnxruntime_NVCC_THREADS` | | Parent target | All archs | Everything else | `onnxruntime_NVCC_THREADS` | ## New Build Options - `--nvcc_threads N` (default 4) — threads for all CUDA targets except flash attention - `--flash_nvcc_threads N` (default: same as `--nvcc_threads`) — threads specifically for flash attention compilation CMake cache variables: `onnxruntime_NVCC_THREADS`, `onnxruntime_FLASH_NVCC_THREADS` ## Testing - Built locally with `CMAKE_CUDA_ARCHITECTURES="75;80;86;89;90;100;120"`, `--nvcc_threads 4 --flash_nvcc_threads 2` - Verified flash attention .cu files compile only for SM80+ (checked `build.ninja` / VS project) - Verified LLM .cu files compile for SM75+ - Ran `onnxruntime_provider_test` — all CUDA EP tests pass - Ran `python test_qmoe_cuda.py` (MoE kernels), flash attention / GQA tests - No link errors in both in-tree provider and plugin EP builds - No nvcc warnings about duplicate `--threads` flags - Plugin CI compile options verified: `--std c++20`, `-Xcudafe --diag_suppress=550,2810`, MSVC `/bigobj` all applied to OBJECT libraries
### Summary Lower ONNX `Sin` and `Cos` to the CoreML ML Program `sin` / `cos` elementwise ops via the existing `UnaryOpBuilder`, registered in the op builder factory. Like `Erf` / `Round` / `Exp`, these have no NeuralNetwork lowering (`UnaryFunctionLayerParams` has no sin/cos), so `IsOpSupportedImpl` rejects them on the NeuralNetwork format. ### Why `Sin` / `Cos` form the sinusoidal timestep embedding of diffusion UNets. Supporting them keeps that prologue on CoreML instead of splitting the graph — a tiny Stable-Diffusion UNet goes from **2 CoreML partitions → 1, zero graph breaks** with this change alone. This PR is **independent** of the rest of the series (it touches only the unary builder) and can be reviewed/merged in any order. ### Tests (`coreml_basic_test.cc`) - `SinCos_MLProgram` — a Sin + Cos graph runs fully on CoreML and matches the CPU reference. - `SinCosNeuralNetworkNotSupported` — the same graph falls back to CPU on the NeuralNetwork format. Doc: `coreml_supported_mlprogram_ops.md` lists `Sin` and `Cos`. ### Series — CoreML EP coverage for transformer / diffusion graphs - microsoft#28595 — Support bool Cast in ML Program *(prerequisite)* - **microsoft#28596 — Add Sin and Cos unary ops** *(this PR — independent)* - microsoft#28597 — Add Where and And builders *(depends on microsoft#28595)* - microsoft#28598 — Add GatherND builder *(depends on microsoft#28595)* Together with microsoft#28278 (scalar-`Gather`), the series takes BERT / GPT-2 / ViT / diffusion-UNet graphs — tiny and full-size — from 2 CoreML partitions to 1, with zero graph breaks. Co-authored-by: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
### Description TRT 11.0 removes many deprecated APIs, so guard TRT-EP code accordingly to support TRT 11 builds. ### Motivation and Context Fixes compilation with TRT 11.0 Signed-off-by: Kevin Chen <kevinch@nvidia.com>
…p-webgpu/MIN_ONNXRUNTIME_VERSION` (microsoft#28687) ### Description <!-- Describe your changes. --> Bake the contents of `plugin-ep-webgpu/MIN_ONNXRUNTIME_VERSION` into the plugin EP library as the `ORT_PLUGIN_EP_MIN_ORT_VERSION` preprocessor definition and pass it to `ApiInit()` so the EP refuses to load against an older ORT runtime. Rework `ApiInit()` to strictly parse the runtime version string as "MAJOR.MINOR.PATCH", optionally enforce a caller-supplied minimum, require MAJOR == 1, and use MINOR as the API version. All failure modes now throw `std::runtime_error` with a descriptive message. ### Motivation and Context <!-- - Why is this change required? What problem does it solve? - If it fixes an open issue, please link to the issue here. --> Enforce minimum ORT version for WebGPU plugin EP as specified in the minimum ORT version file. Previously, the version was hardcoded in `ApiInit()`. Co-authored-by: Copilot <175728472+Copilot@users.noreply.github.com>
…icrosoft#28682) ## Description Use fp32 accumulation in SkipLayerNormalization, SkipSimplifiedLayerNormalization, and EmbedLayerNormalization CUDA kernels to avoid overflow and improve numerical accuracy when processing fp16/bf16 data. The original implementation accumulated mean and variance statistics in the input data type (fp16/bf16), which can overflow for large hidden sizes or when input values have large magnitude. This change promotes all intermediate accumulation (mean, variance, normalization math) to fp32, matching the approach used by TensorRT-LLM's LayerNorm kernels. ## Motivation - fp16 has limited range (max ~65504) and precision (10-bit mantissa). Accumulating `x²/ld` across thousands of elements in fp16 easily overflows or loses precision. - bf16 has even less precision (7-bit mantissa), making accumulation errors more severe. - The fix is straightforward: cast to float before accumulating, compute normalization in float, cast back to the output type. ## Key Changes | File | Change | |------|--------| | `layer_norm.cuh` | Changed `LayerNorm`, `SimplifiedLayerNorm`, `LayerNormSmall`, `SimplifiedLayerNormSmall` to accept and operate on `float` for thread_data, epsilon, mu, rsigma. Removed unused `KeyValuePairSum` overloads for half/bfloat16. | | `skip_layer_norm_impl.cu` | Changed `SkipLayerNormKernel` and `SkipLayerNormKernelSmall` to accumulate in fp32 (`cub::KeyValuePair<float, float>`). Removed `maybe2half` helper (no longer needed). | | `embed_layer_norm_impl.cu` | Changed epsilon from `T` to `float`, accumulation to use `float` thread_data. | | `profile_skip_layer_norm.py` | New profiling script for nsys-based kernel timing analysis. | | `profile_skip_layer_norm.sh` | Shell wrapper for running nsys profiling. | | `parse_nsys.py` | Utility to parse nsys SQLite output and extract CUDA kernel timings. | ## Performance Results Profiled on NVIDIA GPU with nsys (B=1, seq_len=2048, fp16 data, 200 iterations, skip first 5 warmup): | Hidden Size | fp16 accum (μs) | fp32 accum (μs) | Regression | |---|---|---|---| | 768 | 3.81 | 3.81 | **0.0%** | | 1024 | 4.22 | 4.22 | **0.0%** | | 4096 | 13.01 | 13.03 | **+0.15%** (noise) | | 8192 | 28.94 | 28.94 | **0.0%** | **No measurable performance regression.** The kernel is memory-bandwidth-bound, so fp32 arithmetic is completely hidden behind memory latency. ## Testing - Existing unit tests pass (SkipLayerNorm, EmbedLayerNorm ops). - Profiling scripts added for reproducible performance measurement: ```bash cd onnxruntime/test/python/transformers nsys profile -o sln_fp16 --export=sqlite python profile_skip_layer_norm.py --mode fp16 --warmup 5 --repeat 100 python parse_nsys.py sln_fp16.sqlite --skip-first 5 ``` ## Related PRs microsoft#28442 microsoft#15660
This pull request makes a targeted update to the operator schema in the ONNX Runtime codebase, specifically clarifying the optional nature of certain outputs. Schema definition improvements: * Marked the `present_key` and `present_value` outputs as optional in the `ONNX_MS_OPERATOR_SET_SCHEMA` macro within `bert_defs.cc`, making the operator schema clearer and more flexible for consumers.
…rosoft#28260) ### Summary Add per-annotation-ID buffer managers and captured command storage so multiple generators can each capture and replay their own graph independently without cross-contamination Add ReleaseGraph API through the full ORT stack (EP base → C API → InferenceSession → plugin EP) to release captured commands and GPU buffers when a generator is destroyed Replace the single graph_buffer_mgr_ / is_graph_captured_ bool with per_graph_buffer_mgrs_ map and captured_graph_ids_ set keyed by annotation ID Use a std::function getter with cached pointer pattern in GpuBufferAllocator to dynamically route allocations to the active per-graph buffer manager during runs, while keeping Alloc/Free as simple pointer dereferences ### Motivation Edge's Prompt API speed benchmark creates multiple sessions/generators sequentially with graph capture enabled. With the existing single-graph design, the second generator replays the first generator's captured commands with wrong buffers, producing incorrect output and ultimately a QuotaExceededError in the browser. This PR isolates each generator's graph capture state so they don't interfere with each other. ### Related PR The GenAI side change is in microsoft/onnxruntime-genai#2106, which calls SessionReleaseGraph when a generator is destroyed to release the captured graph's GPU buffers. --------- Co-authored-by: copilot-swe-agent[bot] <198982749+Copilot@users.noreply.github.com> Co-authored-by: qjia7 <4221210+qjia7@users.noreply.github.com>
### Description We have internal alerts asking to update protobuf version in response to CVE-2026-0994. The alert asks tools/ci_build/github/linux/docker/inference/aarch64/python/cpu/scripts/requirements.txt to be updated to 5.29.6 while the rest are asked to be set to 6.33.5.
…28280) ### Description ## Summary Adds two WebGPU-only graph fusions and the contrib ops they target, plus a small refactor of the existing `MatMulNBits` dispatch logic so the new fused kernels can share its predicates. | Component | Files | Purpose | |---|---|---| | **`MatMulNBitsMlp` op + kernel** | `contrib_ops/webgpu/quantization/matmul_nbits_mlp.{cc,h}`, `*.wgsl.template` (3) | Fuses the SwiGLU MLP block: optional `(Skip)SimplifiedLayerNormalization` + two `MatMulNBits` projections (gate, up) + optional biases + `Sigmoid`/`Mul` (SiLU) + element-wise `Mul`. Single dispatch instead of 5–7. | | **`MatMulNBitsQkv` op + kernel** | `contrib_ops/webgpu/quantization/matmul_nbits_qkv.{cc,h}`, `*.wgsl.template` | Fuses `(Skip)SimplifiedLayerNormalization` + three `MatMulNBits` projections (Q, K, V) sharing the same input. Single dispatch instead of 4. | | **Op schemas** | `core/graph/contrib_ops/contrib_defs.cc` | `MatMulNBitsMlp` and `MatMulNBitsQkv` contrib op schemas (kMSDomain, opset 1). | | **Graph transformers** | `core/optimizer/matmul_nbits_{mlp,qkv}_fusion.{cc,h}` | Pattern-match the source subgraphs and emit the fused ops. EP-gated to WebGPU only — no impact on other EPs. Registered in `graph_transformer_utils.cc`. | | **Dispatch helpers** | `contrib_ops/webgpu/quantization/matmul_nbits_common.{cc,h}` + `matmul_nbits.cc` | Extracts the "would this dispatch use Subgroup-Matrix / DP4A / WideTile?" predicates into pure functions reusable by the fused kernels. No behavior change in the unfused `MatMulNBits` path. | | **Tests** | `test/optimizer/matmul_nbits_{mlp,qkv}_fusion_test.cc`, `graph_transform_utils_test.cc` | Unit tests for the new transformers (positive + negative cases). | ### Motivation and Context ~25-30% decode TPS throughput improvement on WebGPU + D3D backend on Windows. GPU used: RTX 5060Ti for Qwe3-1.7B. BEFORE (**95 decode TPS**): main branch <img width="344" height="140" alt="image" src="https://github.com/user-attachments/assets/0f5d7cfb-05f9-4f25-acb5-4becb8f5addd" /> AFTER (**120+ decode TPS**): PR branch <img width="359" height="134" alt="image" src="https://github.com/user-attachments/assets/f1254d8e-a400-4dbb-9d06-ab6116f929bb" /> --------- Co-authored-by: github-actions[bot] <41898282+github-actions[bot]@users.noreply.github.com>
…ft#28681) This pull request significantly improves the numerical correctness and robustness of L1 and L2 reduction operations (norms) for integer types on both CPU and CUDA backends. The main changes address integer overflow, undefined behavior, and precision loss in norm calculations, especially for edge cases like minimum representable integers and large accumulations. The changes also ensure consistency between CPU and CUDA implementations, and add detailed documentation for future maintainability. **Numerical correctness and overflow handling for integer norm reductions:** * On CPU, L1 and L2 reductions (`ReduceAggregatorL1` and `ReduceAggregatorL2` in `reduction_ops.h`) now accumulate in double precision to avoid integer overflow and undefined behavior, with Kahan summation for int64+ to minimize precision loss. Results are clamped to the maximum representable value to prevent overflow. [[1]](diffhunk://#diff-ca0c9224442a3c46251b0fb7326aacc1469bdee20ab409b930556f439d560015R722-R805) [[2]](diffhunk://#diff-ca0c9224442a3c46251b0fb7326aacc1469bdee20ab409b930556f439d560015R814-R893) * Introduced a `saturating_abs` function on CPU and a device-side `Impl_SaturatingAbs` kernel on CUDA to safely compute the absolute value for signed integer types, clamping to `max()` if `abs(min())` would overflow. [[1]](diffhunk://#diff-ca0c9224442a3c46251b0fb7326aacc1469bdee20ab409b930556f439d560015R722-R805) [[2]](diffhunk://#diff-f7138acd21464814d1793c9d334bee07d0cbe69719691e67efe3b2e23e4d06c7R516-R566) [[3]](diffhunk://#diff-945ab1deb57e1ff44b790cb2054537c252e23b7c7c374f44da66475361910abdR110-R119) **CUDA backend improvements and consistency:** * For integer reductions on CUDA, the input is cast to double before reduction, and the result is cast back to integer with saturating semantics (using PTX `cvt.sat`), matching the CPU's explicit clamping. This avoids precision loss and undefined behavior. * For no-op reductions (where input and output counts are equal), norm operations now use the saturating absolute value kernel to ensure non-negative results, even for edge-case values like `INT_MIN`. [[1]](diffhunk://#diff-ee5316fc3898058f70e942d9a84de36be4c7da09f144633a2504236430d5d033L209-R217) [[2]](diffhunk://#diff-ee5316fc3898058f70e942d9a84de36be4c7da09f144633a2504236430d5d033L592-R607) [[3]](diffhunk://#diff-ee5316fc3898058f70e942d9a84de36be4c7da09f144633a2504236430d5d033L778-R794) **Documentation and maintainability:** * Added detailed comments explaining the rationale and numerical properties of the new implementations, including why double precision is used, the limitations of float, and the behavior for large reductions and edge cases. [[1]](diffhunk://#diff-ca0c9224442a3c46251b0fb7326aacc1469bdee20ab409b930556f439d560015R722-R805) [[2]](diffhunk://#diff-ca0c9224442a3c46251b0fb7326aacc1469bdee20ab409b930556f439d560015R814-R893) [[3]](diffhunk://#diff-f7138acd21464814d1793c9d334bee07d0cbe69719691e67efe3b2e23e4d06c7R516-R566) [[4]](diffhunk://#diff-ee5316fc3898058f70e942d9a84de36be4c7da09f144633a2504236430d5d033L807-R861) These changes make norm reductions for integer types safe, mathematically correct, and consistent across CPU and CUDA, even for extreme or previously problematic inputs.
ai-fw-intg
requested review from
Jaswanth51,
ankitm3k,
jatinwadhwa921 and
vthaniel
ankitm3k
approved these changes
May 29, 2026
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.