ORT 1.28.0 release cherry-pick round 1#29771
Merged
Merged
Conversation
…on CPU-only Linux (#29590) ### Description `onnxruntime-gpu` 1.27 introduced a hard `NEEDED libcudart.so.13` entry in `onnxruntime_pybind11_state.so`, causing `ImportError` at `import onnxruntime` on CPU-only Linux machines — before any provider is selected. **Root cause:** `cmake/onnxruntime_python.cmake` was changed to compile `fpA_intB_gemm_adaptor.cu` and `fpA_intB_gemm_preprocessors_impl.cu` directly into `onnxruntime_pybind11_state.so` and link `CUDA::cudart` (dynamic). This embeds a load-time CUDA dependency in the Python module itself. **Fix:** Move the CUDA weight-preprocessing entry point (`pack_weights_for_cuda_mixed_gemm`) out of the main pybind module and into a **standalone extension module**, `onnxruntime_cuda_quant_preprocess`, that links `CUDA::cudart` on its own. The main `onnxruntime_pybind11_state.so` no longer compiles or links any CUDA code, so `import onnxruntime` has no `libcudart` dependency. The new module is imported **lazily** by `onnxruntime/python/tools/quantization/cuda_quantizer.py` only when weight prepacking is actually requested — never at `import onnxruntime` time. These preprocessing APIs are **offline-only** helpers: they are used by quantization tooling and model builders to produce prepacked weight initializers ahead of time, and are not part of the inference runtime hot path. Because nothing in the runtime imports them, isolating them into a separate, on-demand DLL has no runtime cost and cleanly keeps CUDA out of the base `import onnxruntime` path. **Why not the provider bridge:** An earlier iteration routed the call through the `ProviderInfo_CUDA` virtual interface (`TryGetProviderInfo_CUDA()`). That does not work for the CUDA-EP-as-plugin build (`onnxruntime_BUILD_CUDA_EP_AS_PLUGIN=ON`): `cuda_provider_factory.cc` is excluded from the plugin sources and there is no provider bridge, so `TryGetProviderInfo_CUDA()` returns `nullptr` and the call throws. The standalone module has no such dependency and works for **both** the legacy in-tree CUDA EP build and the plugin build. ### Key Changes | File | Change | |---|---| | `onnxruntime/python/onnxruntime_pybind_cuda_quant.cc` | **New.** Self-contained `pack_weights_for_cuda_mixed_gemm` (device malloc + transpose/convert + arch permutation) and a `PYBIND11_MODULE(onnxruntime_cuda_quant_preprocess, …)` entry point. | | `cmake/onnxruntime_python.cmake` | Add the `onnxruntime_cuda_quant_preprocess` module target (built when `onnxruntime_USE_CUDA AND NOT WIN32`, compiling the two `fpA_intB` `.cu` files + `CUDA::cudart` + cutlass, hidden visibility) and copy it into `onnxruntime/capi/`. Main pybind module keeps no CUDA sources/links. | | `onnxruntime/python/onnxruntime_pybind_quant.cc` | Remove the `USE_CUDA` `PackWeightsForMixedGemm` and its registration. The CPU-only `pack_fp4_weights_for_cuda_moe_gemm` stays in the main module. | | `onnxruntime/core/providers/cuda/cuda_provider_factory.{h,cc}` | Revert the `PackWeightsForMixedGemm` `ProviderInfo_CUDA` addition (no longer needed; absent in plugin builds). | | `onnxruntime/python/tools/quantization/cuda_quantizer.py` | `_get_pack_weights_for_cuda_mixed_gemm()` now imports `onnxruntime.capi.onnxruntime_cuda_quant_preprocess` lazily; add `has_cuda_weight_prepacking()` capability helper. | | `setup.py` | Package `onnxruntime_cuda_quant_preprocess.so` in the Linux/macOS wheels. | | `onnxruntime/test/python/quantization/test_op_matmulnbits_prepacked_cuda.py` | Point the prepacked-weight parity test and its skip guard at the new module. | | `docs/contrib_ops/cuda/matmul_nbits.md` | Update the offline-packer code snippets to import the new module. | ### Motivation and Context `import onnxruntime` must succeed on CPU-only machines even when the GPU wheel is installed. CUDA dependency errors should surface only when a CUDA provider is explicitly loaded/selected, or when offline CUDA weight prepacking is explicitly requested. This restores the 1.26 behavior where `onnxruntime_pybind11_state.so` had no `NEEDED libcudart.so.*` entry, and — unlike the provider-bridge approach — it also works in the CUDA-EP-as-plugin build. ### Testing Notes - Built both modules in the CUDA build; `readelf -d onnxruntime_pybind11_state.so` shows **no** `libcudart` `NEEDED` entry, while `onnxruntime_cuda_quant_preprocess.so` has `NEEDED libcudart.so.13`. - `import onnxruntime` and lazy loading of `onnxruntime.capi.onnxruntime_cuda_quant_preprocess` both succeed; `has_cuda_weight_prepacking()` returns `True` on a CUDA machine. - `test_op_matmulnbits_prepacked_cuda.py` passes (INT4/INT8 prepacked-vs-runtime parity), confirming the relocated packer produces byte-identical prepacked weights. --------- Co-authored-by: copilot-swe-agent[bot] <198982749+Copilot@users.noreply.github.com> Co-authored-by: Tianlei Wu <tlwu@microsoft.com>
… (sm_120) (#29706) ### Summary GroupQueryAttention's XQA decode kernel failed on consumer Blackwell GPUs (RTX 50-series, sm_120) with `cudaErrorInvalidValue`, while working fine on A100 (sm_80) and H200 (sm_90). This adds a runtime shared-memory capability check so XQA is only selected when the device can actually satisfy the kernel's dynamic shared-memory request, and otherwise falls back to cuDNN SDPA / Flash. A100/H200 behavior is unchanged. ### Root cause XQA bakes its shared-memory layout at compile time from `__CUDA_ARCH__`: - sm_80 / sm_87 / sm_90 use the large K/V-tile layout (`preferedKHeadPartBytes=128`, `cacheVTileSeqLen=64`) → up to ~140 KB of dynamic shared memory for head_size 128/256. - sm_86 / sm_89 / sm_120 use the small layout (64 / 32) → ~78–96 KB. Release/packaging binaries are built with a maximum arch of `90-virtual` (compute_90 PTX only, no native sm_120 SASS). On sm_120 the driver JIT-compiles that sm_90 PTX, so the kernel's `smemSize` carries the Hopper value (~140 KB). `launchMHA` then calls `cudaFuncSetAttribute(..., cudaFuncAttributeMaxDynamicSharedMemorySize, size)`, which exceeds sm_120's ~99 KB per-block opt-in limit (`sharedMemPerBlockOptin`) and returns `cudaErrorInvalidValue`. A100 (163 KB) and H200 (227 KB) have enough room, so they were unaffected. ### Key changes | File | Change | |---|---| | `xqa/xqa_impl_gen.cuh` | Add `GetSmemSize()` host helper that reads the per-kernel `smemSize` device symbol (accurate even for a PTX kernel JIT-compiled for the running SM). | | `xqa/xqa_loader_fp16_impl.cuh` | Add `GetXQAKernelSmemBytes(group_size)` head-dim dispatcher. The non-quantized fp16 footprint is an upper bound for the int8/fp8/bf16 variants (smaller cache element), so one query covers all XQA paths. | | `xqa/xqa_loader_fp16.cu`, `xqa/xqa_loader.h` | Expose `GetXQARequiredSharedMemoryBytes(device_prop, head_size, num_heads, kv_num_heads)`; a single non-templated entry point used by both the fp16 and bf16 GQA kernels. | | `group_query_attention.cc`, `group_query_attention.h` | Gate XQA selection on `required_smem <= device_prop.sharedMemPerBlockOptin`; fall back to cuDNN SDPA / Flash when it does not fit. Result is cached per node (`xqa_shared_memory_ok_`) since head_size/group are constant. | | `xqa/mha_impl.cuh` | Defensive backstop in `launchMHA`: if the requested shared memory still exceeds the device limit, throw an actionable message (which SM to build for / how to disable XQA) instead of the opaque `cudaErrorInvalidValue`. | ### CUDA graph safety `GetXQARequiredSharedMemoryBytes` uses `cudaMemcpyFromSymbol`, which synchronizes and is illegal during CUDA graph capture. The query is: - **cached** per node, so it runs at most once; - **guarded** with `onnxruntime::llm::common::isCapturing(Stream(context))` so the synchronizing call is only issued when the compute stream is not capturing; - resolved during ORT's non-captured warm-up run(s) before capture begins. If the value is somehow still unresolved while capturing, XQA is conservatively skipped for that run (safe fallback) without caching, so a later non-capturing run can resolve it. Warm-up and capture therefore make the same XQA/fallback decision, keeping the captured graph consistent with replay. ### Testing notes - Built the affected TUs (GQA dispatcher + fp16/bf16/int8/fp8 XQA loaders) with `CMAKE_CUDA_ARCHITECTURES="80;90"` (the configuration that reproduces the failure); all compile cleanly. - To validate the fix end-to-end, run a fp16/bf16 GQA decode workload on an sm_120 GPU (e.g. RTX 5090): it should now run (via fallback) instead of returning `cudaErrorInvalidValue`. Set `ORT_ENABLE_ATTENTION_KERNEL_DEBUG_INFO=1` to confirm the selected backend. - To actually run XQA (the fast path) on Blackwell, build with native arch `120` in `CMAKE_CUDA_ARCHITECTURES` (and `100` for datacenter Blackwell). With a native sm_120 cubin the layout is ~80 KB and fits, so XQA is selected.
To avoid hard dependency on nvrtc dll even when it is not used for some models.
Drop 52-real; 90-virtual Add 120-real; 120-virtual Ensure 86-real is included Q: Why not add 100-real to cuda 12.8 build? A: We assume that those machines will have cuda 13.x for best performance. Q: Why drops 52-real A: Many applications require float16 support, while 52-real cannot support it.
…lugin EP (#29620) ### Summary Phase 2 of the CUDA plugin execution provider "no-cuDNN" work. It lets single last-axis `ArgMax`/`ArgMin` run through a small custom CUDA kernel instead of cuDNN, fixes `LogSoftmax` classification in the plugin adapter, and adds a non-throwing cuDNN handle accessor so reduction kernels can fall back gracefully when cuDNN is disabled. ### Key Changes | Area | Change | |---|---| | `reduction_functions.cu` / `.h` | New `arg_min_max_last_axis<TIn, IsArgMax>` kernel (instantiated for `half`, `float`, `double`) that computes ArgMax/ArgMin indices over the last dimension of a row-major matrix without cuDNN. | | `reduction_ops.cc` | In `ReduceComputeCore`, route a single last-axis ArgMax/ArgMin (`CUDNN_REDUCE_TENSOR_FLATTENED_INDICES`) to the custom kernel when shapes fit `int`; otherwise fall through to the existing cuDNN path. `ReduceKernel::ComputeImpl` now uses `TryGetCudnnHandle`. | | `cuda_kernel.h` (native) / `cuda_kernel_adapter.h` (plugin) | Add `TryGetCudnnHandle`, which returns the cuDNN handle when available and `nullptr` otherwise (instead of throwing at handle acquisition). | | `softmax.h` | Detect `LogSoftmax` from `node.OpType()` instead of `info.GetKernelDef().OpName()`, so the plugin EP adapter classifies it correctly. | | `test_cuda_plugin_ep.py` | Add `LogSoftmax` and `ArgMin` tests; drop the `@requires_cudnn` gate from `ArgMax`, `ReduceMean`, `ReduceSum`; reduce over the last axis to exercise the cuDNN-free paths. | | `docs/cuda_plugin_ep/QUICK_START.md` | Drop `ArgMax` and reductions from the list of ops that still require cuDNN. | ### Correctness Notes - `select_last_index == 1` is already rejected on the CUDA EP, so the kernel keeping the first matching index (strict `>` / `<`) is spec-correct for the supported case. - The custom path guards `n > 0`, returns early for `m == 0`, computes the row offset in `int64_t`, and only engages when `m` and `n` fit in `int` (`gsl::narrow_cast`); larger tensors fall back to cuDNN. ### Testing - `python -m pytest onnxruntime/test/python/transformers/test_cuda_plugin_ep.py -k "log_softmax or argmax or argmin or reduce_mean or reduce_sum"` - Plugin no-cuDNN validation: `bash .env/cuda_130_plugin_no_cudnn.sh --build --test_plugin` - `onnxruntime_provider_test --gtest_filter='*Reduce*:*ArgM*'`
### Description This updates the Windows BinSkim-compliant build flags so `/Qspectre` builds also link against the MSVC Spectre-mitigated CRT/STL static libraries. `/Qspectre` only affects ONNX Runtime's own object files; BinSkim BA2024 can still report violations when the default non-Spectre `libcmt.lib`, `libcpmt.lib`, or `libvcruntime.lib` are linked into the final binary. ### Motivation and Context Release validation reported BinSkim BA2024 (`EnableSpectreMitigations`) warnings for `onnxruntime.dll` even when ORT was built with `--use_binskim_compliant_compile_flags`. The warning identified MSVC runtime and STL static libraries as the non-mitigated modules. This change makes the build option select the Spectre-mitigated MSVC library directory when it is available from the Visual Studio toolset. ### Key Changes - Adds `get_msvc_spectre_lib_dir()` to locate `%VCToolsInstallDir%\lib\spectre\<arch>` for the target Windows architecture. - Appends a quoted `/LIBPATH:<spectre-lib-dir>` linker flag whenever Windows BinSkim flags enable `/Qspectre` and AddressSanitizer is not enabled. - Emits a warning when the Spectre-mitigated MSVC libraries cannot be found, with guidance to install the Visual Studio "C++ Spectre-mitigated libs" component. - Preserves the existing ASAN behavior because ASAN libraries do not have Spectre-mitigated variants. ### Testing - `python3 -m ruff check tools/ci_build/build.py` - `python3 -m ruff format --check tools/ci_build/build.py` `lintrunner -a tools/ci_build/build.py` was also attempted. It found the repository config and applied no file changes, but the local environment could not execute the Ruff lintrunner adapters because `python` is not available on PATH; the direct `python3 -m ruff` checks above passed. --------- Co-authored-by: copilot-swe-agent[bot] <198982749+Copilot@users.noreply.github.com>
This pull request updates the build and CI configuration for CUDA-related workflows and the main CMake options. The main changes are the addition of new CMake build flags to enable CUDA quantization preprocessing, improved formatting for build flags, and a change in the default for the CUDA quant preprocess build option. These updates improve clarity, make it easier to customize builds, and ensure that the CUDA quant preprocess module is only built when explicitly requested. **Build configuration changes:** * Added the `--cmake_extra_defines onnxruntime_BUILD_CUDA_QUANT_PREPROCESS=ON` flag to the CUDA build jobs in `.github/workflows/linux_cuda_ci.yml` and `.github/workflows/linux_cuda_plugin_ci.yml`, enabling the CUDA quantization preprocessing module during CI builds. [[1]](diffhunk://#diff-04806013a5e7991ba5606145885d9b0fcd99a7df1f3bb96a2d38fc724ccd9b2aL32-R46) [[2]](diffhunk://#diff-04806013a5e7991ba5606145885d9b0fcd99a7df1f3bb96a2d38fc724ccd9b2aL114-R138) [[3]](diffhunk://#diff-64cd92765a9461e73a80c6b0401fe1960170834725a7e8a2ea153aff6d8f8388R45) * Added the `--cmake_extra_defines onnxruntime_QUICK_BUILD=ON` and `--cmake_extra_defines onnxruntime_USE_FPA_INTB_GEMM=OFF` flags to the CUDA no-cudnn build job for faster builds and to disable a specific GEMM implementation. [[1]](diffhunk://#diff-4e310144ab53bd9b6e48c7ceba29ad2c310724645278a28b85f7ba3a453c4980L35-R47) [[2]](diffhunk://#diff-04806013a5e7991ba5606145885d9b0fcd99a7df1f3bb96a2d38fc724ccd9b2aL114-R138) **Formatting and maintainability:** * Reformatted long `extra_build_flags` strings in workflow YAML files to use multi-line lists for improved readability and easier maintenance. [[1]](diffhunk://#diff-04806013a5e7991ba5606145885d9b0fcd99a7df1f3bb96a2d38fc724ccd9b2aL32-R46) [[2]](diffhunk://#diff-04806013a5e7991ba5606145885d9b0fcd99a7df1f3bb96a2d38fc724ccd9b2aL114-R138) [[3]](diffhunk://#diff-4e310144ab53bd9b6e48c7ceba29ad2c310724645278a28b85f7ba3a453c4980L35-R47) **CMake option default change:** * Changed the default value of the `onnxruntime_BUILD_CUDA_QUANT_PREPROCESS` CMake option from `ON` to `OFF` in `cmake/CMakeLists.txt`, so the CUDA quantization preprocessing module is only built when explicitly enabled.
### Description
When ONNX Runtime is built with the CUDA execution provider as a plugin
(`onnxruntime_BUILD_CUDA_EP_AS_PLUGIN=ON`), the EP-API op-kernel adapter
(`ep::adapter::KernelImpl::PrePackWeightImpl`) received a valid
`OrtAllocator*`
from the framework but discarded it and forwarded a **null**
`AllocatorPtr{}`
into the wrapped kernel's `PrePack()`. Any CUDA kernel that pre-packs a
constant
weight (e.g. `MatMulNBits`, `Conv`, `GroupQueryAttention`, quantized
MoE) then
allocated scratch through that null allocator and crashed during session
initialization:
```
IAllocator::ValidateAllocator(const T&) [with T = std::shared_ptr<onnxruntime::IAllocator>]
allocator != nullptr was false
```
This surfaced end to end as an ONNX Runtime GenAI `og.Model(...)`
failure on a
gpt-oss-20b (`MatMulNBits`) model when the CUDA EP was loaded as a
plugin: the
trivial init session succeeded, but the first real model session crashed
while
pre-packing quantized weights.
### Key Changes
| File | Change |
|---|---|
| `include/onnxruntime/ep/adapter/op_kernel.h` | `PrePackWeightImpl` now
wraps the incoming `OrtAllocator*` and forwards a valid `AllocatorPtr`
to `OpKernel::PrePack` instead of a null `AllocatorPtr{}`. |
| `include/onnxruntime/ep/adapter/allocator.h` | Add a **non-owning**
`IAllocatorWrappingOrtAllocator(OrtAllocator*)` constructor. The
framework owns the pre-pack allocator, so the wrapper must not take
ownership (an owning `Ort::Allocator` would release it on destruction).
Calls now go through the raw `OrtAllocator*` function pointers directly,
preserving the `Reserve`→`Alloc` (version ≥ 18) and
`GetStats`/`AllocOnStream` (version ≥ 23) fallbacks. |
| `onnxruntime/test/python/transformers/test_cuda_plugin_ep.py` | Add
`test_registration_matmul_nbits_prepack`: builds a fp16 `MatMulNBits`
model with a runtime-prepacked (`weight_prepacked=0`) quantized weight,
so weight pre-packing (`MatMulNBits::PrePack_B` →
`IAllocator::MakeUniquePtr(alloc, ...)`) runs during session creation.
This crashed before the fix and now passes. |
### Motivation and Context
The pre-pack allocator is provided and owned by the framework for the
duration
of the `PrePack` call. The legacy (in-tree) CUDA EP received it
correctly; only
the plugin op-kernel adapter dropped it. The non-owning wrapper matches
the
lifetime contract used elsewhere in the adapter (e.g.
`KernelInfoGetAllocator`)
and keeps the CUDA-EP-as-plugin build behaviorally identical to the
in-tree EP.
### Testing
- New `test_registration_matmul_nbits_prepack` in
`test_cuda_plugin_ep.py`
passes on a CUDA-EP-as-plugin build (`ORT_TEST_CUDA_PLUGIN_EP=1`) and
skips
gracefully when the device lacks fpA_intB GEMM support. It re-raises
(fails)
if the `allocator != nullptr` assertion recurs.
- Verified the model that originally reproduced the crash now creates
and runs
end to end through the CUDA plugin EP (gpt-oss-20b, `MatMulNBits`),
including a
100-sample MMLU sanity run.
- Existing `test_cuda_plugin_ep.py` registration tests continue to pass.
…t naming (#28896) ### Description This PR adds Windows ARM64 support to CUDA plugin packaging and updates related build/packaging logic for correctness and consistency across architectures. In response to review feedback, it also: - Corrects CMake comments to match actual Windows ARM64 CUDA toolkit search behavior. - Prioritizes architecture-specific cuDNN DLL search paths before generic fallback paths. - Aligns Windows ARM64 Python artifact naming with the ARM64 CUDA toolkit version. - Extends packaging metadata with per-platform CUDA version fields (including win-arm64) and updates metadata-reading template validation/exports accordingly. ### Motivation and Context Windows ARM64 CUDA plugin packaging requires architecture-aware handling for toolkit/cuDNN discovery and artifact metadata. These updates prevent cross-architecture ambiguity (especially for CUDA 13.x win-arm64 vs x64), improve downstream artifact selection reliability, and keep metadata semantics consistent with produced artifacts. --------- Co-authored-by: copilot-swe-agent[bot] <198982749+Copilot@users.noreply.github.com>
) ## Enable fpA_intB GEMM in CUDA builds and add configurable options ### Summary This PR turns the CUDA fpA_intB (weight-only, FP activation × int weight) MatMulNBits path on by default in CUDA builds, replaces the ambiguous `ORT_FPA_INTB_GEMM` bitmask with a simple on/off flag, and adds session-config keys so the path and its autotuning buckets can be controlled per session. It also makes the CUTLASS tactic profiler CUDA-graph safe and configurable. ### Motivation The fpA_intB kernels were previously gated behind `onnxruntime_USE_FPA_INTB_GEMM=OFF` and an `ORT_FPA_INTB_GEMM` integer bitmask (`0x01=all`, `0x02=GEMV`, `0x04=int4`, `0x08=int8`). The bitmask was ambiguous (e.g. `6` could read as "int4 + GEMV" or "GEMV for int4 and int8") and the GEMM/GEMV kernels actually share one weight layout, so splitting them was never valid. Shipping the kernels by default and exposing a plain enable flag plus per-session config makes the feature usable and tunable without rebuilding. ### Key Changes #### Build enablement | File | Change | |------|--------| | [cmake/CMakeLists.txt](cmake/CMakeLists.txt) | `onnxruntime_USE_FPA_INTB_GEMM` becomes a `cmake_dependent_option`, defaulting **ON** when `onnxruntime_USE_CUDA` is enabled (OFF otherwise). | | [tools/ci_build/github/linux/build_cuda_c_api_package.sh](tools/ci_build/github/linux/build_cuda_c_api_package.sh), [build_linux_python_package.sh](tools/ci_build/github/linux/build_linux_python_package.sh), [build_tensorrt_c_api_package.sh](tools/ci_build/github/linux/build_tensorrt_c_api_package.sh) | Flip packaging builds from `onnxruntime_USE_FPA_INTB_GEMM=OFF` to `ON`. | #### Option simplification (bitmask → boolean) - Removed `kFpAIntBGemmOption_All/Gemv/Int4/Int8` and the old bitmask parsing. - Added `ParseFpAIntBEnabled`: `""` / `"0"` / `"off"` → disabled; any other value → enabled (numeric non-zero still works for back-compat). - The enable flag now only governs nodes **without** prepacked weights. A prepacked weight is already stored in the fpA_intB layout, so the choice was fixed at export time; the constructor forces the path on for prepacked nodes and only `ORT_ENFORCE`s that the shape/hardware actually support it. - GEMV is no longer independently toggleable: it is enabled whenever supported, since GEMM and GEMV share the same weight layout. #### New session-config keys (EP-agnostic, config wins over env) | Config key | Env fallback | Meaning | |------------|--------------|---------| | `ep.cuda.fpa_intb_gemm` | `ORT_FPA_INTB_GEMM` | Enable/disable the fpA_intB path (`0`/`off` vs `1`/`on`). | | `ep.cuda.fpa_intb_profile_m` | `ORT_FPA_INTB_PROFILE_M` | Comma-separated initial profile-M buckets (e.g. `"1,8,64,512"`); empty uses the default bucket set. | These are read by both the built-in CUDA EP and the CUDA plugin EP via `OpKernelInfo::GetConfigOptions()`. #### Profiler: CUDA-graph-safe, in-memory autotuning - Added `getBestConfigOrProfile()` for lazy single-bucket profiling outside CUDA-graph capture; during capture the kernel falls back to a pure lookup (`getBestConfig`) because profiling launches kernels, records/synchronizes events, and allocates scratch — all illegal during capture. - Added configurable profile-M buckets: `ParseProfileMList`, `setProfileMOverride`, `getProfileMBuckets`, plus `kEnvProfileM` and `kDefaultProfileMaxM` (default max M lowered to `2048`). - Clearer error when an M bucket was not profiled before capture ("run a warmup inference outside capture first"). #### Docs - [docs/contrib_ops/cuda/matmul_nbits.md](docs/contrib_ops/cuda/matmul_nbits.md): `ORT_FPA_INTB_GEMM` documented as int/string on/off; clarified prepacked-weight strictness. ### Testing Notes - New: [onnxruntime/test/python/quantization/test_op_matmulnbits_prepacked_cuda.py](onnxruntime/test/python/quantization/test_op_matmulnbits_prepacked_cuda.py) — exercises the prepacked fpA_intB path and the boolean/numeric back-compat values of the enable flag. - To verify locally (CUDA, SM ≥ 75): - Build with CUDA (fpA_intB now defaults ON): `./build.sh --use_cuda ...` - Run: `python -m pytest onnxruntime/test/python/quantization/test_op_matmulnbits_prepacked_cuda.py` - Sanity-check config override: set `ep.cuda.fpa_intb_gemm=0` on a non-prepacked node and confirm the path is skipped; confirm a prepacked node still forces the path on.
Builds are blocked by onnxruntime-github-vs2022-latest. Try unblock our limited PRs for release. Co-authored-by: GitHub Copilot <copilot@example.com>
### Description <!-- Describe your changes. --> Add `OrtErrorCode::ORT_DEVICE_RESET`. ### 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. --> This error code was requested to support cases like QNN EP subsystem restart.
Capture the OrtEpFactory pointer directly in plugin allocator release callbacks instead of capturing transient owner objects. This keeps allocator destruction safe when shared or per-session allocator wrappers outlive the object that created the callback.
tianleiwu
enabled auto-merge (squash)
July 17, 2026 21:18
tianleiwu
requested review from
Wayne-Ch,
baijumeswani,
eserscor,
kunal-vaishnavi,
prathikr and
sanaa-hamel-microsoft
July 17, 2026 22:01
Copies the Java setup step from sibling files `windows_webgpu.yml`, `windows_x64_release_xnnpack.yml`, `windows_x64_release_build_x64_release.yml` to `windows_qnn_x64.yml`. This explicitly ensures a compatible Java version, rather than relying on the JDK already installed on the self-hosted runner. Otherwise there may be breakages when switching the runner. I believe this will fix the failures in CI for the Windows x64 QNN CI Pipeline. I'm seeing it fail in both an unrelated PR: #29728 https://github.com/microsoft/onnxruntime/actions/runs/29509983398/job/87780661649?pr=29728 and in commits to main: https://github.com/microsoft/onnxruntime/actions/runs/29546572913/job/87780090142 https://github.com/microsoft/onnxruntime/actions/runs/29499730505/job/87625355397 This is assuming the chain of causation: 1. #29731 switches the runner pool 2. The `windows_qnn_x64.yml` pipeline finds JDK 8 already on the runner (previous runner pool had JDK 11) 3. The Spotless Gradle plugin v7.2.1 isn't compatible with JDK 8 4. onnxruntime Java project can't be configured 5. Build fails This new workflow action hopefully installs a compatible JDK & all will be well. #29753
hanbitmyths
approved these changes
Jul 18, 2026
nenad1002
approved these changes
Jul 18, 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.
This cherry-picks the following commits for the release: