Repository navigation
Sync with Microsoft ONNX Runtime - 22102026 - #1355
Merged
Merged
Conversation
### Summary - Add native ONNX opset-28 `SpaceToDepth` support for both `DCR` and `CRD` modes on CPU and CUDA, including internal NHWC kernels. - Version-split `SpaceToDepth` and `DepthToSpace` registrations at opset 28 across CPU, CUDA, DML, JS, and WebGPU; add DML ordered `SpaceToDepth` support and make QNN decline unsupported CRD nodes for fallback. - Add opset-28 CPU mode, shape-validation, non-square/blocksize, and native-vs-context-dependent-function parity coverage. ### Validation - `git diff --check` - Source-level DCR/CRD layout checks against the pinned ONNX 1.23 function bodies. - `build.py --config Debug --build --build_dir build\\windows-spacedepth --cmake_generator "Visual Studio 17 2022" --target onnxruntime_provider_test` is blocked by the base branch's unrelated FP6 `TensorProto` datatype-map static assertion in `tensorprotoutils.h`, before producing a test binary. --------- Co-authored-by: Copilot App <223556219+Copilot@users.noreply.github.com> Co-authored-by: Ganesan Ramalingam <grama@microsoft.com>
This pull request introduces improved validation for input tensors in the Paged Attention CUDA operator, ensuring that sequence lengths and block table values are properly checked before computation. It adds a new comprehensive validation function, integrates it into the main compute path, and provides unit tests to verify correctness. Additionally, a minor bug in shape checking is fixed. **Input validation improvements:** * Added a new function `CheckBlockTableAndPastSeqLensValues` in `paged_attention_helper.h` to validate that `cumulative_sequence_length`, `past_seqlens`, and `block_table` values are within expected ranges, preventing invalid memory accesses and logical errors. * Integrated this validation into the main compute function in `paged_attention.cc`, copying relevant device data to CPU and running the checks before computation proceeds. * Fixed a bug in `CheckSequenceLengthTensors` where the logical condition for validating the shape of `seqlens` was incorrect (`&&` replaced with `||`). **Testing:** * Added a new test file `paged_attention_helper_test.cc` with unit tests covering valid and invalid cases for the new validation logic, ensuring robustness against edge cases. --------- Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> Co-authored-by: copilot-swe-agent[bot] <198982749+Copilot@users.noreply.github.com>
This pull request introduces validation for the `token_offset` input in the CUDA implementation of the PackedAttention operator, specifically for the unfused backend. The main goal is to ensure that invalid `token_offset` values are detected early and reported with a clear error message, improving robustness and debuggability. The change includes both the implementation of the validation logic and new tests to verify its correctness. **Validation improvements for PackedAttention:** * Added a CUDA kernel and a host function (`ValidatePackedAttentionTokenOffset`) to check that all `token_offset` values are within the valid range `[0, B * S)` for the unfused backend, and integrated this validation into `PackedAttention<T>::ComputeInternal`. [[1]](diffhunk://#diff-4e4c1980e908e7c273f52041c3dc123cb7a8c4ee20779fb28b4975911556c9dbR31-R84) [[2]](diffhunk://#diff-450dced704ee2fc03ee3611224808eb8a00671f7587e48599ce17c21d5e7c9e8R16-R21) [[3]](diffhunk://#diff-fd7731431d1e80901e297b2ce7f1018fc7305faf58145cb36733bd19b0aeb9d2R270-R279) * Updated relevant CUDA kernels to handle out-of-range `token_offset` values gracefully by skipping computation or zeroing out invalid output regions. [[1]](diffhunk://#diff-4e4c1980e908e7c273f52041c3dc123cb7a8c4ee20779fb28b4975911556c9dbR115-R117) [[2]](diffhunk://#diff-4e4c1980e908e7c273f52041c3dc123cb7a8c4ee20779fb28b4975911556c9dbR371-R374) **Testing enhancements:** * Added a new unit test (`InvalidTokenOffset_Unfused_CUDA`) to verify that invalid `token_offset` values for the unfused backend are correctly detected and result in a descriptive error message.
…p for NVFP4 QMoE (microsoft#33366) > **Stacked on microsoft#33354 (and microsoft#33094).** The base branch is `tlwu/20261013/nvfp4-qmoe-decode-gemv`, so the diff below is this change only. Once microsoft#33094 and microsoft#33354 merge, this PR will be retargeted to `main`. Opened as a draft until then. ## Description Optimizes the non-GEMV kernels and inter-kernel launch serialization along the offline K-packed (`weights_prepacked=1`) NVFP4 `QMoE` decode path on SM90+: 1. **Single-CTA fused decode routing (`FusedSoftmaxTopKMapsKernel`)**: Replaces `LaunchSoftmaxTopK` + `fusedBuildExpertMapsSortFirstToken` in a single CTA launch for `num_rows <= 4`, `num_experts == 512`, and `k in {1, 2, 4, 6, 8, 10}`. One warp per row reproduces the `SoftmaxTopKMergeKernel<T, 128, 4>` reduction trees operation for operation and emits the `(expert, token)`-sorted expert maps bit-identically (`ORT_FP4_GEMV_FUSED_ROUTE=0` disables). 2. **Single-pass `finalizeMoeRoutingKernel`**: Sizes the block to one thread per 128-bit column vector (up to 1024 threads instead of a fixed 256) so hidden dimensions up to 4096 (such as Qwen3.8's `hidden=2560`) reduce in a single pass, and issues batched loads of up to 16 expert rows together while preserving top-k accumulation order (`7.7 us -> 4.0 us` per layer on H200). 3. **Programmatic dependent launch (`PDL`) overlap of `FC1 -> FC2 -> finalize`**: Launches FC2 and `finalizeMoeRoutingKernel` with `cudaLaunchAttributeProgrammaticStreamSerialization` (`ORT_FP4_GEMV_PDL=0` disables; `ORT_FP4_GEMV_FC1_PDL=1` optionally enables PDL for FC1 behind fused routing). FC1 triggers dependents at start, FC2 prefetches its first batch of weights/scales before waiting on FC1's output, and `finalizeMoeRoutingKernel` stages routing metadata into shared memory before waiting on FC2's output (`~2.2 us` saved per layer on H200, preserved under CUDA Graph capture). All three optimizations preserve exact floating-point evaluation order and produce **bit-identical** outputs. ## Summary of Changes | File | Change | |------|--------| | `onnxruntime/contrib_ops/cuda/moe/qmoe_fused_route.cuh` | New `FusedSoftmaxTopKMapsKernel` single-CTA fused softmax + top-k + expert-map kernel (`E=512`, `rows <= 4`, `k in {1,2,4,6,8,10}`) | | `onnxruntime/contrib_ops/cuda/moe/qmoe_kernels.h` / `.cu` | `LaunchFusedRouteMaps` declarations and launcher | | `onnxruntime/contrib_ops/cuda/llm/moe_gemm/moe_kernels.cu` / `moe_util_kernels.h` | Single-pass `finalizeMoeRoutingKernel` (`blockDim.x` up to 1024, `FINALIZE_LOAD_BATCH=16`) and `early_routing_loads` PDL support | | `onnxruntime/contrib_ops/cuda/llm/moe_gemm/moe_gemv_fp4.h` / `.cu` | `griddepcontrol.wait` / `griddepcontrol.launch_dependents` in `MoeGemvFp4DecodeKernel` and `programmatic_dependent_launch` attribute in `LaunchMoeGemvFp4RawNPacked` (`ORT_FP4_GEMV_PDL`) | | `onnxruntime/contrib_ops/cuda/moe/moe_quantization.cc` | Wire `LaunchFusedRouteMaps` and `fc2_finalize_pdl` into `QMoE::ComputeInternal` | | `docs/contrib_ops/cuda/moe_qmoe.md` | Document `ORT_FP4_GEMV_FUSED_ROUTE`, `ORT_FP4_GEMV_PDL`, and `ORT_FP4_GEMV_FC1_PDL` | ## Testing Validated on NVIDIA H200 (`sm_90a`, CUDA 13.0): - `python -m pytest onnxruntime/test/python/transformers/test_qmoe_nvfp4_cuda.py`: **178 passed, 8 skipped, 0 failed** (including `test_nvfp4_offline_row_major_decode_gemv_bit_exact` for FP16 and BF16). - End-to-end Qwen3.8-Flash-Next NVFP4 (`v12b` 3-session CUDA-graph + MTP engine, 48 MoE layers + 1 MTP layer): - **Bit-exactness**: `engine_gate.py --prefill steps` passes **6/6** with 100% identical token sequences (`256/256` tokens on `text2048` and `prose2048`) compared against the baseline without fused routing / PDL. - **End-to-end decode speed (`bench_engine_mtp.py`, 256 generated tokens, median of 3 runs)**: | Context | Before (`microsoft#33354` only) `ms/cycle` | With this PR `ms/cycle` | Cycle speedup | Decode throughput (`tok/s`) | |---:|---:|---:|---:|---:| | 2,048 | 8.274 ms | **7.657 ms** | **-0.617 ms (-7.5%)** | **244.98 tok/s** (was 227.64) | | 8,192 | 8.164 ms | **7.836 ms** | **-0.328 ms (-4.0%)** | **232.49 tok/s** (was 225.81) | | 32,768 | 8.352 ms | **8.027 ms** | **-0.325 ms (-3.9%)** | **237.25 tok/s** (was 229.97) | - `lintrunner` on all touched files: no issues. --------- Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
…icrosoft#33501) ## Description - Update `transformers` requirement pins from `5.3.0` to `5.10.1` across the Python transformer tooling and test `requirements.txt` manifests. - Update `google.golang.org/grpc` to `v1.84.0` in the FlatBuffers Go greeter example `go.mod` files (`grpc/examples/go/greeter/{client,models,server}/go.mod`) via `cmake/patches/flatbuffers/flatbuffers.patch`. ## Motivation and Context Keep Python `transformers` and third-party Go module dependency manifests up to date with recent patched releases. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
This commit packages the D3D12 Agility SDK runtime DLLs in both the standard ONNX Runtime Python wheel with the built-in WebGPU EP and the standalone WebGPU Plugin EP Python wheel.
Dawn updates currently require coordinating the archive pin, dependency hashes, SwiftShader revision, and local patches. Add a single command to prepare, apply, and validate these changes together: ```sh python plugin-ep-webgpu/tools/dawn/update_dawn.py ``` The command defaults to the newest dated Dawn tag. Use `--revision <tag-or-full-commit>` to select a specific version. It records the concrete revision and checksum, updates the reviewed dependency manifest and Linux test SwiftShader pin, and refreshes active patches in CMake order. Refreshed patches must reproduce the expected source with zero fuzz or offsets. The command verifies written contents and automatically validates the manifest against the selected checkout. Unresolved patches or dependency changes needing classification stop the update; writing or final-validation failures trigger restoration of every attempted file, continuing after individual restoration errors and reporting them alongside the original failure. Reports and a diff are saved under `build/dawn-upgrade/`, and repeating the same update makes no further edits. Dependency classifications and the independently packaged DXC release are preserved. DXIL handling is outside the scope. A manifest-only updater remains available for focused maintenance. This PR contains eight tooling, test, and documentation files. Build-inventory generation and packaging-pipeline integration are in microsoft#33448. Dawn version updates generated during local validation are separate from this PR. Validation: - 49 tests and 29 subtests passed with ordinary and aliased Windows temporary paths, including real GNU patch application, latest-tag resolution, metadata preservation, repeat runs, persistent write failures, partial writes, and multiple rollback failures. Recovery was also verified with a real read-only manifest. - Ruff lint/format and Git whitespace checks passed. - An isolated update from `v20261013.193955` to `v20261016.190333` changed nine files and passed automatic manifest and strict patch validation. A second invocation changed zero files. - Earlier pinned comparisons reproduced microsoft#33390's two changed files exactly and microsoft#32686's dependency identities and revisions. The command automates source updates and metadata/patch validation. Build, package, and GPU tests remain necessary to establish runtime compatibility for each Dawn update.
…oft#33492) ### 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. -->
This pull request enhances the robustness and correctness of ONNX Runtime's ORT format model loading, particularly around edge validation in computational graphs. It introduces stricter checks for edge consistency, improves error handling, and significantly expands test coverage to prevent invalid graph structures from being loaded. **Edge validation and canonicalization improvements:** * The `Node::LoadEdgesFromOrtFormat` method was refactored to enforce stricter validation of edge indices, ensure control edges are handled correctly, and guarantee that edges only connect matching `NodeArgs`. This includes checks for out-of-range indices, mismatched arguments, and correct handling of control edges. (`onnxruntime/core/graph/graph.cc`, `include/onnxruntime/core/graph/graph.h`) [[1]](diffhunk://#diff-e231a92b40d89409cc8e82436be0a15bc87ef95c93b303b9feaeab6e50c8835cL892-R955) [[2]](diffhunk://#diff-aaea1507ec81a94c72a1fa72ce320df712156b665f7798573be3f7e439bb4c37L523-R523) * After edge loading, the graph now checks that each destination argument slot has at most one producer, preventing invalid graphs with multiple producers for the same input. (`onnxruntime/core/graph/graph.cc`) **Test coverage expansion:** * New helper functions and tests were added to rigorously verify edge validation logic. These tests cover scenarios such as out-of-range source/destination indices, mismatched node arguments, canonicalization of reciprocal edges, control edge handling, and detection of multiple producers for a single destination slot. (`onnxruntime/test/framework/ort_model_only_test.cc`) [[1]](diffhunk://#diff-091729c3958c3945dfe649e7a1c3313b941a3da9a27662f047cc35c948196e49R90-R160) [[2]](diffhunk://#diff-091729c3958c3945dfe649e7a1c3313b941a3da9a27662f047cc35c948196e49R302-R369) These changes collectively ensure that only valid and well-formed graphs are accepted from ORT format models, reducing the risk of subtle bugs or crashes due to malformed edge definitions. --------- Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
### Description This change adds missing semantic guards in three training graph optimizers: - QDQFusion uses the ONNX UINT8 default when zero point is omitted, leaves shared initializers for graph cleanup, skips a QuantizeLinear whose output is public, and only fuses when the Q/DQ scale, zero-point inputs, and effective axis match. - TransposeReplacement skips the optional default-permutation form instead of indexing a missing attribute. - ScaledSumFusion keeps the original unscaled third input, including graph inputs and outputs of non-scale producers. Regression tests cover each case and retain the existing successful fusion paths. ### Motivation and Context The affected rewrites could change valid model semantics or fail while optimizing valid graphs. Fixes microsoft#32966 Fixes microsoft#32967 Fixes microsoft#32968 Fixes microsoft#32969 Fixes microsoft#32977 Fixes microsoft#32980 ### Testing Built `onnxruntime_test_all` with `--enable_training` in Release mode. After the Q/DQ parameter guard was added, 10 focused tests from two suites passed, covering the new regressions and neighboring successful fusion paths. --------- Signed-off-by: Yuhx141 <155034285+Yuhx141@users.noreply.github.com> Co-authored-by: xadupre <xadupre@microsoft.com> Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
…#33042) ### Description <!-- Describe your changes. --> Add reader-only plumbing for the hardware-neutral session option `session.kernel_dispatch_policy`. The intent is generic across operators and EPs. `KernelDispatchPolicy` and the shared, header-only `TryParseKernelDispatchPolicy` live in `core/framework/kernel_dispatch_policy.h`, with no CUDA or graph dependencies. CUDA GQA is the first consumer, not the owner of the policy. For valid values, **no policy changes runtime dispatch, workspace estimates, or partitioning** in this PR. `latency`, `memory`, and `safe` remain intent metadata for future policy resolution. In particular, `memory` does not reduce the reservation, and `safe` does not select a bounded fallback or provide a no-OOM guarantee. #### Reader contract - Unset, empty, and explicit `auto` resolve to `Auto`. The other valid tokens are exactly `latency`, `memory`, and `safe`. - Unknown, differently cased, whitespace-only, and whitespace-padded non-empty values fail session initialization with `INVALID_ARGUMENT`; they never silently resolve to `Auto`. Validation runs even when partitioning is disabled. - The accountant factory retains the validated raw optional string without changing other settings. The CUDA consumer enforces successful parsing rather than silently accepting an invalid configuration. - `GetGroupQueryAttentionWorkspaceEstimateConfig` accepts `dispatch_policy = KernelDispatchPolicy::Auto` alongside the envelope and initializes both fields in one place. The node-estimator overload and CUDA test bridge forward the policy through that helper. - Setting the policy alone does not enable partitioning or non-windowed GQA estimation without a KV envelope. - This replaces the earlier, unreleased attention-specific option name; there is no compatibility alias. #### Integration and regression coverage Preserve microsoft#33014/microsoft#33016's envelope validation against `[0, INT32_MAX]`, invalid-value diagnostics, and fail-closed GQA placement when no proven estimate is available. Reuse the exported graph-config helper instead of duplicating node parsing. Keep envelopes, eligibility, and sizing local to GQA; other operators and EPs can adopt the shared intent independently. Framework regressions cover valid-token parsing, rejected tokens (including `Memory`, `SAFE`, `memroy`, and ` memory `), unchanged parser output on failure, validation with partitioning enabled and disabled, raw-token retention, unset defaults, and policy-only configuration not enabling partitioning. Existing KV-envelope boundary tests remain intact. CUDA metadata regressions cover default and explicit policy forwarding through the node-config helper/test bridge, identical bytes/backend sets for windowed and non-windowed estimation, and unchanged decline without a non-windowed envelope. #### Validation for this review update - Changed-file lintrunner/clang-format check: passed for all 10 edited C++ files. - `git diff --check`: passed. - No full ORT build or C++ regression suites were run locally for this update, as requested. Earlier-head build evidence does not validate this head; full framework/CUDA validation is left to CI. ### 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. --> Provide a generic policy representation and reader without coupling other consumers to GQA or CUDA. Rejecting typos now prevents future enforcement (including microsoft#33316's committed policies) from silently falling back to `Auto` and bypassing the requested restrictions. This PR builds on microsoft#33016, which is now merged, and targets `main`. Committed policies, operator/EP-specific resolution, route restrictions, runtime coverage/envelope enforcement, bounded fallback, and reduced policy-bound estimates remain separate follow-ups. --------- Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> Co-authored-by: Chi Lo <lochi@microsoft.com> Copilot-Session: b95559c4-5cbe-4787-8dd7-ffa45a2157f9 Copilot-Session: 24b20465-fd4d-4a64-8174-aa0622c5164d
…ing (microsoft#33493) Description: ### Description <!-- Describe your changes. --> Add `com.microsoft::ConvStateReplay` v1, a dedicated WebGPU contrib operator that reconstructs rolling convolution history from an incoming state snapshot and a retained prefix of captured inputs. It performs copy/selection only, without convolution arithmetic. - Supports FP16 and FP32, batch/channel dimensions, configurable history width and capture capacity, and all retained-prefix counts from zero through capacity, including counts greater than the history width. - Accepts full source, capture, and destination backing tensors plus CPU `int64[8]` metadata describing byte offsets, geometry, capacity, and retained count. - Uses `ProgramInput::BufferView` / `ProgramOutput::BufferView` with integer-word access and logical offsets, including two-byte-aligned FP16 subviews. Integer-only copies preserve signed zero, subnormals, and NaN payloads exactly; masked boundary writes preserve neighboring elements. - Validates geometry, alignment, bounds, binding limits, and destination aliasing, and returns after checked GPU completion. The output must be explicitly prebound to the inactive destination backing. - Adds schema/kernel registration, operator documentation, and focused tests against an independent CPU rolling-history reference, with CPU fallback disabled. Validation: built the native WebGPU provider and ran the following from `build/webgpu-pipeline-lifetime/Release` using Mesa lavapipe: XDG_RUNTIME_DIR=/tmp VK_ICD_FILENAMES=/usr/share/vulkan/icd.d/lvp_icd.x86_64.json ./onnxruntime_provider_test --gtest_filter='*ConvStateReplay*:*GatedDeltaNetStateReplay*' Result: **31 tests passed, zero failures or skips**—20 new convolution replay tests and 11 existing GatedDeltaNet replay tests. Coverage includes every retained prefix, batch/channel indexing, segmented buffers, arbitrary offsets, adjacent-state preservation, every FP16 payload pattern, FP32 special values, zero/boundary cases, and invalid metadata/bounds/aliasing. Repository lint passed and generated contrib documentation was verified. ### 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. --> During speculative-decoding partial acceptance, the current GenAI WebGPU fallback downloads convolution state and captured inputs, reconstructs the accepted-prefix history on CPU, then uploads it into the inactive state bank. This fallback accounted for approximately 7.31 seconds of overhead in the reported long Qwen 3.8 DFlash2 width-1 benchmark. Convolution replay only needs to select the final history window from the incoming state followed by the retained captured inputs. This operator provides the GPU primitive needed to eliminate that round trip in a subsequent GenAI integration. This PR is limited to the ORT operator, registration, documentation, and tests. It does not integrate GenAI, change GatedDeltaNet replay behavior, modify the tensor-transfer API, or change production kernel-selection policies. End-to-end performance improvement has not been measured; graph capture is not supported. --------- Co-authored-by: github-actions[bot] <41898282+github-actions[bot]@users.noreply.github.com>
### Description Check whether each input Transpose produces a graph output before folding it into Gemm. This preserves the model's public outputs while allowing the other, internal input Transpose to be fused. The regression tests cover either input being public with all existing `transA`/`transB` combinations, and both inputs being public when an output Transpose triggers the rewrite. They check the surviving nodes, input/output names, and transpose attributes. ### Motivation and Context Fixes microsoft#33441. `SatisfyCondition` can accept Gemm because one input Transpose is safe to fuse. `Apply` then independently processes both inputs, but previously did not repeat the graph-output check. A public input Transpose could therefore be removed, leaving a declared model output without a producer. I reproduced the report with ONNX Runtime 1.31.0 on Windows x64. The unoptimized graph and the pass-disabled control match the reference result; the optimized session fails to resolve the public output. ### Validation - Built current main (`a7a1172c`) and the patch with MSVC on Windows x64, Release configuration. - With the new regressions and the original implementation: 9 existing GemmTransposeFusion tests passed, and the 2 new tests failed. After the fix: all 11 passed. - `GraphTransformationTests.*`: 442 passed, 17 skipped because their GPU execution providers are unavailable in this CPU build. This includes the 11 focused tests. - The original issue's model now runs with optimization enabled and both outputs exactly match `ReferenceEvaluator`. - A local 64-case model matrix covers public-output placement, existing transpose flags, bias, and output transposition. All outputs from the newly built Python module exactly match NumPy. The released 1.31.0 build fails 36 of these cases, while its unoptimized and pass-disabled controls match. - `lintrunner` for both changed files and `git diff --check` pass. AI assistance was used for implementation, regression design, and validation.
…microsoft#28971) ### Description The CUDA `LayerNormalization` / `RMSNormalization` kernels only support broadcasting scale/bias through the encoded `broadcast_param` fast-path; they do **not** implement the generic NumPy-broadcasting path the CPU kernels use. `LayerNormHelper::GetBroadcastParam` required `scale.rank == X.rank` to emit a fast-path value, so a rank-reduced but NumPy-broadcastable scale — e.g. scale `(S, W)` for input `(B, S, W)` with `axis=2` — returned `kLayerNormInvalidInput`. On CPU this falls back to the correct generic path, but the CUDA kernel silently used `broadcast_param = -1`, computing offset `norm_size * (x_row % 1) = 0` and applying `scale[0:W]` to **every** row → wrong results (no error raised). ### Impact This breaks Mamba2 / hybrid SSM models (e.g. **NemotronH**), whose gated `RMSNorm` uses a per-group scale `(n_groups, head_dim)` broadcast over the batch dimension. Such models produce correct logits on CPU but a wrong argmax on CUDA. Minimal repro (fp16): ```python # X:(5,8,960), scale:(8,960), axis=-1 # CPU matches numpy reference; CUDA was off by ~13 (max-abs) before this fix. ``` ### Fix Right-align the scale/bias shapes to `X`'s rank inside `GetBroadcastParam` (NumPy-style, prepending 1s) before matching the `(B,S,...)`, `(B,1,...)`, `(1,S,...)` fast-path layouts. A rank-reduced `(S, W)` scale now maps onto the existing `(1, S, ...)` fast-path and yields the correct per-row offset on CUDA — identical to an equal-rank `(1, S, W)` scale. - Equal-rank shapes are unchanged (right-alignment is the identity). - The CPU generic path is unaffected. ### Testing Adds a CPU+CUDA regression test `RMSNorm_Scale_RankReduced_Sx3_Axis2`. Verified end-to-end: the NemotronH prefill argmax now matches between CPU and CUDA, and a broadcast sweep (RMSNorm + LayerNorm, multiple scale shapes, CPU vs CUDA vs numpy) passes at float epsilon with no regressions. --------- Signed-off-by: Justin Chu <11205048+justinchuby@users.noreply.github.com> Co-authored-by: Justin Chu <11205048+justinchuby@users.noreply.github.com> Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
…ization (microsoft#33357) ### Description Fix variance cancellation in the fused CUDA BERT `SkipLayerNormalization` and `EmbedLayerNormalization` kernels. This is a standalone change against `main`, independent of microsoft#33312. The shared regular LayerNorm helpers currently compute variance as `E[x*x] - E[x]*E[x]`. FP32 accumulation does not make that subtraction stable for high-offset, low-variance inputs. The result can be substantially wrong or lead to NaN outputs. Ordinary CUDA `LayerNormalization` already uses Welford/Chan; this change addresses the separate fused helpers. - Replace raw moments with FP32 Welford updates and Chan/CUB block reduction, including count-zero identities for inactive lanes. - Accumulate and normalize relative to a row origin. Keep the mean in that shifted coordinate system instead of reconstructing a rounded high-offset FP32 mean. - Apply the fix to scalar and vectorized skip kernels and the embedding kernel. - Preserve native-dtype input/bias/skip and word/segment/position addition order and rounding, optional residual/embedding-sum outputs, masks, and simplified RMSNorm arithmetic. - Avoid reading a row origin for an empty hidden dimension; cover the existing CUDA SkipLayerNormalization empty-row behavior. No operator schema, public launch interface, workspace, provider option, or ordinary LayerNormalization changes. ### Motivation and evidence For the actual FP32 values obtained from `[1000.01, 1000.02, 999.99, 1000.0]`, centered FP64 population variance is approximately `0.00012524426`. In the preceding isolated CUDA reproduction, with epsilon `1e-5`, the raw-moments skip kernel returned NaNs at widths 4 and 768. The embedding kernel returned NaNs at width 4 and substantially incorrect finite values at width 768. The new FP32 regression tests fail on the unchanged upstream kernels and pass with this fix. No variance clamp, larger epsilon, or relaxed FP32 tolerance is used. ### Validation Tested on H100 (SM90), CUDA 13, using the compiled CUDA provider and C++ operator test executable. ```bash cmake --build build/Cuda13/Release \ --target onnxruntime_providers_cuda onnxruntime_provider_test -j12 cd build/Cuda13/Release CUDA_VISIBLE_DEVICES=0 ./onnxruntime_provider_test \ --gtest_filter='SkipLayerNormTest.*:EmbedLayerNormTest.*' ``` **44 passed, zero skips.** All six new CUDA regression tests executed. The new tests explicitly disable CPU EP fallback. Coverage includes FP32/FP16/BF16 skip normalization and FP32/FP16 embedding normalization; hidden widths 1, 3, 4, 5, 128, 129, 320, 768, 769, 8192 and 8193; constant rows; padded blocks; scalar/vectorized dispatch; multi-element-per-thread reduction; optional bias/segment inputs; masks; and optional sum outputs. Low-precision residual-cancellation cases verify that native addition rounding is retained. The oracle uses centered FP64 arithmetic over the actual native-rounded fused sums. FP32 output bounds remain absolute/relative `1e-4`; low-precision outputs use the repository's existing dtype-aware tolerances. Optional sum outputs require exact equality. `clang-format --dry-run --Werror` and `git diff --check` pass. Validated source commit: `5ff00337b71aa939e023f8342899575e5dc6ba56`, based on `07daf405c788c300aa9fd2988104b29bcbdede23`. CUDA provider SHA-256: `01826247f90f3b6a64f5fde8b9dcf6bf6e62728af34614494b9949029c38c460`. ### Scope and related work Related: microsoft#32439 proposes Welford as part of stacked CUDA statistics-output and addition-contract work. This PR fixed the onboarding precision issue of Cosmos3 model, onnxruntime/mobius#466 Signed-off-by: rui-ren <15321482+rui-ren@users.noreply.github.com> Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
# Summary Optimize ONNX Runtime's allocation planner to avoid quadratic free-buffer scans during session initialization. The existing planner maintains a linear list of reusable buffers and searches it when assigning memory to intermediate tensors. For large graphs containing many incompatible buffer candidates, repeated linear scans can result in **O(N²) allocation-planning overhead**. This PR introduces a lightweight, compatibility-keyed free-buffer index to reduce unnecessary candidate scans while preserving existing memory reuse semantics. **Key results:** - **283.38× faster CPU allocation planning** on the synthetic stress workload. - **297.03× faster CUDA allocation planning** on the same workload. - **2.91× faster CPU session initialization**. - **2.75× faster CUDA session initialization**. - No material peak host-memory regression observed. These improvements apply to the targeted pathological workload and should not be generalized to all models. ## Motivation ONNX Runtime reuses memory from intermediate tensors whose lifetimes have ended to reduce memory consumption and avoid unnecessary allocations. Previously, the allocation planner searched a linear free-buffer list to identify compatible reuse candidates. For a graph with many intermediate tensors and incompatible buffer shapes, this can result in repeated examination of buffers that cannot be reused. A motivating large decoder initialization exhibited: | Metric | Observed | |---|---:| | Free-list candidate visits | 150,994,944 | | Successful buffer reuse matches | 6 | This indicates substantial unnecessary search work during allocation-plan construction. The objective is to eliminate this pathological scanning behavior without changing which buffers are eligible for reuse or how execution memory is assigned. ## Implementation Introduce a private `FreeBufferIndex` for non-minimal builds. The index uses open addressing with linear probing to map compatibility keys to free-buffer chain heads. Instead of repeatedly traversing the entire free-buffer list, the planner first locates the relevant compatibility group. Key implementation properties: - **Indexed lookup:** Expected O(1) compatibility-group lookup under normal hash-table conditions, rather than repeatedly scanning unrelated free buffers. - **Full-word keys:** Hashes the complete `size_t` key representation to improve distribution for aligned values. - **Collision safety:** Uses full-key equality checks and linear probing to resolve hash collisions. - **Preserved reuse behavior:** Existing compatibility checks and buffer-selection ordering remain authoritative. - **Explicit occupancy tracking:** An empty chain head does not remove a registered key or break a probe chain. - **Bounded load factor:** The table grows before exceeding approximately 75% occupancy. - **Checked capacity growth:** Uses `SafeInt<size_t>` to prevent silent capacity arithmetic overflow. - **Minimal-build compatibility:** `ORT_MINIMAL_BUILD` variants retain the original linear planner to avoid additional binary-size overhead. The optimization changes how candidate buffers are located, not the memory reuse rules themselves. No new runtime configuration flags or execution-provider-specific allocation semantics are introduced. ## Performance Evaluation Benchmarked baseline (`main`) against the optimized implementation. All reported measurements are **medians of 10 runs per build**. ### 1. Synthetic Stress Workload The workload consists of: - 8,192 independent four-Cast chains. - 32,768 total nodes. - Different tensor shapes across chains. - Graph optimizations disabled to preserve the stress workload. This workload exercises the original planner's pathological behavior when many free buffers are incompatible with subsequent allocation requests. | Metric | CPU main | CPU optimized | CUDA main | CUDA optimized | |---|---:|---:|---:|---:| | Allocation/execution-plan construction | 26,723.73 ms | 94.30 ms | 25,533.78 ms | 85.96 ms | | Complete session initialization | 40,914.74 ms | 14,079.38 ms | 39,827.44 ms | 14,494.00 ms | | Peak host RSS | 233.81 MiB | 233.79 MiB | 512.63 MiB | 513.02 MiB | | **Planner speedup** | — | **283.38×** | — | **297.03×** | | **Session initialization speedup** | — | **2.91×** | — | **2.75×** | The results demonstrate that reducing incompatible free-buffer scans substantially improves allocation-planning scalability. More importantly, the planner improvement translates into measurable end-to-end session initialization savings. ### 2. Control Workloads and Decoder Model Additional benchmarks were conducted to determine whether the optimization introduces regressions or provides benefits outside the pathological workload. The following measurements are complete session initialization medians. | Workload | CPU main → optimized | CUDA main → optimized | |---|---|---| | Same-shape control: 8,192 chains | 13,916.98 → 14,083.31 ms | 14,527.28 → 14,565.04 ms | | Random-weight 32-layer Llama decoder | 349.77 → 349.42 ms | 803.88 → 798.69 ms | **Observations:** - The same-shape control does not exhibit the large improvement seen in the different-shape stress workload. - The CPU same-shape control is approximately 1.2% slower, while the CUDA difference is approximately 0.3%. Additional variance analysis would be needed to determine whether these differences are statistically meaningful. - The random-weight Llama decoder shows effectively unchanged initialization performance. - These results indicate that the optimization specifically addresses pathological free-buffer scanning rather than universally accelerating model initialization. ### 3. Peak Host-Memory Usage | Configuration | Main | Optimized | Difference | |---|---:|---:|---:| | CPU | 233.81 MiB | 233.79 MiB | −0.02 MiB | | CUDA | 512.63 MiB | 513.02 MiB | +0.39 MiB | No material peak host RSS regression was observed in the synthetic stress workload. The additional compatibility index has a memory cost, but its contribution did not produce a meaningful increase in measured peak process RSS. Peak host RSS does not directly measure the index's standalone memory footprint or GPU device-memory consumption. ## Correctness and Regression Validation The implementation has been validated using focused correctness, allocation-plan equivalence, and complexity regression tests. | Validation | Result | |---|---| | Full-build test harness | 28 passed, 0 skipped | | Independent test rerun | 28 passed, 0 skipped | | Complete allocation/execution plan comparison | 15 plans identical to baseline | | Linear-algorithm semantic control | 20 passed | | Deterministic complexity regression | Passed | | CUDA CI planner cases | 12 passed, 0 skipped/failures | | Android minimal cross-compile, link, and size gate | Passed | The Android minimal-size gate passed at 1,607,070 bytes against the unchanged 1,607,680-byte limit. Android emulator runtime tests were skipped by the workflow and are not claimed as validated. ### Algorithmic Complexity Validation A deterministic complexity test compared the number of shape compatibility queries performed by the original and indexed planners. | Metric | Original planner | Indexed planner | |---|---:|---:| | Shape compatibility queries | 130,306 | 1,022 | This represents approximately **127× fewer shape compatibility queries** in the focused test. The result validates the intended reduction in unnecessary search work. It is not an end-to-end performance speedup measurement. The allocation-plan equivalence tests additionally verify that the optimized planner preserves the tested allocation and execution decisions. ## Scope and Limitations This PR focuses specifically on allocation-planning efficiency. - It reduces CPU-side planning overhead during session initialization. - It does not change operator implementations, numerical precision, or GPU kernels. - It does not directly optimize steady-state inference latency. - It preserves the existing buffer reuse eligibility rules and is intended to maintain equivalent allocation/execution plans. - Minimal builds retain the original linear implementation. - Hash-table lookups have expected O(1) complexity, but worst-case collisions and candidate-chain traversal can still require additional work. - The synthetic stress-test speedups are not representative of every model. - The random-weight 32-layer Llama decoder benchmark shows no substantial initialization improvement. ## Conclusion This PR eliminates a pathological source of quadratic free-buffer scanning in ONNX Runtime's allocation planner. The benchmark results demonstrate **283–297× faster allocation-plan construction** and **2.75–2.91× faster complete session initialization** on the targeted synthetic workload, without a material increase in measured peak host RSS. Control workloads and a 32-layer Llama decoder show little or no performance change, indicating that the optimization primarily benefits graphs with large numbers of incompatible reusable-buffer candidates. The change improves allocation-planning scalability while preserving the existing memory reuse behavior and keeping minimal builds on the original implementation. --------- Signed-off-by: rui-ren <15321482+rui-ren@users.noreply.github.com> Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
…icrosoft#32898) ### Description When deterministic compute is enabled, CUDA MatMulNBits now selects a stable fpA-intB tactic instead of using the timing-profiled choice. For small row counts it selects the available CUDA GEMV tactic; otherwise it selects a fixed non-split-K GEMM tactic. This path bypasses the profiled tactic cache. With deterministic compute disabled, the existing profiling behavior is unchanged. The PR adds tests for cache independence and exact 2-, 4-, and 8-bit weight reconstruction in deterministic mode. ### Motivation and Context Fresh processes running the same quantized model with greedy decoding sometimes selected different MatMulNBits tile and split-K tactics. Those tactics produced slightly different FP16 outputs from identical inputs, eventually changing the ninth generated token from `680` to `264`. ORT’s existing deterministic-compute flag did not previously govern this tactic selection. On SM80, the isolated MatMul output matched bitwise in 12/12 fresh processes after the change. In a downstream full-model validation, 36/36 deterministic runs produced the same token stream and tactic fingerprint; with the flag disabled, both original branches still appeared. Deterministic selection may be slower than a profiled tactic. Validation does not establish identical results across hardware or batch shapes. ### Qwen INT4 decode specialization This PR now also routes the Qwen Flash-Next language-model projection through the existing M=1 symmetric INT4 CUDA GEMV specialization for the exact shape: - `M = 1` - `N = 248320` - `K = 2560` - block size `32` - no zero points The dispatch remains shape- and format-gated, so other `MatMulNBits` workloads retain their existing selection behavior. The specialization reduces single-token draft-head projection overhead used by Qwen speculative decoding. --------- Co-authored-by: Akshay Sonawane <asonawane@microsoft.com> Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> Co-authored-by: copilot-swe-agent[bot] <198982749+Copilot@users.noreply.github.com>
ai-fw-intg
requested review from
Jaswanth51,
ankitm3k,
jatinwadhwa921 and
vthaniel
October 21, 2026 20:39
hdharpure9922
self-requested a review
October 22, 2026 04:43
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.