Sync with Microsoft ONNX Runtime - 04092026 - #1284
Merged
Merged
Conversation
…t#32300) ### Description <!-- Describe your changes. --> - Update `pytorch/cpuinfo` from `4628dc060ce4e82345dc166bbac875609db4ff69` to `66ee79c038d70dad9f08705b2c9b3e58f6d8f512`, the latest commit on cpuinfo `main` as of August 27, 2026. - Carry the thread-safe, reference-counted initialization and deinitialization changes from [pytorch/cpuinfo#400](pytorch/cpuinfo#400) as one shared ORT patch used by both FetchContent and vcpkg. - Reset Windows ARM64 cache-population state on each initialization so repeated DLL load/unload cycles cannot reuse stale cache indices. - Scope XNNPACK's cpuinfo references to hardware discovery so XNNPACK-enabled ORT builds do not retain unmatched references during DLL unload. The cpuinfo patch can be removed after pytorch/cpuinfo#400, including the ARM64 reinitialization fix, is merged and ORT updates to a revision containing it. The XNNPACK compatibility patch can be removed after the corresponding lifecycle fix is available upstream. ### 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. --> [microsoft#28245](microsoft#28245) integrated `cpuinfo_deinitialize()` after [pytorch/cpuinfo#387](pytorch/cpuinfo#387) added it upstream. That implementation was later reverted by [pytorch/cpuinfo#411](pytorch/cpuinfo#411) because initialization and deinitialization were not safe for multiple consumers. Pinning the latest cpuinfo `main` without an ORT-side patch would therefore make `cpuinfo_deinitialize()` a no-op again. Carrying the corrected implementation keeps ORT independent of the pending upstream review while preserving safe cleanup during dynamic DLL unload. ### Testing - `onnxruntime_cpuinfo_refcount_test` covers sequential consumers, concurrent consumers, and reinitialization after final release. - Verified this test fails against ORT `main`'s pinned cpuinfo revision (`4628dc060ce4e82345dc166bbac875609db4ff69`) and passes against the patched revision (`66ee79c038d70dad9f08705b2c9b3e58f6d8f512`). - In XNNPACK-enabled builds, the refcount test initializes XNNPACK and verifies hardware discovery does not retain a cpuinfo reference. - `onnxruntime_shared_lib_cpuinfo_dlopen_test` loads a small DLL containing ORT's `CPUIDInfo` and cpuinfo, captures a cpuinfo process-heap allocation, unloads the DLL, and verifies that allocation was released. - Verified the FetchContent patch sequences for the default, Linux, and Windows ARM64/ARM64EC paths, plus the vcpkg patch sequence, apply with zero rejected hunks. --------- Co-authored-by: Vineeth Chelur <vchelur@microsoft.com>
…microsoft#32313) ### Description Pool picks between a serial path, where one invocation loops over the whole kernel window, and a workgroup-cooperative path that splits the window across invocations and reduces in shared memory. The choice was made on output size, which does not describe how much of the GPU either path will fill. This selects on occupancy instead: the serial path is only competitive when the output is large enough to keep the device busy on its own, and the cooperative path is what recovers a large kernel window over a small output. ### Motivation and Context EfficientNet-B0 has a global average pool that reduces the entire feature map to one pixel per channel, which is the shape the old heuristic handled worst: a small output that the serial path cannot parallelise over, with a long reduction behind it. Measured on an NVIDIA TITAN V, driver 560.94: - EfficientNet-B0 native end to end: 3.73 ms -> 1.78 ms - Pool as a share of GPU time: 54.6% -> 11.1% Numerically identical output; this only changes how the reduction is scheduled. --------- Co-authored-by: Ananya Anand <t-anaanand@microsoft.com>
microsoft#32315) ### Description `BufferManager::MemCpy` reports misuse with `ORT_ENFORCE`, which throws: an aliased src/dst pair, an undersized destination and a still-mapped buffer all throw rather than returning a Status. The shared data transfer reaches it through `WebGpuDataTransferImpl::CopyTensorsImpl`, which is a `noexcept` C ABI callback and only handled a returned non-OK Status. The exception escaped a `noexcept` frame, so `std::terminate` took the process down with no message and no stack. This wraps the callback body so any throw becomes an `OrtStatus`, which is what the surrounding C API already expects. The two sibling callbacks are deliberately left alone: `CanCopyImpl` only calls C ABI function pointers, and `ReleaseImpl` returns void, so converting a throw there would mean swallowing it rather than reporting it, which is a separate decision from this one. ### Motivation and Context Not currently reachable from Python, because `copy_tensors` rejects every WebGPU OrtValue outright. It becomes reachable as soon as the supported copies are allowed, which is what microsoft#32074 does, so it is worth fixing at the boundary on its own rather than landing inside a larger change. The regression test for it uses the session-scoped allocation APIs under discussion in microsoft#32074 and is held there rather than duplicated here. Co-authored-by: Ananya Anand <4n4ny4@users.noreply.github.com>
…2392) ### Description CUDA 13.0 moved libcu++, CUB and Thrust from <toolkit>/include to <toolkit>/include/cccl. onnxruntime_providers_cuda.cmake compensates for that inside config_cuda_provider_shared_module, but the plugin EP's cmake never did. The plugin target globs the same host .cc files that include CUTLASS headers (contrib_ops/cuda/llm/cutlass_heuristic.cc and others), and those reach <cuda/std/utility>, so on a CUDA 13 toolkit they failed with cutlass/cutlass.h:40:33: fatal error: cuda/std/utility: No such file or directory Move the include handling - and the CUDA 13.3 cudafe++ header workaround it depends on - out of onnxruntime_providers_cuda.cmake and into the new shared cmake/onnxruntime_cuda_cccl.cmake, then call it from both provider cmakes. The include flags of the in-tree provider build are unchanged. The only adjustments made while extracting the code are: * The 13.0 version check becomes an early return instead of wrapping the function body. * The generated-header directory is pinned to ${CMAKE_BINARY_DIR} rather than ${CMAKE_CURRENT_BINARY_DIR}, which in a function expands in the caller's directory scope. Both are the same directory for the two call sites (cmake/CMakeLists.txt is the top-level list file), but a shared module should not generate files in a caller-dependent location. This matches op_reduction_root in onnxruntime_providers.cmake, the closest analogue: a configure-time generated header tree injected per target. * The entry point is named ort_configure_cuda_cccl, since it generates files in the build tree as well as adding include directories. Comments record the two things that are easy to break later: the SM-specific OBJECT libraries of the plugin target inherit these include directories through a generator expression, which is why calling the function after they are created still covers them; and the CUDA 13.3 workaround is UNIX-only because it is untested on Windows, not because MSVC is known to be unaffected. ### Motivation and Context Trying a plugin build failed with: cutlass/cutlass.h:40:33: fatal error: cuda/std/utility: No such file or directory --------- Co-authored-by: Claude Opus 5 (1M context) <noreply@anthropic.com>
…icrosoft#32316) ### Description `perf_test -i` did not forward provider options to the WebGPU EP, so options passed that way were accepted and then silently ignored. This parses them the way every other EP branch in the file already does and hands them to `AppendExecutionProvider`. The comment also records the short key form, because the fully-qualified `ep.webgpuexecutionprovider.<name>` form gets double-prefixed and dropped without a warning. ### Motivation and Context Both were found while trying to sweep WebGPU EP settings from `perf_test`. The failure mode is quiet in both cases: the run completes, reports a plausible number, and uses the default value. A dispatch-window sweep had to be discarded and re-run after hitting the second one. Co-authored-by: Ananya Anand <4n4ny4@users.noreply.github.com>
… 1.5 hours. (microsoft#32399) ### Description <!-- Describe your changes. --> ### 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. -->
### Description <!-- Describe your changes. --> ### 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. -->
`SkipLayerNormalization` can be decomposed by `LayerNormalization` + `Add`. > output=LayerNorm(X+skip+bias)
### Description ShaderHelper::Init() special-cased shaders whose dispatch grid has a single row of workgroups (dispatch_group_size_y_ == 1 && dispatch_group_size_z_ == 1) by computing global_idx directly from `@builtin(global_invocation_id).x` and workgroup_idx from `@builtin(workgroup_id).x`. This is only correct when the workgroup itself is also 1D (workgroup_size_y == 1 && workgroup_size_z == 1). For a program dispatched as a single row of 2D/3D workgroups (e.g. workgroup_size = (8, 8, 1)), global_id.x omits the contribution of local_invocation_id.y/z that local_invocation_index folds in, so global_idx (and workgroup_idx-derived offsets) come out wrong and the shader reads/writes incorrect elements. Remove the special-cased fast path so every non-indirect dispatch uses the general num_workgroups-based formula, which is correct regardless of workgroup or dispatch shape. Also drop the now-unnecessary is_1d_dispatch distinction from the program cache key (program_cache_key.h/.cc, webgpu_context.cc), since generated shader source no longer varies with dispatch dimensionality. ### Motivation and Context See above.
…oft#31828) ## Summary - release CUDA provider option values against the `jstring` objects they came from - apply the same correction to the TensorRT provider option binding Fixes microsoft#31706 ## Testing - `git diff --check` - A JNI/Gradle build could not be run locally because this Windows environment has no JDK or C compiler. The change is limited to the two provider-option release loops and is ready for repository CI. ## AI assistance disclosure I identified and reproduced the issue, then used OpenAI Codex to assist with the implementation and regression-test work. I personally reviewed the resulting diff and ran the validation commands listed above to verify the fix. Any full-suite or local-environment limitations are documented in the validation section. --------- Co-authored-by: Scott McKay <skottmckay@gmail.com>
### Description Fix the TensorRT 10.11+ shape-tensor optimization profile setup by passing the configured optimum and maximum values to their matching kOPT and kMAX profile selectors. ### 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. --> Swapping OPT and MAX can make an valid profile invalid, causing TensorRT engine construction to fail when the configured optimum and maximum values differ. --------- Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
…icrosoft#28902) ### Description CUDA EP's `Slice` always allocated an output buffer and launched the per-element `_SliceKernel`, even when the output is a contiguous subregion of the input (e.g. `tensor[i:j, :]`). This adds a fast path that copies such slices with a single `cudaMemcpyAsync` (DeviceToDevice) from `input_ptr + offset`, skipping the per-element index decomposition. - **`TryComputeContiguousSliceOffset`** (`core/providers/cuda/tensor/slice.cc`): detects the contiguous case and returns the start element offset. Conditions: all steps == 1; scanning from the rightmost axis, axes are fully included up to the first trimmed (pivot) axis; every axis left of the pivot selects exactly one element. Offset is computed from row-major strides. - **`Slice::CallSliceImp`**: takes the memcpy path when the slice is contiguous, otherwise falls back to the existing kernel. Placed in `CallSliceImp` rather than the shared `ComputeInternal`, so training `SliceGrad` (which overrides `CallSliceImp` with scatter semantics) is unaffected. - **Tests** (`slice_op.test.cc`): added `Slice3D_ContiguousLeadingAxis` and `Slice3D_ContiguousSingleLeadingIndex`, exercising the path across CPU/CUDA EPs. Covered patterns: leading-axis slicing, batch splitting on axis 0, and single-index leading selection. Strided, middle-axis, and inner-axis slices remain on the kernel path. ### Motivation and Context Models with repeated leading-axis slicing (e.g. per-layer embedding extraction) incur one kernel launch per `Slice`. Replacing the trivially-contiguous cases with a single memcpy removes the per-element index computation and launch overhead, matching how `Reshape`/`Squeeze`/`Flatten` avoid copies. Full zero-copy aliasing would require allocation-planner support for offset aliases; this is the lower-risk memcpy alternative noted in the issue. --------- Co-authored-by: copilot-swe-agent[bot] <198982749+Copilot@users.noreply.github.com> Co-authored-by: Justin Chu <justinchuby@users.noreply.github.com> Co-authored-by: Copilot App <223556219+Copilot@users.noreply.github.com>
## Description Fix the CUDA PagedAttention build when FlashAttention is disabled. Two values used only by the FlashAttention path otherwise trigger `-Werror=unused-variable` and `-Werror=unused-but-set-variable`. ## Summary of Changes - Define `flash_block_size_ok` only when `USE_FLASH_ATTENTION` is enabled. - Mark `max_kv_len_lower_bound` unused in builds without FlashAttention. - Leave runtime behavior unchanged for both configurations. ## Testing - Built `onnxruntime_providers_cuda` with `onnxruntime_USE_FLASH_ATTENTION=OFF` and `-Werror`. - Compiled the rebased `paged_attention.cc` directly with the same CUDA Debug configuration. ## Checklist - [x] Tests not required for this build-only fix - [x] No breaking changes - [x] Documentation not required
ai-fw-intg
requested review from
Jaswanth51,
ankitm3k,
jatinwadhwa921 and
vthaniel
September 3, 2026 20:37
hdharpure9922
self-requested a review
September 4, 2026 03:44
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.