Sync with Microsoft ONNX Runtime - 31082026 - #1275
Merged
Merged
Conversation
### Description - Replace dense `VarlenCausalConvWithState` checkpoint output with the compact three-output ABI `[output, final_state, state_update]`. - Add optional per-request `capture_count` input and bounded `state_update_capacity` attribute. - Return only the appended activation values required to replay accepted speculative prefixes, shaped `[batch_size, state_update_capacity, channels]`. - Remove the retired `max_checkpoints` attribute and `prefix_states` output. - Add schema validation, CUDA shape/runtime checks, and focused compact-capture tests. ### Motivation and Context Dense checkpoint output scales the full convolution state with speculative width. The convolution transition only needs the newly appended value at each accepted position, so compact capture avoids duplicating the full `[channels, kernel_size - 1]` state while preserving exact replay from committed state. This change is coordinated with the compact GatedDeltaNet operator and ONNX Runtime GenAI replay integration: - microsoft#32282 - microsoft/onnxruntime-genai#2472 ### Validation - `ContribOpVarlenCausalConvWithStateTest.*`: 39/39 passed on H200 with CUDA 13.0 and cuDNN 9.23. - Coverage includes sequential-prefix equivalence, all-ones decode, adjacent requests, omitted output at zero capacity, capture-count requirements and shape validation, and capacity bounds. - File-scoped lintrunner and `git diff --check` passed for all five changed files. - Generated `docs/ContribOperators.md` is intentionally excluded and left to CI generation.
The subgroup-matrix code was compiled out of WASM builds via `#if !defined(__wasm__)` guards, because emdawnwebgpu did not expose the Dawn subgroup-matrix API. Bump the Dawn dependency to v20260818.211311, which includes the needed support for the Dawn subgroup-matrix API in WASM builds.
### Description `PagedAttention` hard-codes a bottom-right causal mask on every backend. Block drafters submit their whole query block in a single step and each row of that block has to attend to the rest of the block, which a causal mask forbids. This adds an `is_causal` INT attribute, default `1`. Every existing graph is byte-for-byte unchanged. When `is_causal=0`: - the value is forwarded to FlashAttention's `mha_varlen_fwd`, giving a mask that is unbounded on the right; - `local_window_size` still bounds the mask on the left (`window_size_left = local_window_size - 1`), so sliding-window drafters keep working; - the paged-decode and CUTLASS `MemoryEfficientAttention` backends are excluded, because both bake the causal mask into the kernel. Rather than silently returning a causal result, the operator returns `INVALID_ARGUMENT` naming the requirement. ### Motivation and Context Needed to export the DFlash 2 block drafter for Qwen3.8-27B. The drafter's checkpoint sets `is_causal: false`: it predicts a block of `block_size` tokens at once, so its attention over the query block is bidirectional while the cached context stays strictly to the left. ### Testing `onnxruntime/test/python/transformers/test_paged_attention.py` gains four cases in `TestPagedAttentionFeatures`, comparing against the existing `attention_ref` with the right window opened: - `test_non_causal` — bidirectional query block, no local window - `test_non_causal_local_window` — left bound still honoured - `test_non_causal_with_rotary_and_packed` - `test_non_causal_rejected_without_flash_attention` — asserts the `INVALID_ARGUMENT` message All four pass on H200 (SM90, CUDA 13.0). The surrounding `TestPagedAttention` / `TestPagedAttentionFeatures` suites are unaffected (81 passed; the one failure, `test_fp8_cache_0_per_tensor`, is a pre-existing `torch` → `numpy` `Float8_e4m3fn` conversion issue in the harness, present before this change). --------- Co-authored-by: github-actions[bot] <41898282+github-actions[bot]@users.noreply.github.com>
## Summary - Preserve an existing equivalent custom allocator when EP registration does not request replacement. - Add provider-neutral coverage that verifies the custom allocator is still returned and performs the allocation. - Add WebGPU plugin allocator coverage and disable the per-device distinct allocator case with a TODO until WebGPU EP exposes distinct allocator memory info per `OrtEpDevice`. This change is limited to allocator preservation during EP registration. It does not change allocator ownership or behavior when an EP library is unregistered. Related to microsoft#32164. ## Solve issues 1. If a user registers a custom allocator before registering a plugin EP, plugin EP registration may inadvertently remove the existing custom allocator. 2. When registering the WebGPU plugin EP with multiple GPU adapters, creating the allocator for the second adapter removes the allocator created for the first adapter. This repeats for subsequent adapters and can leave the environment without any shared WebGPU allocator. ## Testing - `onnxruntime_autoep_test --gtest_filter=SharedAllocators.*:WebGpuPluginSharedAllocatorRegistrationTest.*:WebGpuPluginSharedAllocatorTest.*` (8 passed, 2 disabled) - Lintrunner on all three changed files --------- Co-authored-by: Copilot App <223556219+Copilot@users.noreply.github.com>
…2259) ### Description - Create the WGPUInstance used by WebGPU.importJsDevice() with WGPUInstanceFeatureName_TimedWaitAny. - Export the instance creation helper from the Wasm API. - Add a browser E2E test that passes a user-created GPUDevice, runs the existing model, validates output, and destroys the device. [Test page](https://xiaofeihan1.github.io/ort-webgpu-device-version-test/) ### Motivation and Context Emscripten associates an imported JavaScript GPUDevice with the provided WGPUInstance. WebGPU EP uses wgpuInstanceWaitAny() for synchronous GPU-to-CPU downloads. Without TimedWaitAny on that instance, inference with a user-provided device fails with: BufferManager::Download ... Failed to wait for the operation:3. Fixes microsoft#32257 ### Testing - New browser E2E test fails against unmodified main with BufferManager::Download ... Failed to wait for the operation:3. - The same browser E2E test passes after the fix (1 passed) and validates MatMul output. - node --check onnxruntime/wasm/post-webgpu.js - node --check js/web/test/e2e/browser-test-webgpu-custom-device.js - node --check js/web/test/e2e/run-data.js - git diff --check
### Description - Add the CUDA `GatedDeltaNet` contrib operator with recurrent decode and tensor-core chunked prefill paths. - Support ragged and rank-4 inputs, fused Qwen gate normalization, and native GDN arithmetic from raw `A_log`. - Return the compact three-output ABI `[output, final_state, state_update]`. The FP32 `state_update` capsule packs decay, shared-key, and delta transitions needed to replay accepted speculative prefixes without materializing dense recurrent checkpoints. - Add schema and shape validation, CUDA registration, focused tests, a microbenchmark, and an authored operator guide. ### Motivation and Context Dense recurrent checkpoints scale the full FP32 state with draft width. At the Qwen3.8 geometry, a four-slot window across 48 GDN layers consumes 576 MiB. Compact transition capture keeps one committed state and records only the information needed to reconstruct an accepted prefix. The schema intentionally exposes only native arithmetic and does not include the experimental `arithmetic_mode` attribute. Models exported with the retired experimental ABI must be re-exported. Companion ONNX Runtime GenAI integration: microsoft/onnxruntime-genai#2472 ### Validation - `./build/cu130/Debug/onnxruntime_provider_test --gtest_filter='*GatedDeltaNet*'` - 26/26 focused tests passed on H200 with CUDA 13.0 and cuDNN 9.23. ### Performance and Quality On H200, context 2048, generation 256, with five paired fresh-process repetitions per batch: - Native arithmetic won all 40 MTP and DFlash2 throughput pairs versus the retired compatibility experiment. Median native/compatibility ratios were 1.0519 for MTP and 1.0543 for DFlash2. - Paired quality differences were not statistically significant: MMLU-Pro 83.75% native vs. 83.38% compatibility (McNemar p=0.73); GPQA 81.82% vs. 78.28% (p=0.23). - Against the retired separate-factor representation, the packed capsule preserved exact tokens and replay work. It won 20/20 MTP pairs and was at parity for DFlash2 batch 16, with a small-batch benefit.
### Description <!-- Describe your changes. --> Fix Graph::ToGraphProtoInternal to clear the destination GraphProto before populating it, rather than clearing the graph's backing proto. Add regression coverage for Compile API output using both an output-model write callback and a custom initializer-location callback, including: - Models with no initializers. - Embedded initializers. - External initializers. - Reloading the emitted model and running inference. - Verifying inputs, outputs, nodes, and initializers are serialized exactly once. ### 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. --> When compilation produced no EPContext nodes, the Compile API emitted a plain optimized ONNX model. If an output write callback and custom initializer-location callback were both configured, serialization appended graph fields to an already-populated destination. This duplicated nodes, inputs, outputs, and value information. CompileModel returned success, but loading the emitted model failed with: Error: Duplicate definition-site for (X). Clearing the destination proto before repopulating it ensures the emitted model remains valid while preserving existing embedded and external initializer handling. --------- Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
…s off Win2025. (microsoft#32281) ### 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. --> --------- Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> Co-authored-by: copilot-swe-agent[bot] <198982749+Copilot@users.noreply.github.com> Co-authored-by: eserscor <247253654+eserscor@users.noreply.github.com>
### Description <!-- Describe your changes. --> Add native AArch64 build and SwiftShader test lanes, package the plugin EP for Python, NuGet, and Foundry Local, and share the Linux WebGPU Docker context across architectures. Add environment variable `ORT_WEBGPU_EP_ALLOW_SOFTWARE_ADAPTER` to specify creation of an `OrtEpDevice` for the WebGPU EP and the CPU device which allows the WebGPU EP to be selected via EP device if no GPU is available. This allows the packaging test build to run with the Vulkan SwiftShader software implementation on a machine with no GPU. ### 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. --> Add WebGPU plugin EP Linux AArch64 package variants. --------- Co-authored-by: Copilot <175728472+Copilot@users.noreply.github.com>
) Add a Windows-only build option to download and configure the Agility SDK for Dawn's D3D12 backend. PTAL, thanks! @jchen10
### Description Remove the internal documentation URL from the Guardian baseline metadata while retaining the expected empty properties object. ### Motivation The security team requested that internal URLs be removed from the public repository. ### Validation - Parsed .config/guardian/.gdnbaselines successfully as JSON - Ran git diff --check
### Description Prevents `MatMulIntegerToFloat` fusion from processing overlapping patterns and removes scheduled nodes by index. Adds regression coverage for chained overlapping candidates. ### Motivation and Context Overlapping candidates could schedule the same node for removal more than once, leaving the transformed graph invalid. Co-authored-by: Daniel Song <danielsong@microsoft.com>
### Description - Register `BFloat16` for the GatedDeltaNet schema and CUDA kernel. - Route BFloat16 execution through the recurrent engine while retaining FP16-only tensor-core paths. - Add BFloat16 numerical and planner coverage. - Guard `mma.sync.m16n8k16` generation to SM80+, allowing pre-Ampere builds to compile. - Update operator documentation. ### Motivation and Context GatedDeltaNet supported float and float16 CUDA inputs but not BFloat16. Its unguarded SM80 MMA instruction also caused `ptxas` failures when compiling for older targets such as `compute_61`. --------- 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: Tianlei Wu <tlwu@microsoft.com>
## Description Adds the first phase of checked workspace estimation for CUDA packed attention operators. The implementation introduces a graph-free recipe shared by `PackedAttention` and `PackedMultiHeadAttention`, migrates runtime workspace sizing and offsets to that recipe, and preserves the existing allocation topology and byte totals for Flash, TensorRT fused attention, memory-efficient attention, and unfused attention. The recipe explicitly models packed token count (`T`) versus padded capacity (`B*S`), planar versus interleaved QKV layouts, projection and attention workspace components, producer vectorization width, backend scratch regions, and checked CUDA ABI narrowing. Part of microsoft#29775. Design: microsoft#29775 (comment) ## Changes - Add pure checked workspace problem/recipe types without graph or CUDA runtime dependencies. - Use the recipes in CUDA PackedAttention and PackedMultiHeadAttention runtime allocation and workspace views. - Preserve legacy workspace byte parity and two-allocation PackedAttention behavior. - Widen CUTLASS attention-bias stride arithmetic to avoid intermediate `int` overflow. - Add independent formula, overflow, layout, containment, header-isolation, and route-observed tests. ## Validation - 43 packed workspace recipe/parity tests. - PackedAttention TRT, memory-efficient, and unfused routes with `T < B*S`. - PackedMultiHeadAttention Flash, TRT, memory-efficient, unfused, and invalid-head fallback routes. - CUDA provider compile/link, feature-disabled compile, and isolated plugin-boundary header compile. ## Deferred Graph adapters and L1/L2 estimation remain for the next phase. Single planned-root/preallocation integration remains a later phase. Device-side validation of `token_offset` and cumulative sequence values is not added here because it requires a separate CUDA graph/capture-safe contract. ## Documentation - [CUDA Attention workspace estimation roadmap](https://github.com/microsoft/onnxruntime/blob/titaiwangms/attention-workspace-recipes/docs/annotated_partitioning/attention_workspace_estimation.md) - Updates the CUDA workspace inventory to distinguish runtime-exact sizing from conditional AOT estimation. --------- Co-authored-by: copilot-swe-agent[bot] <198982749+Copilot@users.noreply.github.com> Copilot-Session: c04148cc-7ace-4cf4-b981-0ddc92334e78
…el (microsoft#32197) ### Description Level2+ fusion transformers assign a fused node the execution provider of the nodes it replaces without checking that a kernel exists for the fused op. For example, an fp16 Add and an fp16 Gelu both have CPU kernels, but fusing them produces a com.microsoft.BiasGelu node that the CPU EP only implements for float, so session initialization fails when its kernel is looked up. Detect these nodes (IsFp16NodeOnCpuWithoutKernel) and route them through the existing isolated-fp16-node fp32 fallback in InsertCastTransformer, regardless of whether they're otherwise "isolated" or produce a graph output, since running them in fp16 isn't an option to begin with. Track which nodes had their CPU assignment already recorded by the partitioner so the partition-assignment callback isn't fired twice. Also fixes two gaps in the isolated-node check: the no-fp16-input bailout wasn't skipped for these no-kernel nodes (unlike its output-side twin), and the kernel-less check only looked at input types, missing nodes whose fp16-ness is only on the output. ### Motivation and Context Proposed fix for this issue: microsoft#32186 the unit test tried using "Abs" pretending it has a FP16 kernel, but the new check detects a missing kernel. Replacing with "Round" which actually has FP16 generally preserves the original intent of the check. --------- Co-authored-by: Claude Sonnet 5 <noreply@anthropic.com> Co-authored-by: Copilot Autofix powered by AI <175728472+Copilot@users.noreply.github.com>
…lay (microsoft#32121) ## Description `CudaAsyncBuffer::CopyToGpu` writes a kernel's pointer array into a pinned host buffer, issues a `cudaMemcpyAsync` H2D, and immediately hands the buffer back to the allocator via `AddDeferredReleaseCPUPtr`. Under stream capture that copy is not a copy, it is a *node*, and the node records the **host address**: every replay re-reads whatever that block holds at the time. A pure decode-replay loop never notices, because nothing recycles the block. Run one differently-shaped inference in between — a prefill — and it takes the block, and the next replay copies the prefill's bytes into what `Concat` believes is its array of input pointers. ## Summary of Changes | File | Change | |------|--------| | `onnxruntime/core/providers/cuda/cuda_kernel.h` | `CudaAsyncBuffer` retains its pinned host buffer for the provider's lifetime when the stream is capturing; defer-release as before when it is not. | | `onnxruntime/core/providers/cuda/cuda_execution_provider.h` | Declare `RetainBufferForGraphCapture`, plus the mutex and buffer list backing it. | | `onnxruntime/core/providers/cuda/cuda_execution_provider.cc` | Implement it. | This covers every `CudaAsyncBuffer` user: `Concat`, `Split`, `ScatterND`, batched `MatMul`, `NonMaxSuppression`, and `CudnnRnnBase`. ## Testing `compute-sanitizer` on a 4-layer DeepSeek-V4 export, prefill interleaved with graph-replayed decode, before the fix: ``` Invalid __global__ read of size 4 bytes at _ConcatKernel<int>(..., const void **, int)+0x300 Access to 0xfffffffffffffc81 is misaligned Host Frame: cudaGraphLaunch ``` Clean after the fix. Non-capturing runs are unaffected — the deferred-release path is unchanged, so there is no new retention in the common case. ## Motivation and Context Buffers are retained only while capturing, which is a bounded, one-time cost per captured graph rather than per `Run`. One deliberate non-change: the *destination* scratch is still freed at the end of `Compute`, so a replay writes into memory the arena considers free. That is the address-stability assumption graph capture already makes throughout this EP, and replays never overlap another run, so it is left alone here rather than widened in this PR. ## Checklist - [x] Tests added/updated (verified with `compute-sanitizer` on a capture/replay workload) - [x] No breaking changes - [ ] Documentation updated (not applicable)
## Summary - Tune tensor-core NVFP4 GEMV dispatch to require multiple waves before selecting wide column tiles. - Preserve a larger K split for long reductions while avoiding excessive K splitting on Qwen gate/up shapes. - Add shape-scoped benchmark overrides, boundary tests, and CUDA contrib documentation. ## Motivation On H200, the previous dispatch could select a configuration with too few blocks for Qwen MTP shapes such as N=17408, K=5120. The updated policy keeps KSplit=8 for the longer K=8192 reduction while using KSplit=2 for K=5120, and requires sufficient grid waves before selecting wide column tiles. This PR is the tiling follow-up to microsoft#32128 and contains no duplicate vectorized NVFP4 dequantization changes. It can be rebased/stacked on microsoft#32128 after that PR merges. ## Validation - CUDA 13.0 / SM90 build passed. - `git diff --check` passed. - Added Qwen boundary coverage for the selected tensor-core tilings. - The local build configuration had provider unit tests disabled, so the new gtest was not executed locally.
### Description Extends the WebGPU Conv+activation fusion allowlist by eight kinds: QuickGelu, HardSwish, Elu, Gelu, Gelu(tanh), Softplus, ThresholdedRelu and Erf, each with a WGSL snippet for the generated-shader path and a matching branch in the im2col template. QuickGelu at alpha == 1 is a distinct shader because the multiply drops out entirely, so it carries its own `QuickGeluUnitAlpha` cache-key term to stop the two variants colliding in the pipeline cache. This also fixes the QuickGelu alpha fallback, which was `1.0f` rather than the schema default `1.702f`: attributes are materialized onto nodes during `Graph::Resolve()`, so the fallback was unreachable and could not change model output, but it was still wrong on paper and inconsistent with the standalone WebGPU QuickGelu kernel. Four of the eight (`Elu` and the three contrib GELU variants) also need a transpose-optimizer handler to fuse end to end, and since that map is shared cross-EP infrastructure rather than WebGPU code it lives in microsoft#32118; until that lands those four still execute correctly, just unfused. Covered by 35 new tests, including negative controls for Selu and the CPU EP, execution parity against unfused results, and one that strips QuickGelu's alpha after `Resolve()` so it actually fails without the fix. ### Motivation and Context These activations are common after convolutions. QuickGelu is how SiLU/Swish reaches the graph and HardSwish appears throughout MobileNet-class models, but each one previously forced a separate dispatch and a round trip through global memory. --------- Co-authored-by: Copilot App <223556219+Copilot@users.noreply.github.com> Co-authored-by: Ananya Anand <t-anaanand@microsoft.com> Co-authored-by: Ananya Anand <4n4ny4@users.noreply.github.com>
### Description - Extend the CUDA small-N GEMV path from short decode steps to speculative verification batches with up to 64 rows. - Keep eligible FP16/BF16 MatMul shapes on the fused GEMV path instead of falling back to the general GEMM path. - Extend block-scaled FP8 and NVFP4 weight-only GEMV launchers across row-tile boundaries by splitting larger speculative steps into supported sub-launches. - Add coverage for row counts across the dispatch boundaries, including 9, 17, 33, and 64 rows. ### Motivation and Context Multi-token speculative verification produces matrix row counts beyond the original decode-only range. Falling back at those boundaries adds dequantization, workspace, and general GEMM overhead to a latency-sensitive path. This change keeps those shapes on the specialized kernels while preserving the existing fallback for ineligible dimensions. ### Validation - CUDA EP internal tests: 78 passed, 6 skipped across 20 suites. `MatMulSmallNGemvOpTest.DispatchesEligibleShapesWhenEnabled` passed and covers 8, 9, 33, and 64 rows. - `MatMulBlockQuantizedFp4WeightOpTest.GemvTensorCoreTilesFp16` - `MatMulBlockQuantizedFp4WeightOpTest.GemvTensorCoreTilesBf16` - `MatMulBlockQuantizedFp8WeightOpTest.GemvTensorCoreTilesFp16` - The three focused block-scaled tests passed in 19.0 seconds. - File-scoped lintrunner and `git diff --check` passed for all 12 changed files.
ai-fw-intg
requested review from
Jaswanth51,
ankitm3k,
jatinwadhwa921 and
vthaniel
August 30, 2026 20:36
hdharpure9922
self-requested a review
August 31, 2026 04:30
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.