Sync with Microsoft ONNX Runtime - 26072026#1223
Open
ai-fw-intg wants to merge 22 commits into
Open
Conversation
### Description - Replace the function-local `unordered_map` used by `GetElementSizeOfTensor` with a directly indexed `constexpr` table. - Validate at compile time that every numeric `TensorProto_DataType` has a nonzero entry while `UNDEFINED` and `STRING` remain unsupported. - Make `GetElementSizeOfTensor` usable in constant expressions and safely return zero for out-of-range enum values. - Add an independent exhaustive `static_assert` test covering every current ONNX tensor data type. ### Motivation and Context The previous map performed dynamic first-use initialization and heap-backed hash lookup for a fixed enum-to-size relation. A constexpr table moves completeness checking to compilation, removes runtime initialization, and reduces the generated framework object and library size. Local Windows RelWithDebInfo measurements, using a stabilized build and forcing only `tensorprotoutils.cc` stale: | Metric | Before | After | Change | |---|---:|---:|---:| | Incremental build | 14.01 s | 12.04 s | -14.1% | | `tensorprotoutils.obj` | 6,500,915 B | 6,217,508 B | -283,407 B (-4.36%) | | `onnxruntime_framework.lib` | 196,933,848 B | 196,570,904 B | -362,944 B (-0.18%) | ### Validation - `lintrunner -a onnxruntime/core/framework/tensorprotoutils.h onnxruntime/core/framework/tensorprotoutils.cc onnxruntime/test/framework/tensorutils_test.cc` - `cmake --build build/Windows/RelWithDebInfo --config RelWithDebInfo --target onnxruntime_test_all --parallel 8` - `onnxruntime_test_all.exe --gtest_filter="TensorProtoUtilsTest.*"` (10/10 passed) --------- Co-authored-by: Gopalakrishnan Nallasamy <gopalakrishnan.nallasamy@microsoft.com>
) ### Description - Make the ONNX `TensorProto_DataType` to ORT C API `ONNXTensorElementDataType` conversion usable in constant expressions. - Preserve the shared numeric prefix and explicitly map the final three values, whose ordering differs after ONNX inserted `FLOAT8E8M0` before `UINT2` and `INT2`. - Add a `consteval` validator requiring a complete one-to-one mapping between the two enums. - Add an independent exhaustive `static_assert` test, including invalid values. ### Motivation and Context The previous conversion was a runtime 26-case switch. More importantly, there was no compile-time guard ensuring future ONNX and ORT enum additions remained complete and bijective. This change turns that compatibility contract into a compilation invariant while preserving invalid-value fallback to `UNDEFINED`. Local Windows RelWithDebInfo measurements, using a stabilized build and forcing only `onnxruntime_map_type_info.cc` stale: | Metric | Before | After | Change | |---|---:|---:|---:| | Incremental build | 7.294 s | 6.535 s | -10.4% | | `onnxruntime_map_type_info.obj` | 765,269 B | 762,123 B | -3,146 B (-0.41%) | | `onnxruntime_framework.lib` | 196,933,848 B | 196,931,236 B | -2,612 B | These measurements are toolchain-specific and primarily guard against regressions; the compile-time compatibility check is the main benefit. ### Validation - `lintrunner -a onnxruntime/core/framework/onnxruntime_map_type_info.h onnxruntime/core/framework/onnxruntime_map_type_info.cc onnxruntime/test/framework/type_info_test.cc` - `cmake --build build/Windows/RelWithDebInfo --config RelWithDebInfo --target onnxruntime_test_all --parallel 8` - `onnxruntime_test_all.exe --gtest_filter="TypeInfoTests.*"` (4/4 passed) --------- Co-authored-by: Gopalakrishnan Nallasamy <gopalakrishnan.nallasamy@microsoft.com>
### Description This updates WebGPU `Reshape` to follow the existing `enable_int64` factory-registration pattern (used by `Unsqueeze`/`Expand`/etc.) instead of static macro registration, and adds explicit int64 WebGPU test coverage for `Reshape`. ### Motivation and Context PR [microsoft#27478](microsoft#27478) added int64 support for the Expand and Unsqueeze (Squeeze) operators in the WebGPU EP. However, the WebNN EP maps a number of those shape-manipulation operators — including Squeeze/Unsqueeze — onto Reshape (WebNN spec doesn't support Squeeze/Unsqueeze). As a result, when a model uses int64 data through these WebNN-mapped ops, execution falls back because Reshape itself does not yet accept int64 tensors, undermining the int64 coverage that microsoft#27478 was intended to enable. Fix microsoft#29756 --------- Co-authored-by: copilot-swe-agent[bot] <198982749+Copilot@users.noreply.github.com> Co-authored-by: Copilot Autofix powered by AI <175728472+Copilot@users.noreply.github.com>
…begin padding (microsoft#29638) ### Description The CUDA MaxPool custom index kernel (`max_pool_with_index.cu`) clamped a negative window start to `0` via `_Max<int64_t>(start, 0)`, discarding the dilation phase. For any axis with begin-pad > 0 **and** dilation > 1, the loop began at input index `0` (padding) instead of the first valid dilated tap, silently producing wrong `Y` values and wrong `Indices` (also corrupting downstream `MaxUnpool`). CPU is unaffected. Repro (`k=2, s=1, pads=[1,1], dilations=[2]`, `x=[1,9,2,8,3,7,4,6,5]`): CUDA returned `Y[0]=1`/`idx[0]=0` vs CPU `9`/`1`. **Changes** - **Kernel** (`onnxruntime/core/providers/cuda/nn/max_pool_with_index.cu`): replace each `X_start = _Max<int64_t>(X_start, 0)` with a phase-preserving advance to the first non-negative dilated tap. `*_end` is computed from the original start, so it stays correct, and the `maxval` seed reads a valid in-bounds tap. ```cpp // instead of: h_start = _Max<int64_t>(h_start, 0); if (h_start < 0) h_start += ((-h_start + dilation_h - 1) / dilation_h) * dilation_h; ``` Applied to `d_start`, `w_start`, `h_start`, so all of 1D/2D/3D are covered. - **Tests** (`onnxruntime/test/providers/cpu/nn/pool_op_test.cc`): un-excluded the CUDA legs from `MaxPool_10_DilationPadding_1d`/`_2d` (previously excluded due to this bug), and added `MaxPool_DilationPadding_1d_Indices` / `_2d_Indices` covering both value and index outputs with CPU as the oracle. ### Motivation and Context Dilated MaxPool on CUDA always lands on this custom-kernel path (`MaxPool<8>` selects it whenever `I != nullptr || !default_dilations`, and `dilation > 1` forces `!default_dilations`), so the wrong-result was unavoidable for dilated + padded MaxPool. Found as an adjacent finding during the AvgPool `ceil_mode` / asymmetric-pad work; the CUDA AvgPool kernel does not share this bug. --------- Co-authored-by: copilot-swe-agent[bot] <198982749+Copilot@users.noreply.github.com> Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
- Add a cooperative subgroup-matrix Gemm kernel for Intel Xe2/Xe3 (f16), mirroring the MatMul path with an alpha/beta*C epilogue and transA/transB support. - Add Gemm edge-case correctness tests.
…soft#29755) ### Description <!-- Describe your changes. --> Skip DXGI device discovery when Win32k system calls are disabled. DXGI internally calls `gdi32!DdQueryDisplaySettingsUniqueness`, which reads from the GDI shared memory section. Under Win32k lockdown this section may not be mapped (depending on the host's gdi32 initialization strategy), causing an access violation. ``` gdi32!DdQueryDisplaySettingsUniqueness+0x7: 00007ff8`30cf45f7 8b8090001800 mov eax,dword ptr [rax+180090h] ds:00000000`00180090=???????? # Child-SP RetAddr Call Site 00 00000072`4b5fc6f8 00007ff8`2b0f583e gdi32!DdQueryDisplaySettingsUniqueness+0x7 01 00000072`4b5fc700 00007ff8`2b0f433f dxgi!CDXGIFactory::SampleAdapters+0xae 02 00000072`4b5fc780 00007ff8`2b0f3f6a dxgi!CDXGIFactory::Initialize+0x10f 03 00000072`4b5fc870 00007ff8`2b0f34b8 dxgi!CreateDXGIFactoryImpl+0x9a 04 00000072`4b5fc8f0 00007fff`a9e08856 dxgi!CreateDXGIFactoryActualImpl2+0x58 05 00000072`4b5fc930 00007fff`a9e075c3 onnxruntime!onnxruntime::`anonymous namespace'::GetDeviceInfoD3D12+0x86 [device_discovery.cc @ 337] 06 00000072`4b5fcc10 00007fff`a9dec8a1 onnxruntime!onnxruntime::DeviceDiscovery::DiscoverDevicesForPlatform+0x323 [device_discovery.cc @ 584] 07 00000072`4b5fd000 00007fff`a9ded12e onnxruntime!`onnxruntime::DeviceDiscovery::GetDevices'::`2'::<lambda_1>::operator()+0x61 [device_discovery_common.cc @ 20] 08 00000072`4b5fd4c0 00007fff`a93b6c46 onnxruntime!onnxruntime::DeviceDiscovery::GetDevices+0x7e [device_discovery_common.cc @ 19] 09 00000072`4b5fd4f0 00007fff`a93b2d7d onnxruntime!onnxruntime::`anonymous namespace'::SortDevicesByType+0x36 [environment.cc @ 781] ``` ### 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. --> WebNN is migrating ORT graph compilation to a sandboxed process with PROCESS_MITIGATION_SYSTEM_CALL_DISABLE_POLICY.DisallowWin32kSystemCalls enabled. This prevents calls through GDI for security hardening. This change has no effect on normal (unsandboxed) processes.
Speed up build with cuda fpA_intB gemm by splitting the huge generated cu files to multiple smaller ones.
### Description CPU GroupQueryAttention now supports FP16 inputs/outputs with INT8 and INT4 quantized KV caches. The FP16 QK and SV paths use architecture-specific SIMD kernels and reuse FP32 accumulation where it is faster, avoiding full cache conversion while keeping FP16 performance near or above FP32. - **MLAS FP16 KV kernels** - Add direct FP16 quantization, QK, and SV APIs. - Dispatch FP16 quantized KV GEMMs to AVX2, AVX-512, and NEON kernels, with a scalar fallback. - Convert FP16 query tiles once where profitable and reuse the cache-friendly FP32 SV accumulation kernel before vectorized FP16 output conversion. - **CPU GQA integration** - Enable FP16 quantized-cache prompt, decode, flash, and non-flash paths. - Preserve the existing FP32 quantized-cache implementation and dispatch. - **Coverage and benchmarking** - Replace FP16 rejection coverage with FP16 INT8/INT4 prompt and decode tests. - Add direct FP16 QK/SV MLAS benchmarks and extend the CPU GQA benchmark with FP32/FP16 and 4/8-thread comparisons. - Document the experiments in `docs/contrib_ops/cpu/gqa_fp16_experiments.md`. ```cpp MlasQKGemmFp16(..., query_fp16, ..., quantized_kv_cache, ..., scores); MlasSVGemmFp16(..., scores, ..., quantized_kv_cache, ..., output_fp16); ``` ### Motivation and Context Quantized CPU GQA previously accepted only FP32 activations. This enables FP16 model execution while retaining quantized KV-cache storage and avoiding full-cache conversion overhead. KleidiAI `matmul_clamp_f16_f16_f16p` was evaluated but is not a direct fit: it requires packed FP16 RHS data and FP16 output, while these APIs consume INT4/INT8 caches and QK produces FP32 scores. Dequantizing and packing the entire cache for every attention call would discard the fused quantized-cache advantage. ### MLAS Microbenchmark Measured with `onnxruntime_mlas_benchmark` on an AMD EPYC 7763 (AVX2), using decode `M=1`, head size 128, 10 repetitions, and a 0.5-second minimum per repetition. Values are median real time. The improvement compares the optimized FP16 kernel with the original FP16 implementation; the last column compares optimized FP16 with FP32 in the final build. | Operation | Quantization | FP16 improvement | Optimized FP16 vs. FP32 | |---|---|---:|---:| | QK, length 512 | S8 per-tensor | 5.9% | 6.1% faster | | QK, length 512 | S4 per-tensor | 8.8% | 1.0% faster | | QK, length 2048 | S8 per-tensor | 6.5% | 6.1% faster | | QK, length 2048 | S4 per-tensor | 9.0% | 0.9% faster | | SV, length 512 | S8 per-tensor | 50.7% | 8.4% slower | | SV, length 512 | S4 per-tensor | 2.3% | 0.9% slower | | SV, length 2048 | S8 per-tensor | 51.7% | 9.1% slower | | SV, length 2048 | S4 per-tensor | 1.5% | 1.0% slower | Across all per-tensor and per-channel cases, QK FP16 improves by 4.7-9.0% and is faster than FP32. SV FP16 improves by 44.9-51.7% for INT8 and finishes within 9.1% of FP32 across all measured quantization modes. ### End-to-End Thread Scaling Measured with `benchmark_gqa_cpu_flash.py` using batch size 1, 16 query heads, 8 KV heads, head size 128, per-tensor INT8 KV cache, and decode sequence length 1. Four-thread runs use `taskset -c 0-3`; eight-thread runs use `taskset -c 0-7`. Each value is the mean of two runs with 50 warmup and 500 measured iterations. | Threads | Total length | FP32 flash (ms) | FP16 flash (ms) | FP16 vs. FP32 | |---:|---:|---:|---:|---:| | 4 | 513 | 0.225 | 0.254 | 13% slower | | 4 | 1025 | 0.414 | 0.426 | 3% slower | | 4 | 2049 | 0.739 | 0.714 | 3% faster | | 4 | 4097 | 1.191 | 1.402 | 18% slower | | 8 | 513 | 0.191 | 0.238 | 25% slower | | 8 | 1025 | 0.377 | 0.289 | 23% faster | | 8 | 2049 | 0.676 | 0.502 | 26% faster | | 8 | 4097 | 0.918 | 0.925 | 1% slower | From 4 to 8 threads, FP16 latency improves by 30-34% at lengths 2049-4097, compared with 8-23% for FP32. The two-run variation reaches 10-18% for the noisiest long-context measurements, so the controlled MLAS microbenchmarks are the primary evidence for the kernel gains. ### Testing - `onnxruntime_mlas_test --gtest_filter='KVQuant.*'` - CPU GQA FP16 INT8/INT4 prompt, decode, flash, and non-flash operator tests - AVX2 and AVX-512 builds - `lintrunner` NEON source diagnostics are clean; an Arm64 build was not available on the benchmark host because it does not have an AArch64 cross compiler or sysroot. --------- Co-authored-by: copilot-swe-agent[bot] <198982749+Copilot@users.noreply.github.com> Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> Co-authored-by: Tianlei Wu <tlwu@microsoft.com>
…osoft#29834) Add a ClipInt64 kernel/shader for opsets 12-12 and 13, so int64 Clip stops falling back to the CPU EP (which aborts session init under the compile-only flow). int64 (like all integer types) is only a valid Clip element type from opset 12 on -- the opset-11 Clip schema constrains T to float types -- so there is no 11-11 registration. The WebNN-inserted index Clip before ArgMax in the SD-Turbo text encoder is int64. The 4-byte templated Clip cannot cover it (sizeof(T)==sizeof(float) static_assert). WebGPU has no native 64-bit integer type; the EP stores int64 as vec2<u32> and operates on the truncated low 32 bits as i32 (shader_variable.cc GetByOffset/SetByOffset). Because Clip is monotonic, ClipInt64 clamps that i32 value (min/max saturated into the i32 range) and sign-extends on write, consistent with the EP's existing int64 semantics. This is a split CL from microsoft#29682 to fix part of microsoft#29756
## Summary ARM64 has no NEON implementation for fp32 rotary embedding. `MlasRopeDispatchNeon` only ever assigns `HRope`, and only when FP16 vector acceleration is available; `SRope` is never assigned, so `MlasRotaryEmbedOneRow<float>` always falls back to the scalar reference implementation on ARM64 regardless of hardware. `test_rope.cpp`'s shared test registration guard reflects this too, excluding `MLAS_TARGET_ARM64` entirely. This adds a NEON fp32 kernel (`RopeKernel_Fp32`, non-interleaved and interleaved variants) to `rotary_embedding_kernel_neon.cpp`/`.h` and assigns it to `d.SRope` unconditionally in `MlasRopeDispatchNeon`, outside the `MlasFp16AccelerationSupported()` guard since fp32 doesn't depend on FP16 vector support. The interleaved path uses `vld2q_f32`/`vst2q_f32` to deinterleave/reinterleave the real/imag pairs directly, instead of the shuffle/permute sequence the AVX2 implementation uses for the same case. `test_rope.cpp`'s ARM64 guard now registers the existing 14 fp32 cases (dim 6/16/24/32/42/64/70 x interleaved). The fp16 half of that guard is left AMD64/RVV-only, since ARM64 fp16 already has its own dedicated coverage in `test_rope_neon_fp16.cpp` — this change is scoped to fp32 only. ## Testing `onnxruntime_mlas_test`, full suite, no regressions: - Apple M1 (macOS): 27975/27975 passed - Neoverse-N1 (Oracle Cloud A1, Ubuntu): 34769/34769 passed `RoPE_fp32/*`: 14/14 passed on both. ## Benchmark (Neoverse-N1, Oracle Cloud A1, median of 3 runs, `RoPE<float>`) | dim | interleaved | before (scalar) | after (NEON) | speedup | |------|-------------|-----------------|--------------|---------| | 128 | no | 514 ns | 24.4 ns | 21.1x | | 256 | no | 1028 ns | 43.1 ns | 23.9x | | 512 | no | 2067 ns | 80.6 ns | 25.6x | | 1024 | no | 4161 ns | 156 ns | 26.7x | | 128 | yes | 264 ns | 44.4 ns | 5.9x | | 256 | yes | 521 ns | 85.0 ns | 6.1x | | 512 | yes | 1035 ns | 165 ns | 6.3x | | 1024 | yes | 2064 ns | 329 ns | 6.3x | Numbers are from Neoverse-N1 only; it's a dedicated instance and gives a cleaner signal than a shared laptop.
…29824) ## Description Adds a native block-scaled **FP4×FP4 (W4A4) grouped-GEMM prefill path** for the `nvfp4` QMoE quantization mode on Blackwell (SM120+). Previously `nvfp4` always dequantized E2M1 weights to FP16/BF16 and ran the dense A16 MoE runner; now prefill shapes route through the native CUTLASS block-scaled tensor-op, while decode/small-M shapes stay on the fused GEMV / dequant fallback. Also fixes the CUDA `no kernel image` (error 209) failure on SM120 by making the LLM object library build native `sm_120a` SASS on Linux. ## Summary of Changes ### Native NVFP4 FP4×FP4 prefill (QMoE) | File | Change | |------|--------| | `onnxruntime/contrib_ops/cuda/moe/moe_quantization.cc` | Enable native NVFP4 CUTLASS path on SM120+ when shape-supported; construct the FP4×FP4 runner plus a dense A16 fallback runner; wire NVFP4 block scales / global scales through `QuantParams::FP4`; route prefill (`num_rows >= ORT_FP4_PREFILL_MIN_TOKENS`) to native and decode to GEMV/fallback; profile the `kFP4` tactic for native prefill. Replaces manual `getenv` parsing with `ParseEnvironmentVariableWithDefault`. | | `onnxruntime/contrib_ops/cuda/moe/moe_quantization.h` | Add native-NVFP4 state (`enable_nvfp4_cutlass_gemm_`, `fp4_prefill_min_tokens_`, `fp4_native_max_tokens_per_expert_`, dense fallback runner). | | `onnxruntime/contrib_ops/cuda/llm/moe_gemm/moe_gemm_kernels_fp4_fp4.cu` | New TU instantiating the FP4×FP4 grouped-GEMM template dispatch. | | `onnxruntime/contrib_ops/cuda/llm/moe_gemm/moe_kernels.cu` | Support the FP4-weight × FP4-activation instantiation. | ### Build / SM120 native SASS | File | Change | |------|--------| | `cmake/onnxruntime_providers_cuda.cmake` | MSVC-gate `EXCLUDE_SM120_REAL` on the LLM object library: Linux now builds native `sm_120a` SASS (fixes CUDA error 209 on real-only arch lists); MSVC keeps the virtual `compute_120` PTX to dodge the CCCL `tcgen05` host-compile failure; `USE_FP4_QMOE` keeps `120-real` everywhere. | | `cmake/onnxruntime_cuda_source_filters.cmake` | Include the new FP4×FP4 source in the CUDA build graph. | ### Routing controls (environment variables) - `ORT_ENABLE_NVFP4_CUTLASS_GEMM` (default 1) — master switch for the native NVFP4 path. - `ORT_FP4_PREFILL_MIN_TOKENS` (default 64) — prefill/decode routing threshold. - `ORT_FP4_NATIVE_MAX_TOKENS_PER_EXPERT` (default 0 = no bound) — upper bound above which prefill uses the dense A16 fallback. ### Tests & docs | File | Change | |------|--------| | `onnxruntime/test/python/transformers/test_qmoe_nvfp4_cuda.py` | Add a native-prefill routing helper and apply a looser tolerance **only** to natively-routed shapes (W4A4 activation quantization), keeping strict bounds for decode/fallback shapes. | | `docs/contrib_ops/cuda/moe_qmoe.md` | Document native prefill accuracy & routing (§9b.5), the LLM object library / SM120 native SASS arch filtering and the CUDA 209 gotcha (§14.2), plus a Limitations note and TOC entries. | ## Testing - `CUDA_VISIBLE_DEVICES=0 python -m pytest onnxruntime/test/python/transformers/test_qmoe_nvfp4_cuda.py` → 18/18 pass on RTX 5060 Ti (SM120). Native-routed shapes (tokens 64/128) show `max_diff ≈ 0.197`; decode/GEMV shapes are near-exact. - Native-arithmetic correctness proof: forcing the same shapes through the dequant fallback with `ORT_FP4_PREFILL_MIN_TOKENS=100000` drops `max_diff` to ≈ `1e-3`, confirming the ~0.2 error is legitimate W4A4 activation quantization, not a kernel bug. - No regressions in the existing MoE/QMoE suites (`test_moe_cuda.py`, `test_qmoe_cuda.py`, `test_qmoe_fp4_cuda.py`). ## Motivation and Context Finish remaining works of microsoft#29697. `nvfp4` QMoE previously had no native execution path, so all shapes paid the dequant-to-A16 cost. This PR enables the Blackwell block-scaled FP4×FP4 tensor-op for prefill while keeping the latency-sensitive decode path on the fused GEMV, and resolves the SM120 `no kernel image` (CUDA 209) failure that occurred when the LLM object library shipped no loadable SM120 image. ## Checklist - [x] Tests added/updated - [x] No breaking changes (native path is opt-out via `ORT_ENABLE_NVFP4_CUTLASS_GEMM=0`) - [x] Documentation updated
…icrosoft#29684) ### Description Adds `onnxruntime/test/python/transformers/benchmark_onnx_attention_vs_gqa.py`: a single-stream decode benchmark comparing ONNX `Attention` (CUDA EP) with contrib `GroupQueryAttention`, as requested by the "Baseline Benchmark Recipe" TODO in microsoft#28352. Five arms at matched config (S_q=1, causal, no RoPE/mask/softcap; Llama-3-8B shape by default): GQA on its XQA / Flash / cuDNN SDPA kernel tiers, ONNX Attention with past/present inputs, and ONNX Attention opset-24 external cache (in-place `TensorScatter` + `nonpad_kv_seqlen`). Includes a `--sanity` cross-arm output-parity check, a `--profile` NVTX mode for Nsight Systems captures, and GPU/driver/wheel provenance in every log. ### Motivation and Context The decode-latency figures in microsoft#28352 came from informal profiling with no reproducible context; this script is the recipe to reproduce them. Benchmark-only change — no ORT source modifications, not collected by CI. Verified on RTX PRO 6000 (SM120): all arms produce identical outputs on identical inputs (max |diff| ≤ 8e-6 fp16); fp16/bf16 sweeps repeatable within ±2 µs. Measured results will be posted on microsoft#28352. *Developed with AI assistance (Claude Code); measurements taken on real hardware and reviewed by the author.* --------- Co-authored-by: Claude Fable 5 <noreply@anthropic.com>
Add non-Windows 1DS telemetry with privacy-safe device identity, test/CI suppression, package integration, and documented runtime controls. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> Copilot-Session: 5a640c3d-b813-4671-84cf-ad98db91b480
### Description - skip_layer_norm_fusion.cc: make the FindPath matcher paths function-local `static const std::vector` (they are compile-time constants). Built once instead of per node iteration, and the destructors run at exit rather than inlined into ApplyImpl, sidestepping a -Wfree-nonheap-object false positive. - test/unittest_util/conversion.h: convert the float<->half / int cast helpers from Eigen Map cast-assignments to std::transform. The Eigen path vectorizes to _mm512_loadu_ps, which GCC 15 mis-flags with -Warray-bounds when inlined into callers passing small fixed-size buffers. Semantics are unchanged. - zipmap_test.cc: take std::initializer_list<T> and materialize the class-list vector inside TestHelper, keeping the vector out of the inlined TEST body where GCC 15 emits a -Wfree-nonheap-object false positive. ### Motivation and Context GCC 15 with -march=native emits several false-positive warnings that fail the -Werror build. Fixed the underlying code so the warnings no longer fire, with no functional change: --------- Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com>
### Description Tile is a data-movement-only operator, so switching the type constraint from WebGpuSupportedFloatTypes() to WebGpuSupportedNumberTypes() enables int32 and uint32 support without any shader changes. ### Motivation and Context The WebGPU EP's Tile kernel only supported float types (WebGpuSupportedFloatTypes() → float, MLFloat16), so int32/uint32 Tile nodes fell back to CPU EP. This hurts performance (integer Tile runs on CPU with extra WebGPU↔CPU copies) and fragments graph partitioning, e.g. [movenet-multipose](https://huggingface.co/Xenova/movenet-multipose-lightning).
### Description `Trilu` was missing in the native C++ WebGPU EP, causing kernel lookup failures for models, e.g. [sd-1.5-text-encoder](https://huggingface.co/microsoft/stable-diffusion-v1.5-webnn/blob/main/text-encoder.onnx). This PR adds a WebGPU `Trilu` kernel in `onnxruntime/core/providers/webgpu/` and wires it into EP registration and provider-targeted coverage. ### Motivation and Context The native C++ WebGPU EP did not implement ONNX `Trilu`, so graphs containing `Trilu` failed provider kernel resolution. This change closes that gap in the WebGPU EP (C++ path under `onnxruntime/core/providers/webgpu/`) and enables those models to execute without the previous “kernel not found” failure. --------- Co-authored-by: copilot-swe-agent[bot] <198982749+Copilot@users.noreply.github.com> Co-authored-by: Copilot Autofix powered by AI <175728472+Copilot@users.noreply.github.com>
…ns (microsoft#29835) Register missing opset ranges for existing WebGPU kernels so these nodes stop falling back to the CPU EP (extra GPU<->CPU copies; session-init abort under the compile-only flow). * AveragePool (pool.cc, webgpu_execution_provider.cc): register opsets 11-18, 19-21 and 22, replacing the single opset-11 registration. ORT clamps an open-ended SinceVersion(11) registration to [11, 18] (not [11, INT_MAX]), so MobileNet (opset 21, AveragePool SinceVersion 19) was not claimed and fell back to CPU; the CPU-only NchwcTransformer (Level3) then amplified the single pool into 3 CPU nodes. * Clip int32/uint32 (unary_elementwise_ops.cc, webgpu_execution_provider.cc): register the 4-byte templated Clip for int32/uint32 at opsets 12-12 and 13, via a dedicated WEBGPU_CLIP_KERNEL_FROM_12 macro. Integer types are only valid Clip element types from opset 12 on (the opset-11 Clip schema constrains T to float types), so there is no 11-11 integer registration. Used by shape/index/mask subgraphs (e.g. CLIP text position ids). * AveragePool count_include_pad divisor (pool.cc): a pre-existing WebGPU pool-kernel bug exposed by the opset-19 registration. With count_include_pad the shader divided by the full kernel size, which is wrong under ceil_mode (the last window can extend past the padded region; those overflow cells must not be counted). Now count is accumulated in the kernel loop and, under count_include_pad, excludes cells beyond the padded region (signed arithmetic, since u32 window indices underflow for begin padding, which is inside the region). This is a split CL from microsoft#29682 to fix part of microsoft#29756
Add uint8 support to the WebGPU Cast in both directions, so uint8 Cast nodes stop falling back to the CPU EP. uint8 was in neither the input (T1) nor output (T2) type constraint, so both casts *to* uint8 (e.g. a terminal bool->uint8 at a graph output) and casts *from* uint8 (back to int32/float/bool) fell back. WebGPU stores uint8 packed as Uint8x4 (4 per u32, lane 0 -> low byte), the same layout already used for bool. * to uint8 (cast.cc, shader_variable.cc): add uint8 to the output (T2) type constraint, a to-uint8 shader expression, and a Uint8x4 packing path in SetByOffset. * from uint8 (cast.cc, shader_variable.cc): add uint8 to the input (T1) type constraint, a Uint8x4 unpacking path in GetByOffset (reads the 4 bytes into a vec4<u32>, inverse of the packing), and include uint8 in is_from_unsigned so uint8->int64 zero-extends. This is a split CL from microsoft#29682 to fix part of microsoft#29756
…icrosoft#29864) ### Description Widens the fp16 direct C output path on the CompInt8 `MatMulNBits` compute op from 4 bit only to 2, 4 and 8 bit. microsoft#29791 added that path for 4 bit, where the int8 kernel writes each N strip into a small per thread fp32 scratch and the strip is converted straight to fp16, so the full fp32 result never lands in a caller buffer. `SQ8BitGemm_CompInt8` and `SQ2BitGemm_CompInt8` had no such branch, so an 8 bit or 2 bit fp16 model still allocated an `M * N` fp32 output buffer and then ran a separate conversion pass over it. The scratch handling and the conversion loop are now a single helper, `CompInt8StripToFp16`, which the 4 bit path was changed to use as well, so there is one implementation rather than three copies. The per thread scratch lives in a non template accessor so a thread keeps one buffer regardless of how many bit widths it runs. The helper calls the kernel exactly once over the whole M range with the scratch as a `CountN` strided buffer, which is what the fp32 branch does, so the converted values are bitwise identical to the fp32 path. `MlasQNBitGemmFp16DirectCOutputSupported` now accepts 2, 4 and 8. ### Motivation and Context Follow up to microsoft#29791, which noted in a comment that the C epilogue is portable and would widen as other bit widths gained the branch. This is that change, so 8 bit and 2 bit fp16 models get the same saving 4 bit already has: one fewer `M * N` fp32 allocation and one fewer full conversion pass over the output. `test_sqnbitgemm_fp16_quant_a.cpp` already checks the fp16 C output against the fp32 result converted to fp16, per bit width, and skips that check when `MlasQNBitGemmFp16DirectCOutputSupported` is false, so widening the gate turns on byte identity coverage for 2 bit and 8 bit without a new test. Locally the fp16 quantize A suite passes 15 of 15, the `*NBitGemm*` tests pass 13535 of 13535, and the full MLAS unit test suite passes 28180 of 28180.
### Description <!-- Describe your changes. --> Fix for device discovery for cards which allow for node and compute partitioning. Currently device discovery relies on pcie device attributes that dont populate sysfs items for certain cards. ### 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. --> Need this to support cards that support XCP functionality. --------- Co-authored-by: Ahsan Saghir <ahsaghir@amd.com> Co-authored-by: Copilot Autofix powered by AI <175728472+Copilot@users.noreply.github.com>
microsoft#29688) ### What When a thread pool's size is not set programmatically (`thread_pool_size <= 0`), size it from `ORT_INTRA_OP_NUM_THREADS` / `ORT_INTER_OP_NUM_THREADS` before falling back to the machine-sized default. Explicit `SessionOptions` / global threading options are completely unaffected. - **Strict parsing** — a non-integer or negative value fails loudly (`ORT_ENFORCE`). - **`= 0` escape hatch** — an explicit `0` opts back into the machine-sized default (useful when a platform sets these but you want host-sized pools). - **Nothing set → current behavior, unchanged.** Implements microsoft#29510. ### Why ORT's default intra-op pool sizes to the host's physical cores (`GetNumPhysicalCpuCores()`). In CPU-limited containers (Kubernetes, Modal, Cloud Run, …) the cgroup CPU reservation is invisible to core detection, so a container reserved N cores on a much larger host gets an N′-sized pool sized to the host — oversubscribing the reservation and degrading tail latency badly under concurrency, with no environment-level way to bound ORT today. The only current fix is passing `SessionOptions` at every `InferenceSession` call site, which is unreachable when sessions are created inside third-party libraries. ### Scope note (re: @skottmckay's feedback on microsoft#29510) The original proposal also honored `OMP_NUM_THREADS` as an intra-op fallback. Per the discussion on microsoft#29510, this PR is deliberately scoped to **ORT-specific variables only** — that keeps the semantics under ORT's control and avoids the "why not `OPENBLAS_NUM_THREADS` / `MKL_NUM_THREADS` next" ambiguity. An `OMP_NUM_THREADS` fallback can be a follow-up if there's demand. ### Tests Adds 4 `ThreadPoolTest` cases (all pin the env vars via `ScopedEnvironmentVariables` so ambient values can't leak): intra + inter sizing from the ORT vars, explicit size wins over env, the `0` escape hatch, and invalid-value-throws. Comparisons are made against a reference pool of the intended explicit size because `DegreeOfParallelism` scales by a granularity factor on hybrid CPUs. I was not able to build/run the ORT test suite locally on this machine, so I'm relying on CI to validate the build + `ThreadPoolTest`.
ai-fw-intg
requested review from
Jaswanth51,
ankitm3k,
jatinwadhwa921 and
vthaniel
July 25, 2026 20:35
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.