Skip to content

Sync with Microsoft ONNX Runtime - 25072026#1222

Open
ai-fw-intg wants to merge 16 commits into
ovep-developfrom
sync_msft_25072026
Open

Sync with Microsoft ONNX Runtime - 25072026#1222
ai-fw-intg wants to merge 16 commits into
ovep-developfrom
sync_msft_25072026

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.

GopalakrishnanN and others added 16 commits July 23, 2026 14:24
### 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).
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.