Sync with Microsoft ONNX Runtime - 20082026 - #1264
Merged
Merged
Conversation
This shall not be merged until 0.1 is released.
### Description Adds CPU and WebGPU contrib kernels for Qwen-3.5 operators, including GatedAdd, GatedRMSNorm, and LinearAttentionGate. Also fixes WebGPU input/output binding order and removes an unused CPU include. ### Motivation and Context This adds CPU and WebGPU support for Qwen-3.5 ops. The PR is opened from the renamed branch `copilot/add-cpu-webgpu-kernels-for-qwen-3-5` to accurately reflect the implementation scope. --------- Co-authored-by: copilot-swe-agent[bot] <198982749+Copilot@users.noreply.github.com> Co-authored-by: kunal-vaishnavi <115581922+kunal-vaishnavi@users.noreply.github.com>
## Description Enable FP4 QMoE for Blackwell SM120 in Windows CUDA builds. Previously, MSVC builds rejected this configuration because compiling the full LLM CUDA source set for native SM120 exposed `tcgen05` compilation failures. ## Changes - Split the FP4 QMoE sources that require native SM120 instructions into dedicated CUDA object targets: - `onnxruntime_providers_cuda_llm_fp4` - `onnxruntime_providers_cuda_plugin_llm_fp4` - Compile those FP4 targets for native `sm_120a` (`120-real`) so `expandInputRowsKernel` can use `cvt.e2m1x2`. - Compile the general LLM CUDA targets as `compute_120` PTX (`120-virtual`) on MSVC, avoiding `tcgen05` failures in unrelated translation units. - Enable the CUTLASS SM120 grouped TMA kernels needed by FP4 QMoE on Windows. - For CUDA 13 and MSVC, apply `/Zc:__cplusplus-` only to the host compilation of the SM120 TMA object target. This avoids MSVC `C2719` errors in generated host stubs for by-value, 128-byte-aligned `CUtensorMap` parameters while preserving C++20 for CUDA device compilation. - Apply the same source split and workaround to both the in-tree CUDA provider and the CUDA plugin provider. - Update the FP4 QMoE documentation with the Windows build behavior and constraints. ## Validation - Built the native FP4 QMoE object target for `sm_120a` with MSVC and CUDA 13. - Built the native SM120 CUTLASS TMA object target with the target-local host compiler workaround. - The CUDA plugin packaging pipeline verified this change: the Windows x64 build passed. --------- Co-authored-by: copilot-swe-agent[bot] <198982749+Copilot@users.noreply.github.com>
ai-fw-intg
requested review from
Jaswanth51,
ankitm3k,
jatinwadhwa921 and
vthaniel
August 19, 2026 20:36
hdharpure9922
self-requested a review
August 20, 2026 04:48
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.