Sync with Microsoft ONNX Runtime - 04082026 - #1241
Merged
Merged
Conversation
microsoft#29882) ### Description `GetOutputTensor` in the OpenVINO EP resolved a constant-folded output's OpenVINO friendly name back to an ONNX output by truncating at the first `/` and doing an exact-string lookup. OpenVINO appends a `/sink_port_0` suffix to the friendly name, and ONNX output names may themselves contain `/`, so a name like `D/x/sink_port_0` was truncated to `D` and matched different output named `D`. The constant was then written into the wrong output's slot; with a pre-bound fixed-size output buffer this is an out-of-bounds write. This PR: - Matches the full friendly name first, and only strips the single trailing `/`-suffix if that fails, so names containing `/` resolve correctly. - Adds a byte-size bounds check in `FillOutputHelper` before the `std::copy`, so any residual size mismatch throws instead of overrunning the buffer. - Adds regression tests covering the same-shape mis-route and the mismatched-shape overrun.
### Description Adds the `GRU` operator to the WebGPU execution provider, following the existing WebGPU `LSTM` kernel. Registered for opsets 7-13 and 14+, with float (`T`) and int32 (`T1`) type constraints. Each timestep is computed in two passes because, for `linear_before_reset = 0`, the recurrent term `(r (.) H_prev) * Rh^T` mixes reset-gate values across hidden units and so cannot be produced by a single per-unit thread: - `GruGateProgram` computes the update (`z`) and reset (`r`) gates for the whole `[batch, hidden]` tensor. For `linear_before_reset = 0` it emits `r (.) H_prev` directly; for `linear_before_reset = 1` it emits `r` so the reset is applied after the recurrent matmul. - `GruHiddenProgram` computes the hidden gate and the new state `Ht = (1 - z) (.) h + z (.) H_prev`. Supported: bias, forward / reverse / bidirectional directions, `sequence_lens` masking, the `layout` attribute, `clip`, and both `linear_before_reset` modes. Activations are limited to Sigmoid/Tanh/Relu (as in the WebGPU LSTM kernel). ### Motivation and Context Resolves microsoft#29452. GRU was the natural follow-up to the recently added WebGPU LSTM support, letting models with GRU nodes run on the WebGPU EP. ### Testing Coverage comes from the existing GRU operator tests in `deep_cpu_gru_op_test.cc`, which now also execute against the WebGPU EP (`base_tester` iterates the WebGPU EP when built with `--use_webgpu`). Cases using activations the kernel does not implement (e.g. `LeakyRelu`/`ScaledTanh`) are excluded from the WebGPU run via the shared test helper. The gate math and buffer indexing were cross-checked against an independent ONNX-spec GRU reference for both `linear_before_reset` modes, with and without bias (match to ~1e-16). Validation on real WebGPU hardware is left to CI. --------- Co-authored-by: Isabel Wu <231155141+wuisabel-gif@users.noreply.github.com>
### Description Speeds up MXFP4/NVFP4 QMoE weight dequantization (`LaunchQMoEDequantizeFp4Weights` / `LaunchQMoEDequantizeNvfp4Weights`) by fixing the memory access pattern. The existing scalar kernels map one thread to one output element with `k` varying fastest. Packed weights are stored `[E, K, N/2]` (n-packed) while the output is `[E, N, K]`, so consecutive lanes read packed bytes `packed_n` apart — a separate 32-byte sector per lane for half a byte of payload, plus a 64-bit div/mod per element. Measured at ~5% of HBM peak on a 40-layer NVFP4 MoE prefill (3.8 ms per launch, 307 ms per forward). This PR adds `QMoEDequantizeFp4WeightsVecKernel`, which tiles the *output* instead: - a block covers 64 consecutive rows x 64 consecutive `k`; `threadIdx.x` selects the k group, `threadIdx.y` the row; - each thread emits 8 values as one 16-byte store, so 8 lanes of the same row issue one 128-byte contiguous store (a warp covers 4 rows in 4 requests instead of 32); - `kTileN = 64` rows is exactly 32 packed bytes, so a block fully consumes every packed sector it touches; - the block-scale byte and per-expert global scale are read once per 8 values; - the index decomposition is pure grid arithmetic — no integer division. The mapping matters more than the vector width. An earlier revision that gave each thread 16 (then 32) consecutive `k` of a *single* row reported ~95% of `Max Bandwidth` at only ~26% DRAM throughput in Nsight Compute: bound by memory *requests*, not DRAM. Widening the per-lane store leaves requests-per-byte unchanged, and measurement confirmed no improvement from 16 -> 32 values per thread. One template covers both codecs: `kE4M3Scale` selects the NVFP4 scale codec (Float8E4M3FN, block 16) over the MXFP4 one (Float8E8M0, block 32). Also in this PR: - `DecodeFp4E2M1` is made branch-free by assembling the float bits directly. The previous runtime-indexed local table compiles to a constant-bank load that the hardware replays once per distinct address in a warp, and neighbouring weights rarely share a code. - `DecodeFloat8E4M3FN` is moved earlier in the file (no behavior change) so the shared kernel can use it. - Added the missing `cudaGetLastError` check after the scalar dequantize launches. ### Correctness The tiled kernel reproduces the scalar kernels' indexing exactly (weight nibble selection, block-scale index, output index); the scalar kernels remain the fallback for any shape the tiled mapping cannot cover (odd `n`, `k` not a multiple of 64, or grid dimensions over 65535). Every shape produced by the QMoE quantizers takes the new path, so the existing `test_qmoe_fp4_cuda.py` and `test_qmoe_nvfp4_cuda.py` end-to-end tests cover it. ### Motivation and Context Prefill of NVFP4/MXFP4 MoE models spends a significant fraction of time in weight dequantization; this removes it as a bottleneck.
WebNN has no dedicated `LpNormalization` operator, so decompose it into:
`output = input / max(Lp_norm(input, axis), eps)`
- p==1: `norm = reduceL1(input, {axis})`
- p==2: `norm = reduceL2(input, {axis})`
…float initializers (microsoft#31138) ### Description `get_float_initializer_data` / `set_float_initializer_data` assumed a float initializer always stores its value in the typed `float_data` field. When the value lives in `raw_data` instead, `float_data(0)` reads out of bounds and `set_float_data(0, ...)` writes out of bounds — both undefined behaviour. Both functions now select the field that actually holds the data: use `float_data` when it is non-empty, otherwise read/write `raw_data` (guarded by size checks). ### Testing Adds `openvino_ov_protobuf_utils_test.cc` covering get/set against both `float_data` and `raw_data` backed float scalars. Because the OpenVINO EP is built as a shared-library module with hidden symbols, `ov_protobuf_utils.cpp` is compiled directly into `onnxruntime_provider_test` so the tests can link.
…-graph initializers in subgraph (microsoft#31141) ### Description When a `MatMulNBits` node lives inside a subgraph (e.g. `If` branch) and its quantized weights `B` and `scales` are initializers of the **enclosing** graph, `TryGetConstantInput` fails to resolve those tensors during kernel construction and `PrePack`. On ARM64 with `accuracy_level=4` (KleidiAI path), this causes a `nullptr` dereference → segfault at session initialization. **Root cause:** `KernelRegistryManager::CreateKernel` builds `OpKernelInfo` using `session_state.GetConstantInitializedTensors()`, which only contains the *current subgraph's* constants. Parent-scope initializers are absent, so `TryGetConstantInput(scales)` silently returns null, and `SQ4BitGemmPackQuantBDataAndBlkSum` asserts `QuantBScaleBegin != nullptr`. **Changes:** - **`session_state.h`** — adds `GetConstantInitializedTensorsForKernelCreation()` method + two private fields (`outer_scope_augmented_constant_tensors_`, `outer_scope_augmented_map_built_`). - **`session_state.cc`** — in `FinalizeSessionStateImpl`, between `CleanInitializedTensorsFromGraph()` and `CreateKernels()`, builds an augmented constant-tensor map for subgraph session states by iterating `parent_node->ImplicitInputDefs()` and re-indexing each parent-scope constant into the current subgraph's `OrtValueNameIdxMap`. The parent's own augmented map is consulted (already built, since parent finalizes before child), so arbitrary nesting depth is handled correctly. This map is intentionally **not** used by `PrepackConstantInitializedTensors` to avoid double-prepacking outer-scope tensors. - **`kernel_registry_manager.cc`** — one-line change: `GetConstantInitializedTensors()` → `GetConstantInitializedTensorsForKernelCreation()` in `CreateKernel`. - **`matmul_4bits_test.cc`** — regression test `MatMulNBits.SubgraphParentScopeInitializers` that constructs the exact topology from the bug report (If node, MatMulNBits in both branches, B/scales as parent-graph initializers, `accuracy_level=4`) and asserts session initialization and inference succeed. ### Motivation and Context On ARM64, ORT segfaults during session initialization when loading a model with a `MatMulNBits` node inside a subgraph whose quantized weights and scales are initializers of the parent graph with `accuracy_level=4`. Reported against v1.28.0 on macOS/ARM64 and reproducible with any model following this pattern (e.g. decoder models with quantized weights shared across `If` branches). <!-- START COPILOT CODING AGENT SUFFIX --> - Fixes microsoft#31137 --------- Co-authored-by: copilot-swe-agent[bot] <198982749+Copilot@users.noreply.github.com> Co-authored-by: tianleiwu <30328909+tianleiwu@users.noreply.github.com> Co-authored-by: Hariharan Seshadri <shariharan91@gmail.com> Co-authored-by: github-actions[bot] <41898282+github-actions[bot]@users.noreply.github.com>
…ft#31571) ### Description Adds bfloat16 support to the CUDA collective ops `AllReduce`, `AllGather` and `AllToAll`. Three things blocked bf16 today: 1. **`GetNcclDataType` had no `BFloat16` case**, so any bf16 tensor reaching a collective failed at runtime with `Tensor type not supported in NCCL`. Now mapped to `ncclBfloat16`. 2. **The `AllReduce` CUDA kernel used `DataTypeImpl::AllIEEEFloatTensorTypes()`**, which excludes `BFloat16` because bf16 is not an IEEE format. Replaced with an explicit `float` / `double` / `MLFloat16` / `BFloat16` list. 3. **The op schemas did not list `tensor(bfloat16)`.** Added to all three. Note that the `AllGather` and `AllToAll` kernels already register `DataTypeImpl::AllTensorTypes()`, so for those two the schema was the *only* thing rejecting bf16 — no kernel change was needed. ### Motivation and Context bf16 is the native dtype of most current LLMs, so a tensor-parallel graph built from one is bf16 end to end. Hitting an fp32-or-fp16-only `AllReduce` at every layer boundary forces either a cast pair around each collective (extra kernels and bandwidth on the critical path, twice per layer) or exporting the whole model in fp16. This came up while bringing up a tensor-parallel + expert-parallel DeepSeek-V4 export on 8×H200, where each of the 43 layers issues two `AllReduce` on bf16 activations. ### Testing Validated end to end rather than by unit test, since the distributed collective tests require a multi-GPU host and are opt-in: - 8-rank TP=8 / EP=8 inference of a 284B bf16 model, 2 × bf16 `AllReduce` per layer × 43 layers, producing correct token sequences. - A full 800-sample MMLU-Pro run over that deployment. Happy to add a bf16 case to the existing distributed collective test if reviewers would like it in this PR. ### Note for reviewers `ncclBfloat16` requires **NCCL ≥ 2.10** (mid-2021). CMake currently only version-checks NCCL for the `USE_NCCL_P2P` define (`≥ 2.7`) and does not enforce a floor, so in principle a build against NCCL 2.7–2.9 would now fail to compile this file. Every CUDA 12/13 toolchain ships far newer NCCL, so I left it unguarded rather than adding `#if NCCL_VERSION_CODE >= NCCL_VERSION(2,10,0)` clutter — but let me know if you would prefer the guard or a hard CMake floor.
Fix build error: ``` tmpxft_0000783c_00000000-7_flash_fwd_split_hdim32_bf16_sm80.compute_120.cudafe1.cpp C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v13.0\include\curand_poisson.h(642): error microsoft#20199-D: unrecognized #p ragma in device code [C:\Work\onnxruntime\build_plugin\Release\onnxruntime_providers_cuda_plugin_llm.vcxproj] __pragma(warning(push)) __pragma(warning(disable:4996)) __pragma(nv_diagnostic push) __pragma(nv_diag_suppress 1444) ^ Remark: The warnings can be suppressed with "-diag-suppress <warning-number>" ```
ai-fw-intg
requested review from
Jaswanth51,
ankitm3k,
jatinwadhwa921 and
vthaniel
August 3, 2026 20:35
hdharpure9922
self-requested a review
August 4, 2026 03:53
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.