Skip to content

Sync with Microsoft ONNX Runtime - 04092026 - #1284

Merged
hdharpure9922 merged 16 commits into
ovep-developfrom
sync_msft_04092026
Sep 4, 2026
Merged

Sync with Microsoft ONNX Runtime - 04092026#1284
hdharpure9922 merged 16 commits into
ovep-developfrom
sync_msft_04092026

Conversation

@ai-fw-intg

Copy link
Copy Markdown

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

crvineeth97 and others added 16 commits September 2, 2026 13:52
…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

@hdharpure9922 hdharpure9922 left a comment

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

LGTM

@hdharpure9922
hdharpure9922 merged commit acc6951 into ovep-develop Sep 4, 2026
7 of 9 checks passed
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.