Sync with Microsoft ONNX Runtime - 29082026 - #1273
Open
ai-fw-intg wants to merge 16 commits into
Open
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>
ai-fw-intg
requested review from
Jaswanth51,
ankitm3k,
jatinwadhwa921 and
vthaniel
August 28, 2026 20:36
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.