diff --git a/.github/workflows/windows-web-ci-workflow.yml b/.github/workflows/windows-web-ci-workflow.yml
index 79ca29fe21480..24be32171cc6f 100644
--- a/.github/workflows/windows-web-ci-workflow.yml
+++ b/.github/workflows/windows-web-ci-workflow.yml
@@ -158,7 +158,6 @@ jobs:
# WebGPU EP tests
- name: Run ort-web tests - WebGPU EP
- continue-on-error: true
shell: cmd
run: |
mkdir ${{ runner.temp }}\web\test\07
diff --git a/.github/workflows/windows_webgpu.yml b/.github/workflows/windows_webgpu.yml
index af528a912b4b6..77a99473a47d9 100644
--- a/.github/workflows/windows_webgpu.yml
+++ b/.github/workflows/windows_webgpu.yml
@@ -161,7 +161,6 @@ jobs:
run: python tools\ValidateNativeDelegateAttributes.py
shell: cmd
working-directory: ${{ github.workspace }}\csharp
- continue-on-error: true
webgpu_test_x64_RelWithDebInfo:
name: webgpu_test_x64_RelWithDebInfo (${{ matrix.vcpkg_option }})
diff --git a/cmake/external/cuda_configuration.cmake b/cmake/external/cuda_configuration.cmake
index 47d6eef45302c..a7cf51ebbbd0a 100644
--- a/cmake/external/cuda_configuration.cmake
+++ b/cmake/external/cuda_configuration.cmake
@@ -121,16 +121,19 @@ macro(setup_cuda_architectures)
message(WARNING "Detecting native CUDA compute capability - fail")
message(WARNING "CUDA compute capability detection failed, compiling for all optimized architectures")
unset(CMAKE_CUDA_ARCHITECTURES)
+ unset(CMAKE_CUDA_ARCHITECTURES CACHE)
else()
message(STATUS "Detecting native CUDA compute capability - done")
set(CMAKE_CUDA_ARCHITECTURES "${CUDA_ARCH_OUTPUT}")
endif()
elseif(CMAKE_CUDA_ARCHITECTURES STREQUAL "all")
unset(CMAKE_CUDA_ARCHITECTURES)
+ unset(CMAKE_CUDA_ARCHITECTURES CACHE)
message(STATUS "Setting CMAKE_CUDA_ARCHITECTURES to all enables a list of architectures OnnxRuntime optimized for, "
"not all architectures CUDA compiler supports.")
elseif(CMAKE_CUDA_ARCHITECTURES STREQUAL "all-major")
unset(CMAKE_CUDA_ARCHITECTURES)
+ unset(CMAKE_CUDA_ARCHITECTURES CACHE)
message(
STATUS "Setting CMAKE_CUDA_ARCHITECTURES to all-major enables a list of architectures OnnxRuntime optimized for, "
"not all major architectures CUDA compiler supports.")
diff --git a/cmake/onnxruntime_providers_cuda_plugin.cmake b/cmake/onnxruntime_providers_cuda_plugin.cmake
index 87d3a3af13975..8450471b65cd9 100644
--- a/cmake/onnxruntime_providers_cuda_plugin.cmake
+++ b/cmake/onnxruntime_providers_cuda_plugin.cmake
@@ -486,6 +486,12 @@ endif()
message(STATUS "CUDA Plugin EP: cuDNN include: ${CUDA_PLUGIN_CUDNN_INCLUDE_DIR}")
message(STATUS "CUDA Plugin EP: cuDNN runtime library: ${CUDA_PLUGIN_CUDNN_LIBRARY}")
+# With cuDNN in /usr drop CUDA_PLUGIN_CUDNN_INCLUDE_DIR; we do not want
+# "-isystem /usr/include" ahead of libstdc++, breaking #include_next mechanism.
+if(CUDA_PLUGIN_CUDNN_INCLUDE_DIR IN_LIST CMAKE_CXX_IMPLICIT_INCLUDE_DIRECTORIES)
+ set(CUDA_PLUGIN_CUDNN_INCLUDE_DIR "")
+endif()
+
# Include directories — only public ORT headers + CUDA toolkit + cuDNN + internal headers for adapter
target_include_directories(onnxruntime_providers_cuda_plugin PRIVATE
${REPO_ROOT}/include
diff --git a/docs/ContribOperators.md b/docs/ContribOperators.md
index d45736fe9e6f1..2589e5878a808 100644
--- a/docs/ContribOperators.md
+++ b/docs/ContribOperators.md
@@ -133,6 +133,7 @@ Do not modify directly.*
* com.microsoft.SparsePagedAttention
* com.microsoft.SparseToDenseMatMul
* com.microsoft.SpeculativePrefixCommit
+ * com.microsoft.SwiGLU
* com.microsoft.Tokenizer
* com.microsoft.TorchEmbedding
* com.microsoft.TransposeMatMul
@@ -8360,6 +8361,62 @@ This version of the operator has been available since version 1 of the 'com.micr
+### **com.microsoft.SwiGLU**
+
+ Clamped SwiGLU, the gated activation of a feed-forward or MoE expert block:
+ G = clamp(gate, max=limit)
+ L = clamp(up, min=-limit, max=limit)
+ output = G * Sigmoid(activation_alpha * G) * (L + activation_beta)
+ with the arithmetic done in float regardless of T, widened to double for the product when `activation_alpha * G` is below -87 so that a negative tail is not lost to float exp underflow. A `limit` of zero or less disables both clamps.
+
+ The result is rounded to T once, after the multiply by (L + activation_beta). With `round_activation` set, G * Sigmoid(activation_alpha * G) is first rounded to T, as a QuickGelu node followed by a Mul in T would, which keeps a fused QuickGelu + Mul bit-identical to the graph it replaces.
+
+ This uses the same activation_alpha / activation_beta formula and positive-limit clamping rule as MoE and QMoE. The no-clamp sentinel differs: here any `limit` of zero or less disables clamping, while MoE/QMoE use infinity when `swiglu_limit` is omitted and clamp finite values, including an explicit zero.
+ With `up` supplied and the default `limit` and `activation_beta`, it equals the ONNX SwiGLU operator (opset 28) with `alpha` = `activation_alpha`.
+
+ `up` is optional. When it is omitted, `gate` carries both halves of one `[.., 2 * inter]` projection -- gate first, then up -- and the split is done internally, so a fused sibling GEMM needs no `Split` node.
+
+#### Version
+
+This version of the operator has been available since version 1 of the 'com.microsoft' operator set.
+
+#### Attributes
+
+
+- activation_alpha : float
+- Multiplier inside the sigmoid.
+- activation_beta : float
+- Value added to the linear half. A zero value leaves it unchanged, including a negative zero.
+- limit : float
+- Clamp applied to both halves; zero or less disables it.
+- round_activation : int
+- If non-zero, round G * Sigmoid(activation_alpha * G) to T before multiplying by (L + activation_beta). Zero rounds only the final result.
+
+
+#### Inputs (1 - 2)
+
+
+- gate : T
+- Gate half of the projection, or both halves concatenated.
+- up (optional) : T
+- Up half of the projection, same shape as gate.
+
+
+#### Outputs
+
+
+- output : T
+- Activated tensor, shaped like one half of the projection.
+
+
+#### Type Constraints
+
+
+- T : tensor(float16), tensor(bfloat16), tensor(float)
+- Constrain the activation type to float tensors.
+
+
+
### **com.microsoft.Tokenizer**
Tokenizer divides each string in X into a vector of strings along the last axis. Allowed input shapes are [C] and [N, C].
diff --git a/docs/OperatorKernels.md b/docs/OperatorKernels.md
index 21f0365fc90a2..cc0442284693a 100644
--- a/docs/OperatorKernels.md
+++ b/docs/OperatorKernels.md
@@ -703,6 +703,9 @@ The **OpSet Version** column uses the following notation:
|||14|**T** = tensor(double), tensor(float), tensor(float16)
**U** = tensor(double), tensor(float), tensor(float16)|
|||[9, 13]|**T** = tensor(double), tensor(float), tensor(float16)|
|||[7, 8]|**T** = tensor(double), tensor(float), tensor(float16)|
+|BitShift|*in* X:**T**
*in* Y:**T**
*out* Z:**T**|28+|**T** = tensor(int16), tensor(int32), tensor(int64), tensor(int8), tensor(uint16), tensor(uint32), tensor(uint64), tensor(uint8)|
+|||[11, 27]|**T** = tensor(uint16), tensor(uint32), tensor(uint64), tensor(uint8)|
+|BitwiseXor|*in* A:**T**
*in* B:**T**
*out* C:**T**|18+|**T** = tensor(int16), tensor(int32), tensor(int64), tensor(int8), tensor(uint16), tensor(uint32), tensor(uint64), tensor(uint8)|
|Cast|*in* input:**T1**
*out* output:**T2**|28+|**T1** = tensor(bfloat16), tensor(bool), tensor(double), tensor(float), tensor(float16), tensor(float4e2m1), tensor(float8e4m3fn), tensor(float8e5m2), tensor(int16), tensor(int32), tensor(int64), tensor(int8), tensor(uint16), tensor(uint32), tensor(uint64), tensor(uint8)
**T2** = tensor(bfloat16), tensor(bool), tensor(double), tensor(float), tensor(float16), tensor(float4e2m1), tensor(float8e4m3fn), tensor(float8e5m2), tensor(int16), tensor(int32), tensor(int64), tensor(int8), tensor(uint16), tensor(uint32), tensor(uint64), tensor(uint8)|
|||[25, 27]|**T1** = tensor(bfloat16), tensor(bool), tensor(double), tensor(float), tensor(float16), tensor(float4e2m1), tensor(float8e4m3fn), tensor(float8e5m2), tensor(int16), tensor(int32), tensor(int64), tensor(int8), tensor(uint16), tensor(uint32), tensor(uint64), tensor(uint8)
**T2** = tensor(bfloat16), tensor(bool), tensor(double), tensor(float), tensor(float16), tensor(float4e2m1), tensor(float8e4m3fn), tensor(float8e5m2), tensor(int16), tensor(int32), tensor(int64), tensor(int8), tensor(uint16), tensor(uint32), tensor(uint64), tensor(uint8)|
|||[23, 24]|**T1** = tensor(bfloat16), tensor(bool), tensor(double), tensor(float), tensor(float16), tensor(float4e2m1), tensor(float8e4m3fn), tensor(float8e5m2), tensor(int16), tensor(int32), tensor(int64), tensor(int8), tensor(uint16), tensor(uint32), tensor(uint64), tensor(uint8)
**T2** = tensor(bfloat16), tensor(bool), tensor(double), tensor(float), tensor(float16), tensor(float4e2m1), tensor(float8e4m3fn), tensor(float8e5m2), tensor(int16), tensor(int32), tensor(int64), tensor(int8), tensor(uint16), tensor(uint32), tensor(uint64), tensor(uint8)|
@@ -1179,6 +1182,7 @@ The **OpSet Version** column uses the following notation:
|SparseAttentionIndexer|*in* query:**T**
*in* key:**T**
*in* query_norm_weight:**T**
*in* key_norm_weight:**T**
*in* cos_cache:**T**
*in* sin_cache:**T**
*in* mask:**TB**
*in* past_key:**T**
*in* gate:**T**
*in* position_bias:**T**
*in* head_weights:**T**
*in* position_ids:**I**
*in* past_sequence_length:**M**
*in* past_proj_buffer:**T**
*out* selected_indices:**M**
*out* present_key:**T**
*out* present_proj_buffer:**T**|1+|**I** = tensor(int64)
**M** = tensor(int32)
**T** = tensor(bfloat16), tensor(float), tensor(float16)
**TB** = tensor(int64)|
|SparsePagedAttention|*in* query:**T**
*in* key:**T**
*in* value:**T**
*in* key_cache:**T_CACHE**
*in* value_cache:**T_CACHE**
*in* cumulative_sequence_length:**S**
*in* past_seqlens:**S**
*in* block_table:**S**
*in* slot_mapping:**S**
*in* selected_indices:**S**
*in* selected_counts:**S**
*in* auxiliary_key:**T**
*in* auxiliary_value:**T**
*in* auxiliary_lengths:**S**
*in* cos_cache:**T**
*in* sin_cache:**T**
*in* head_sink:**T**
*in* q_norm_weight:**T**
*in* k_norm_weight:**T**
*in* k_scale:**T_KV_SCALE**
*in* v_scale:**T_KV_SCALE**
*in* attention_metadata:**S**
*out* output:**T**
*out* key_cache_out:**T_CACHE**
*out* value_cache_out:**T_CACHE**|1+|**S** = tensor(int32)
**T** = tensor(bfloat16), tensor(float16)
**T_CACHE** = tensor(bfloat16), tensor(float16), tensor(int8)
**T_KV_SCALE** = tensor(float)|
|SpeculativePrefixCommit|*in* next_tokens:**I64**
*in* input_ids:**I64**
*in* status:**I32**
*in* ctl:**I64**
*in* ctx_commit:**I64**
*in* prefixes:**I64**
*in* ptrs:**I64**
*out* head_c:**I32**
*out* updated_ctl:**I64**
*out* updated_ctx_commit:**I64**|1+|**I32** = tensor(int32)
**I64** = tensor(int64)|
+|SwiGLU|*in* gate:**T**
*in* up:**T**
*out* output:**T**|1+|**T** = tensor(bfloat16), tensor(float), tensor(float16)|
|TransposeMatMul|*in* A:**T**
*in* B:**T**
*out* Y:**T**|1+|**T** = tensor(bfloat16), tensor(double), tensor(float), tensor(float16)|
|Trilu|*in* X:**T**
*in* k:**tensor(int64)**
*out* Y:**T**|1+|**T** = tensor(bfloat16), tensor(bool), tensor(double), tensor(float), tensor(float16), tensor(int16), tensor(int32), tensor(int64), tensor(int8), tensor(uint16), tensor(uint32), tensor(uint64), tensor(uint8)|
|UnfoldTensor|*in* input:**T**
*out* output:**T**|1+|**T** = tensor(bfloat16), tensor(bool), tensor(double), tensor(float), tensor(float16), tensor(int16), tensor(int32), tensor(int64), tensor(int8), tensor(string), tensor(uint16), tensor(uint32), tensor(uint64), tensor(uint8)|
diff --git a/docs/Pattern_Rewriter.md b/docs/Pattern_Rewriter.md
index 333cfe6eadfd7..464bee2871e0c 100644
--- a/docs/Pattern_Rewriter.md
+++ b/docs/Pattern_Rewriter.md
@@ -1245,6 +1245,75 @@ before entering the transformer or after an error. Callers must not concurrently
graph. Options and compiled rules are immutable; match state and statistics are
per call.
+### Restricting rewriting to a selected region
+
+The scoped overload accepts an explicit set of live node IDs:
+
+```cpp
+InlinedHashSet region(selected_nodes.begin(), selected_nodes.end());
+ORT_RETURN_IF_ERROR(rewriter.Run(graph, statistics, region));
+ORT_RETURN_IF_ERROR(graph.Resolve());
+```
+
+Only roots in the region are visited. Every node consumed by a match must also
+belong to the region; a match crossing its boundary is rejected before either
+rule callback runs. Inputs may come from outside the region, and returned values
+may feed outside consumers. Their `NodeArg` identities are preserved, although
+boundary edges are rewired to new producers. Existing deletion-boundary and EP
+checks still apply. Region-scope and deletion-boundary rejections both report
+`FailureStage::Boundary`; the diagnostic reason distinguishes them.
+
+Each successful rewrite removes its source IDs from `region` and adds all of its
+replacement IDs. New nodes participate in the same ordered, iterative matching;
+unmatched region nodes remain in the returned set. This lets a capability
+callback reconcile its node list without rescanning unrelated nodes. The
+overload neither selects an EP nor performs kernel lookup, placement, or resource
+accounting. Callers must still validate those contracts for the resulting region.
+
+An empty region does no work. Invalid input IDs return an error before matching
+or mutation. Match-only mode leaves the set unchanged. Failures do not roll back
+earlier rewrites: the set reflects completed replacements, not a transaction.
+An ordinary staging failure leaves that match's source IDs in the set and does
+not add its staged replacement IDs.
+As with unscoped `Run`, resolve the graph explicitly before later consumers.
+
+### Function expansion and capability application
+
+The rewriter does not decide when ORT expands functions or assigns nodes to an
+EP. A pattern can match a function-call node, but this does not inspect or
+rewrite its model-local function definition.
+
+Capability discovery and rewrite application are separate operations. In
+ordinary ONNX initialization, AOT function inlining can query `GetCapability`
+before normal partitioning. Reporting a capability can preserve a function
+call for later partitioning; it is not permission to run its optimization
+callback during that query. Discovery must be read-only and repeatable.
+Normal partitioning checks availability before invoking the selected
+capability's optimization callback and placing its updated node set.
+
+`PatternCapabilityTests.PreservesFunctionUntilAcceptedRewrite` and
+`PatternCapabilityTests.ExpandsUnclaimedFunctionAndRediscoversBody` exercise
+these paths with a test-only CPU EP and existing CPU kernels:
+
+| Function-call candidate | Expected integration behavior |
+|---|---|
+| Recognized with an executable rewrite | Discovery leaves the call intact. After acceptance, scoped rewriting replaces the call and updates the parent capability before placement. |
+| Not recognized | Existing function expansion exposes the body. A later capability query discovers its Transpose chain and returns a fresh region for rewriting. |
+
+Both tests run with AOT inlining enabled and disabled. With AOT disabled, the
+unclaimed call expands after a partitioning pass, and partitioning repeats.
+No borrowed matches or capabilities from the function-call graph are retained
+for matching the expanded body. Assertions cover the original call's removal,
+callback and rewrite counts, final node types and CPU assignment, output names,
+and numerical equality with ordinary CPU execution.
+
+This is a fixed-function integration probe, not general function-equivalence
+checking or a production registration interface. Recognizing only a function
+name is insufficient for arbitrary model-local definitions. The probe does not
+change expansion policy or cross-EP priority, compare whole-function and
+expanded-body placement alternatives, or establish ORT-format export,
+resource-accounting, or general N-to-M placement contracts.
+
### Resolved metadata and mutation boundaries
Call `Graph::Resolve` after constructing a source graph and before `Run` or
@@ -1290,9 +1359,10 @@ remain unsupported; identity elimination is a separate transformation.
`Options::max_iterations` is a positive upper bound, defaulting to 10.
Each iteration processes a fixed candidate batch, initially all live nodes in
-topological order. Replacements and affected existing nodes are queued for the
-next batch, never appended to the current batch. An existing later candidate
-can still match an earlier replacement in the same batch. Processing stops
+topological order. For scoped `Run`, the initial batch contains only region nodes.
+Replacements and affected existing nodes are queued for the next batch, never
+appended to the current batch. An existing later candidate can still match an
+earlier replacement in the same batch. Processing stops
when a batch performs no rewrites, there are no pending candidates, or the bound
is reached. Match-only mode processes exactly one batch without mutation.
diff --git a/docs/design/Extensible_Runtime_Fusion.md b/docs/design/Extensible_Runtime_Fusion.md
new file mode 100644
index 0000000000000..1089b5a596e56
--- /dev/null
+++ b/docs/design/Extensible_Runtime_Fusion.md
@@ -0,0 +1,987 @@
+# EP-registered runtime fusion for canonical ONNX
+
+Status: **engineering design for review, not approved for implementation** (revised October 28, 2026).
+Tracks [#33403](https://github.com/microsoft/onnxruntime/issues/33403).
+No native build, dependency installation, GPU allocation, model execution, or runtime probe was
+performed for this design. Test names, commands, numerical gates, and performance budgets below are
+proposed acceptance criteria, not results. Maintainer review and acceptance are still required.
+
+## 1. Relationship to the original proposal
+
+This document refines the
+[pinned Extensible Runtime Fusion proposal](https://github.com/microsoft/onnxruntime/blob/587ef1bf3467935b083740a52ac8fcba3785b2a6/docs/design/Extensible_Runtime_Fusion.md).
+That file is absent from the default-branch base used here; the proposal branch is a reference,
+not an implementation dependency or a branch to merge.
+
+The north star remains:
+
+> Mobius exports canonical standard ONNX. Olive owns quantization policy. Compiler EPs retain
+> visibility and compile their regions. Implementation packages own optional kernel-fusion patterns
+> and adapters. ORT owns generic registration, validation, ownership enforcement, and rewriting.
+
+The canonical artifact is authoritative; a saved optimized model is a target-specific derivative.
+Existing contrib artifacts and private transformer registration remain supported. Kernel-internal
+improvements behind an unchanged operator contract already need no re-export. The additional proof
+required here is a **matcher/adapter package update with identical canonical artifact hashes**.
+
+### Scope
+
+Start with existing CUDA kernels: FP16 dense dynamic-cache GQA and stateful convolution, then a
+separately gated whole-`Scan` recurrence milestone. No new EP, kernel, quantization format,
+Flash-Next/MTP support, universal optimizer migration, automatic offload/defusion search, or
+transparent retry after state mutation. CPU and native WebGPU are architectural extensibility
+requirements, not first-release implementations. Compiler EPs need not adopt this mechanism.
+Qwen3.8-27B is a pilot, not a model-specific runtime API.
+
+### Revised integration direction
+
+Fusion patterns register on the EP side. Each EP owns their priority order and exposes eligible
+execution candidates through the capability path. ORT chooses compatible candidates in that order,
+subject to ownership, semantic, kernel, and, when enabled, memory-feasibility checks. There is no
+per-session speed competition and no separate post-partition package-selection pipeline.
+
+Separate candidate description from application so memory estimation can be connected without
+rerunning a mutating matcher. Reuse `GetCapability`, `ComputeCapability`, and the capability
+optimization application model first; define only the missing adapters. Independent package delivery
+remains a goal, not a reason to freeze a second graph-inspection ABI or require a competing EP.
+These are the proposed design direction, not a claim of maintainer approval or existing support.
+
+## 2. Evidence ledger and baseline discipline
+
+All runtime claims below refer to **R**, unless a difference is explicitly noted. Links are immutable
+research references, not claims that proposed interfaces exist.
+
+| Key | Exact source |
+|---|---|
+| R | [ORT main `1988019b6e23ac4def4628cddd603c1bfb504fe8`](https://github.com/microsoft/onnxruntime/tree/1988019b6e23ac4def4628cddd603c1bfb504fe8) |
+| M | [Mobius main `b1e9c8e9dec1dfa206183f957e77d07cfb734df2`](https://github.com/onnxruntime/mobius/tree/b1e9c8e9dec1dfa206183f957e77d07cfb734df2) |
+| G | [GenAI main `991e3560e63b8a26806d4ec7d702c04e75fb6e51`](https://github.com/microsoft/onnxruntime-genai/tree/991e3560e63b8a26806d4ec7d702c04e75fb6e51) |
+| F | [Fusion proposal `587ef1bf3467935b083740a52ac8fcba3785b2a6`](https://github.com/microsoft/onnxruntime/blob/587ef1bf3467935b083740a52ac8fcba3785b2a6/docs/design/Extensible_Runtime_Fusion.md) |
+| P | [Placement proposal `2f96ad081fdbe1d76a2a731262a998b3cd4c9c0d`](https://github.com/microsoft/onnxruntime/blob/2f96ad081fdbe1d76a2a731262a998b3cd4c9c0d/docs/annotated_partitioning/PartitioningWithAnnotationsAndMemoryConstraints.md), refreshed October 26 for the proposed recipe/planning integration only |
+| W | [Composable rewrite rules, gramalingam/onnxruntime#16 at `6d707086e04fe287f162b9dc2baf58b935d8e2bc`](https://github.com/gramalingam/onnxruntime/blob/6d707086e04fe287f162b9dc2baf58b935d8e2bc/docs/Pattern_Rewriter.md); optional matcher implementation, not an ORT-main dependency |
+| V | [OpenVINO plugin `7463f66e48d09fe5150af8c76520fbf75e08c490`](https://github.com/onnxruntime/onnxruntime-ep-openvino/tree/7463f66e48d09fe5150af8c76520fbf75e08c490); compiler-EP comparison only |
+
+The PR base is default-branch commit `327d11ab6c22e387acf793dd9de8caf406055ca9`.
+Source comparison with R found no changes in the cited session hook, public API, registry selection,
+Scan registration, or GQA/LinearAttention/GatedDeltaNet kernel files. Relevant differences are the
+accountant's new `INT32_MAX` validation for the CUDA GQA total-sequence bound and updates in the GQA
+workspace bounds/estimate helpers. These do not supply target re-estimation for arbitrary rewrites.
+Do not combine P's proposed reconciliation with R and describe the result as implemented.
+Any implementation PR must repeat the comparison against its actual base and publish its build pin.
+The October 26 revision also inspected the PR head's `execution_provider.h:152-179`,
+`compute_capability.h:27-44`, and `graph_partitioner.cc:1145-1181` for the capability integration below.
+Refreshing P and W does not refresh R or turn either proposal into shipped runtime functionality.
+V's `CreateEpFactories`/`ReleaseEpFactory` lifetime patterns are useful references. Its
+GetCapability/Compile and AddNodesToFuse path owns compiled regions; it is not a public extension
+that borrows existing CUDA kernels. Creating a competing EP solely to reach those kernels is rejected.
+
+### Current behavior versus missing mechanism
+
+| Evidence at R | Implication |
+|---|---|
+| [`inference_session.cc:1590-1602,1896-1952,2015-2087`](https://github.com/microsoft/onnxruntime/blob/1988019b6e23ac4def4628cddd603c1bfb504fe8/onnxruntime/core/session/inference_session.cc#L1896-L1952) | L1 precedes partition; mutation accounting is installed after partition; L2+ and casts repeat; copies follow. |
+| [`graph_transformer_utils.cc:413-426,464-467`](https://github.com/microsoft/onnxruntime/blob/1988019b6e23ac4def4628cddd603c1bfb504fe8/onnxruntime/core/optimizer/graph_transformer_utils.cc#L413-L426), [`graph_transformer_mgr.cc:38-59`](https://github.com/microsoft/onnxruntime/blob/1988019b6e23ac4def4628cddd603c1bfb504fe8/onnxruntime/core/optimizer/graph_transformer_mgr.cc#L38-L59) | Repeated optimization loops must not independently reselect a capability's implementation. |
+| [`execution_provider.h:152-179`](https://github.com/microsoft/onnxruntime/blob/3207aa12931ac9de054fdf0de183f7f036c01269/include/onnxruntime/core/framework/execution_provider.h#L152-L179) | `GetCapability` can return overlapping source-node sets; ORT decides assignment. The signature includes optimizer-registry and resource-accountant inputs. |
+| [`compute_capability.h:27-44`](https://github.com/microsoft/onnxruntime/blob/3207aa12931ac9de054fdf0de183f7f036c01269/onnxruntime/core/framework/compute_capability.h#L27-L44), [`graph_partitioner.cc:1145-1181`](https://github.com/microsoft/onnxruntime/blob/3207aa12931ac9de054fdf0de183f7f036c01269/onnxruntime/core/framework/graph_partitioner.cc#L1145-L1181) | Child optimization callbacks can update an available parent capability before placement/fusion. This is an application hook, not a complete alternative-recipe planner. |
+| [`group_query_attention_fusion.cc:413-440`](https://github.com/microsoft/onnxruntime/blob/1988019b6e23ac4def4628cddd603c1bfb504fe8/onnxruntime/core/optimizer/group_query_attention_fusion.cc#L413-L440) | Existing GQA fusion starts from contrib GQA, not canonical Attention. |
+| [`attention_fusion.cc:1030-1046`](https://github.com/microsoft/onnxruntime/blob/1988019b6e23ac4def4628cddd603c1bfb504fe8/onnxruntime/core/optimizer/attention_fusion.cc#L1030-L1046), [`session_state.cc:242-310`](https://github.com/microsoft/onnxruntime/blob/1988019b6e23ac4def4628cddd603c1bfb504fe8/onnxruntime/core/framework/session_state.cc#L242-L310) | Rewrites can retain EP assignment; finalization later binds registered kernels. |
+| [`inference_session.cc:1113-1129`](https://github.com/microsoft/onnxruntime/blob/1988019b6e23ac4def4628cddd603c1bfb504fe8/onnxruntime/core/session/inference_session.cc#L1113-L1129) | Private C++ registration is not an external binary contract. Preserve its existing enable behavior. |
+| [`onnxruntime_c_api.h:6130-6653,8210-8244`](https://github.com/microsoft/onnxruntime/blob/1988019b6e23ac4def4628cddd603c1bfb504fe8/include/onnxruntime/core/session/onnxruntime_c_api.h#L6130-L6653) | Reuse public `OrtGraph`/`OrtNode`/`OrtValueInfo` inspection where possible; Model Editor alone does not supply the capability recipe/application contract. |
+| [`onnxruntime_experimental_c_api.inc:21-37`](https://github.com/microsoft/onnxruntime/blob/1988019b6e23ac4def4628cddd603c1bfb504fe8/include/onnxruntime/core/session/onnxruntime_experimental_c_api.inc#L21-L37) | Experimental discovery exists, but no independently deployed EP-pattern registration interface is established here. |
+| [`kernel_registry_manager.cc:60-104`](https://github.com/microsoft/onnxruntime/blob/1988019b6e23ac4def4628cddd603c1bfb504fe8/onnxruntime/core/framework/kernel_registry_manager.cc#L60-L104) | Custom registries have precedence; an operator name alone proves neither implementation nor eligibility. |
+
+## 3. Integration and graph lifecycle (D1)
+
+### EP-owned registration and capability discovery
+
+Each EP owns a registry of enabled fusion patterns and an explicit ordered pattern list. A pattern
+can be EP-bundled or supplied by an optional independently updated package, but both enter the same
+registry and selection path. EP options express permission and priority, not proof of eligibility.
+Keep existing cross-EP priority; this design does not let an external package claim nodes on behalf
+of an unrelated provider.
+
+During `GetCapability`, the EP matches the normalized graph and describes executable candidates.
+A candidate identifies a source-node set and a prepared implementation recipe, not a second
+complete model that ORT must execute to compare performance. Different patterns may cover the same
+source nodes or emit different replacement DAGs. Standalone execution remains a candidate where
+supported; a region executable only after fusion need not first have standalone target-EP kernels.
+
+Use the existing `ComputeCapability` application model, including `nodes_to_optimize` and an
+ORT-controlled `optimization_func` adapter, rather than add an unrelated post-partition callback.
+Existing registry optimizers can be shared through `GraphOptimizerRegistry`. Neither `Graph&`
+nor `std::function` crosses a future external ABI.
+
+**This is reuse, not an assertion that the current API is sufficient unchanged.** The current
+partitioner checks availability, runs child optimizations, and places/fuses the updated parent.
+Explicit candidate ordering, nonmutating recipe descriptions, arbitrary replacement-DAG capability
+updates/splitting, graph-revision validation, and external registration still need supported adapters.
+The first structural probe must verify that emitted adapters and existing contrib nodes receive
+normal kernel registration/finalization without incorrectly becoming one compiled synthetic node.
+
+### Proposed lifecycle and ordering
+
+The logical lifecycle is:
+
+```text
+canonical load -> shared preprocessing / L0-L1
+ -> EP GetCapability: describe candidates from EP-ordered patterns
+ -> ORT checks availability, eligibility, and ordered conflicts
+ -> [budgeted integration: query recipe requirements before final selection]
+ -> provisionally select compatible candidates
+ -> apply selected recipes through capability optimization adapters
+ -> reconcile graph, capabilities, kernels, and resource information
+ -> finalize placement -> kernel creation / PrePack -> Run
+```
+
+Within an EP, sort candidates first by the EP's pattern-list ordinal, then by stable graph traversal,
+anchor node ID, and candidate ID. Package registration/load order must not override this list.
+Process valid candidates greedily in that order; an accepted candidate excludes overlapping ones,
+including parent whole-Scan versus descendant rewrites. A declined candidate does not claim nodes
+or consume reservations, so a lower-priority feasible candidate can still be considered.
+Priority never overrides schemas, semantics, ownership, kernel availability, or memory feasibility.
+There is no original-versus-fused speed tournament, online benchmark, or optimal-combination search.
+
+Pattern selection requires explicit enablement and optimization level at least L2. Absent, disabled,
+and Basic-only configurations preserve ordinary behavior. Required-pattern policy at Basic fails
+configuration validation. Initial scope is ordinary ONNX inference initialization; ORT-format
+loading, compile-only sessions, training, online repartition, and optimized-cache replay need
+separate support. Optional use on those paths is diagnosed as not applicable; required use fails.
+
+`GetCapability` can be called again as graphs/providers/layouts change. Discovery must therefore be
+read-only and repeatable, not globally once-only. Key prepared candidates by graph revision,
+ownership view, EP instance, effective configuration and workload bounds. Do not reuse stale matches
+or double-charge resources on another pass. Apply only a selected candidate valid for that revision;
+its application callback must execute the chosen recipe, not rerun matching to select another one.
+Install the necessary graph/accounting update support before application; do not depend solely on
+callbacks currently installed after partition.
+
+Publish positive and near-match fixtures at capability discovery, not only untouched exporter
+topology. Preserve unrelated private transformer behavior. Later transforms may further optimize
+emitted nodes, but must not silently invalidate selected kernel/state/resource contracts.
+
+### Memory estimation and placement planning
+
+Baseline delivery can first exercise ordered capability selection without claiming a memory bound.
+The next integration connects target-specific estimators to those same recipes, not a separate
+post-partition fusion pipeline. A budgeted selector evaluates memory feasibility before accepting
+a candidate, then considers the next EP-ordered candidate if the higher-priority one is infeasible.
+The baseline order-only selector does not need a full CPU/GPU planner to demonstrate fusion.
+
+For the future planner in [#33372](https://github.com/microsoft/onnxruntime/pull/33372), supply tentative
+EP ownership without final assignment. Describe logical graph effects and allocation requirements
+before selecting placement; the fused implementation may fit where unfused source estimates do not.
+Queries must not mutate the live graph, construct kernels, upload weights, or materialize trial
+device resources. Apply a provisional winner on a candidate graph, reconcile transfers/casts and the
+execution/backing plan, verify, then commit final placement before materialization.
+Placement-sensitive L2+ passes must participate as selected recipes or be gated in this mode.
+This planner is proposed follow-on work, not current ORT behavior or a prerequisite for baseline.
+
+### Responsibility boundaries for future offloading
+
+Fusion supplies eligible implementation recipes and their resource contracts; P owns graph-level
+CPU/GPU region and boundary selection. An expert-offload controller can manage residency and CPU/GPU
+execution inside an admitted node without assigning separate graph nodes to the CPU EP, as proposed
+by the QMoE [controller](https://github.com/microsoft/onnxruntime/pull/33451) and
+[backend](https://github.com/microsoft/onnxruntime/pull/33452). Execution/capture consumes the
+reconciled placement, as in the [partitioned-capture prototype](https://github.com/microsoft/onnxruntime/pull/32718).
+These proposals are integration references, not implemented compatibility or first-release dependencies.
+Not every offloading decision belongs in `GetCapability`.
+
+Reuse the existing L1/L2 transformer infrastructure and preserve existing fusions; adopting the
+rewrite engine does not require migrating them to capability discovery. The fusion adapter must not
+become a second global placement planner or budget owner. A node-local controller operates within
+resources admitted by the shared resource owner; it cannot independently spend the same headroom.
+
+## 4. Recipe contracts and independent delivery (D2)
+
+### Logical descriptors, not a frozen ABI
+
+Define the capability integration before fixing external function tables. The previous draft's
+parallel `OrtRF*` inspection/registration table is not the selected interface for this revision.
+Reuse `OrtGraph`, `OrtNode`, and `OrtValueInfo` for external read-only inspection where applicable,
+adding only missing revision, EP-instance, workload, and bounded-constant-access contracts.
+No new public API names, binary layouts, experimental factory, or loading entry point are approved
+by this document. The first prototype can use internal adapters, but cannot claim independent
+binary deployment on that basis.
+
+The following describes required information, not new C/C++ structs:
+
+| Descriptor | Required information |
+|---|---|
+| EP pattern registration | Pattern/contract ID and revision, implementation/package identity and revision, enablement, explicit EP-owned priority ordinal, retained callback state. |
+| Candidate context | Graph/revision, source-node IDs, anchor, target EP instance and current/tentative ownership view, immutable effective session configuration and workload bounds. |
+| Prepared recipe | Candidate ID, complete consumption closure, boundary inputs/outputs/captures, logical replacement DAG or compiler-private implementation description, attributes/types/opsets/constants, dependencies, and retained application state. |
+| Requirement query | Selected recipe/kernel/build/device, dtype and workload/alias/concurrency bounds; target-owned requirements with an explicit unknown result, not package-asserted byte counts. |
+| Selection/application | Exact candidate/recipe identity and revision, selected contracts, validated boundary mapping, and an application handle that cannot substitute another implementation. |
+
+ORT must bound candidate count, descriptor storage, and constant reads; recommended initial caps
+remain 8,192 candidates, 64 MiB descriptor storage, and 1 MiB new/read constant payload.
+Metadata inspection must not materialize 27B weights. Existing initializers can be referenced without
+copying payloads; new constants become ORT-owned. Preserve shared storage and every observable
+boundary value identity. A consumed CPU compute node is not made eligible by a CPU-resident input.
+
+Freeze effective configuration and the ordered registry for each session; another session may use
+a different order or workload envelope without sharing mutable candidate state. Reject conflicting
+duplicate pattern identities rather than let package loading silently replace priority.
+Required-pattern policy means at least one accepted candidate of that qualified pattern, not a
+promise that every occurrence fuses. A candidate discovered or retained is not yet selected.
+
+### Delivery, ownership, and lifetime requirements
+
+EP-bundled patterns and optional independent packages must obey the same recipe/selection model.
+An external package registers into the target EP's registry; it does not become a competing EP,
+claim nodes directly, or establish priority by load order. No exporter or model weights are required
+in the runtime package. EP-owned registration does not require matcher updates to ship with the EP.
+
+Before external delivery, specify and qualify a version-negotiated registration/recipe adapter,
+including which existing C/EP APIs it reuses and which minimal additions are necessary. Do not expose
+`Graph&`, STL, protobuf objects, provider pointers, CUDA streams, or kernel factories across that
+boundary. Compiler-private lowering and host-visible rewrites must normalize to the same selected
+recipe contract; compiler EPs need not adopt the pilot to continue working.
+
+Retain the following requirements when defining that adapter:
+
+- Application-controlled static registration and explicit library loading; no directory scanning or
+ executable paths taken from model metadata. Native packages are trusted application code.
+- Atomic registration and cleanup on failure; copied descriptors and explicit retain/release rules
+ for package, candidate, and recipe state. Graph inspection handles are callback-scoped borrows;
+ retained recipes must not contain raw pointers that survive a graph revision.
+- SessionOptions cloning/configuration isolation and immutable per-session registry snapshots.
+ Release candidate/session state before package callbacks or libraries can be unloaded.
+- Serialized, non-reentrant callbacks within a session and safe concurrent initialization across
+ sessions. No asynchronous callback outlives its call; shared immutable registrations are thread-safe.
+- No exceptions crossing a C boundary, explicit `OrtStatus` ownership, unchanged out-parameters on
+ failure, and release callbacks that cannot fail.
+- Explicit errors for incompatible versions, failed loads/registration, malformed recipes, and
+ callback errors. Optional absence is an application choice, not permission to swallow errors.
+ Missing optional patterns leave ordinary execution; required unavailability fails.
+
+Independent DSO conformance and a supported ORT/EP/build/device/package matrix are delivery gates,
+not prerequisites to choosing the capability architecture. Exact public signatures remain a
+follow-on approval item; this revision does not promise indefinite cross-release ABI stability.
+
+### Workload eligibility
+
+Workload bounds are application-declared in both resource modes. Validate `0 <= min <= max`,
+input names, uniqueness, and ranks before discovery. Matchers and requirement queries see the same
+immutable validated session view, not mutable options or another session's shared package state.
+The eventual implementation must reject out-of-envelope feeds/bindings before any kernel launches.
+Graph symbolic dimensions alone are not bounds and missing bounds remain unknown.
+
+Eligibility must hold for the entire fixed-shape or bounded envelope. No package Run callback,
+dynamic match/defuse, or data-dependent fallback is added. For example, a session allowing
+`B=2,P>0,S>1` cannot receive GQA just because its first call is prefill. Use B=1, or separate
+B=2 prefill (`P=0`) and decode (`S=1`) envelopes with the same canonical file.
+Mask contents and arbitrary length-tensor values cannot be certified by shape bounds.
+
+### Matcher choice is a demo-stage decision
+
+W's composable pattern engine is an optional implementation, not the system interface. Its
+standalone `Pattern::MatchAt`/`VisitMatches` can inspect matches without applying a rule; standalone
+matching does not enforce the rule engine's EP/opset policy. `Replacement::Instantiate` constructs
+a fragment, not a complete boundary-preserving replacement transaction. Those components may be
+adapted behind the EP registry, but do not supply memory admission or an external ABI.
+
+Do not make adopting that DSL or merging W a system-design prerequisite. During the Qwen demo,
+first use small exported Attention/convolution/Scan fixtures to measure pattern expressiveness,
+adapter effort, topology survival, and matching overhead; then validate the full model.
+A different matcher must produce the same recipe/selection contract. Demo results decide practical
+adoption, not whether the architecture needs a second optimizer pipeline.
+
+## 5. Host validation and commitment (D3)
+
+1. Freeze the session's EP registry/order and configuration. Discover bounded candidate descriptions
+ without graph mutation or resource charges; retain them only for their graph/ownership revision.
+ Enforce count/byte caps and unique candidate identities within that revision.
+2. Validate IDs, current revision, graph scope, acyclicity, schemas/import versions, type/shape
+ compatibility, constants, and every use crossing the region boundary. Preserve all escaping
+ outputs, including graph outputs and descendant captures. Missing observable state is not dead
+ just because it has no ordinary local consumer. Do not raise an existing domain import and
+ reinterpret surviving nodes; reject incompatible requests. Adding an otherwise absent target
+ domain import is allowed only after validating all emitted schemas.
+3. Require the complete consumed executable region to be available for the target EP instance.
+ Unassigned nodes can be claimed if the proposed implementation supports them; final assignment
+ of every source node is not a discovery prerequisite. A tentative ownership view can establish
+ same-EP eligibility. Never steal work committed to another EP, enter compiler-owned opaque
+ regions, or use host-resident constants as permission to consume CPU-assigned compute.
+4. Resolve replacement kernels through the **actual session registry precedence**, without kernel
+ construction. Validate schema/type constraints, build pruning, host-input requirements, device
+ capabilities, and package-qualified operator contract. V1 declines any custom-registry override
+ of an emitted kernel, even one with the same op name; it does not bypass normal lookup. Core
+ records the selected built-in registration privately and checks it again at finalization.
+5. In budgeted mode, obtain target-specific coverage before selection/commit (section 8).
+6. Preserve ordinary EP priority, then process candidates in the EP-owned pattern order and stable
+ tie-break order from section 3. Choose compatible, non-overlapping feasible candidates. Consumption
+ closure, including Scan descendants, defines overlap. Rejected candidates do not reserve nodes
+ or bytes. Lower-priority alternatives are available before application, not as recovery from
+ an already failed mutation. No per-session benchmark or cost tournament is added.
+7. Validate the accepted **combined** replacement against a disposable graph representation before
+ touching the live graph. This catches cycles and boundaries introduced by interactions between
+ individually valid recipes. Through a validated capability adapter, ORT materializes selected
+ nodes, updates/splits the parent capability as needed, and preserves boundary values. New compute
+ nodes must have target kernels; host adapters are explicit and checked, not hidden callbacks.
+8. Apply only the prepared recipes, update capabilities/identities/accounting, and Resolve. Reconcile
+ the actual implementations before final placement/resource commitment. Unexpected application or
+ Resolve failure aborts baseline initialization; the pilot promises no live-graph rollback or Run
+ retry. A future planner may reject an isolated candidate graph under P's separate contract, but
+ cannot continue with a partially mutated live graph or recurrent state.
+
+Unsupported kernels, masks, ownership, or resource coverage detected before commitment leave the
+original region intact with a reason, unless a required pattern cannot be satisfied. If neither a
+selected recipe nor ordinary execution covers the region, initialization fails; leaving it unchanged
+does not manufacture missing kernel support.
+Distinguish `NO_MATCH`, `EP_MISMATCH`, `OPAQUE_REGION`, `UNSUPPORTED_MASK/STATE/TYPE`,
+`KERNEL_MISSING/OVERRIDDEN`, `RESOURCE_UNKNOWN/INSUFFICIENT`, `OVERLAP`, `STALE_SNAPSHOT`,
+`INVALID_PROPOSAL`, and `COMMIT_FAILED`. Stale/malformed proposals and callback failures are errors,
+not silent non-matches.
+
+Diagnostics identify artifact digest, package/revision, pattern contract, graph/revision, consumed
+and emitted nodes, EP pattern priority, target instance diagnostic ID, selected registration,
+rejection reason, and later transformations. Default logging includes no tensor payload, library
+search path, or weight data. An opt-in trace records timing and resource provenance.
+
+## 6. Pilot source-to-target contracts (D4)
+
+The mappings below are package pattern contracts, not model names in the core API. They require
+structural and semantic qualification; registration of a kernel establishes feasibility only.
+Canonical ONNX inputs/outputs remain unchanged, including all cache and recurrent state values.
+Adapters are ordinary existing operations in the replacement DAG, not hidden package execution.
+These are demo/implementation qualification gates, not prerequisites to selecting the capability
+architecture or a particular matcher DSL.
+
+### 6.1 Canonical Attention to GroupQueryAttention
+
+Use M's `_attention.py::_apply_attention` and
+[`rewrite_rules/_group_query_attention.py::AttentionToGQA:680-833`](https://github.com/onnxruntime/mobius/blob/b1e9c8e9dec1dfa206183f957e77d07cfb734df2/src/mobius/rewrite_rules/_group_query_attention.py#L680-L833)
+as prior art, not as a semantic oracle.
+Qwen35Attention's offset Q/K normalization, partial MRoPE, projection packing, and sigmoid output
+gate stay outside the first region. Target `com.microsoft::GroupQueryAttention:1`, `do_rotary=0`,
+no packed QKV, no local window, no quantized cache, no absorbed normalization.
+CUDA activation registration is FP16/BF16, not FP32
+([R registration](https://github.com/microsoft/onnxruntime/blob/1988019b6e23ac4def4628cddd603c1bfb504fe8/onnxruntime/contrib_ops/cuda/bert/group_query_attention.cc#L56-L84)).
+V1 qualifies **FP16 only**; FP32 remains a baseline/non-match. BF16 is a later independently gated
+package contract even though the kernel is registered.
+
+| Canonical boundary | Emitted boundary / condition |
+|---|---|
+| Q `[B,Hq,S,D]`, K/V `[B,Hkv,S,D]` or proven equivalent 3-D form | Transpose/reshape to query `[B,S,Hq*D]`, key/value `[B,S,Hkv*D]`; inputs 0/1/2. Require `Hq % Hkv == 0`, equal K/V head dimensions, proven head ordering. |
+| `q_num_heads`, `kv_num_heads`, Attention scale | GQA `num_heads=Hq`, `kv_num_heads=Hkv`; materialize exact scale, including default `1/sqrt(D)`. Do not change scaling or softcap/softmax semantics. |
+| Dynamic past K/V `[B,Hkv,P,D]` | Inputs 3/4; present outputs 1/2 `[B,Hkv,P+S,D]`. No paged/static/windowed/sliding/borrowed cache, quantization, or fixed-capacity tail semantics. |
+| Per-row total valid length `L_b` | `seqlens_k[b] = L_b - 1`, INT32 device input 5; scalar `total_sequence_length = max_b L_b`, INT32 CPU input 6. These are not a mask tensor or past lengths. |
+| Causal current chunk | GQA attribute `causal=1` (not the internal `is_unidirectional` name); prove query position i attends keys through `P+i`. Reject top-left/bottom-right alignment disagreement. |
+| Y and cache outputs | Output 0 is `[B,S,Hq*D]`; restore canonical Y layout and preserve exact present-K/V value IDs. |
+| Optional Attention diagnostics | Reject an observed QK/softmax/LSE or other unsupported output. CUDA GQA cannot return QK; do not erase observable outputs. |
+
+**First mask envelope:** uniform dense rows, one shared `P`, all `S` query tokens valid, `L_b=P+S`,
+and no explicit mask or a structurally proven redundant causal mask. Allow prefill `P=0` and
+single-token decode for dense uniform B; continuation chunks `P>0,S>1` require **B=1** under the
+[existing helper constraint](https://github.com/microsoft/onnxruntime/blob/1988019b6e23ac4def4628cddd603c1bfb504fe8/onnxruntime/contrib_ops/cpu/bert/group_query_attention_helper.h#L429-L445).
+`B>1,P>0,S>1` is a non-match even with equal lengths. Prove INT32 arithmetic before narrowing: `0 <= P`,
+`1 <= S`, `P+S <= INT32_MAX`, and all shape/offset products fit their kernel interfaces.
+Derive host placement from the selected kernel definition; copies are inserted by normal ORT
+logic and included in resource/performance accounting. Never pull arbitrary device data to the
+host during discovery or infer runtime mask contents from a symbolic shape. A graph input mask that
+happens to contain ones in a test is not structurally proven redundant. V1 declines arbitrary
+runtime mask inputs; supporting a checked runtime predicate before any state mutation would require
+a separately approved adapter/error contract. A workload promise alone must not erase mask semantics.
+
+For padding, map a standard prefix-valid mask to `L_b` only under a **separate qualified contract**
+that proves query alignment, cache padding, output values for padded queries, and per-row lengths.
+V1 rejects different per-row valid lengths, left/interior padding, arbitrary additive bias,
+noncausal mode, unknown mask provenance, and fully masked rows. Finite large-negative values are
+not automatically equivalent to `-infinity`. Mask conversion must preserve original error/bounds
+semantics; `Cast(INT64->INT32)` alone is not range validation. Unproven cases remain unfused.
+
+The smallest blocker is a tiny capability-discovery Attention fixture with `P>0,S>1`, plus a ragged
+near-match, recording source ownership and causal/output equality. Acceptance of uniform dense
+GQA does not approve broader padding, BF16, cache aliases, or absorbed rotary.
+
+### 6.2 Causal convolution
+
+M's [`functions/causal_conv.py:21-145`](https://github.com/onnxruntime/mobius/blob/b1e9c8e9dec1dfa206183f957e77d07cfb734df2/src/mobius/functions/causal_conv.py#L21-L145)
+supplies the source history/Conv/Swish form. Target existing
+`com.microsoft::CausalConvWithState:1`:
+
+| Source | Target |
+|---|---|
+| Current x `[B,C,S]`, history h `[B,C,K-1]`, `Concat(h,x)` along sequence | input 0 = x; optional input 3 = h; missing history means zero history only if source agrees. |
+| Depthwise Conv weight `[C,1,K]`, optional bias `[C]` | inputs 1/2; `group=C`, stride 1, dilation 1, and exactly the source causal padding/cropping. |
+| `z * Sigmoid(z)` after bias | `activation="swish"`; do not absorb differently ordered gating. |
+| Trailing `K-1` values of **raw** concatenated history/current input | present_state output 1, not post-convolution or activated values. Output 0 keeps x's shape. |
+
+V1 requires `ndim=1,channels_last=0,dilation=1,state_window=0`, dense FP16, no extra padding/stride
+or side uses of consumed intermediates. A Concat already providing history plus a valid Conv
+must not also receive another left pad. Check every Slice axis/start/end/step and Conv attribute,
+not just the operator names. For `S INT32_MAX`:
+decline GQA during matching without allocating tensors at those maximum dimensions.
+
+### Reduced decoder configuration and export recipe
+
+Use the pinned M config fixture, not a remote model identifier, through
+`transformers.Qwen3_5Config.from_dict`. For an exact small **text-only** test configuration,
+override the following `text_config` fields, retaining all other values from that fixture:
+
+```json
+{
+ "hidden_size": 256,
+ "intermediate_size": 512,
+ "num_hidden_layers": 2,
+ "layer_types": ["linear_attention", "full_attention"],
+ "full_attention_interval": 2,
+ "num_attention_heads": 2,
+ "num_key_value_heads": 1,
+ "head_dim": 256,
+ "linear_num_key_heads": 2,
+ "linear_num_value_heads": 4,
+ "linear_key_head_dim": 32,
+ "linear_value_head_dim": 64,
+ "linear_conv_kernel_dim": 4,
+ "mamba_ssm_dtype": "float32",
+ "vocab_size": 256,
+ "bos_token_id": 1,
+ "eos_token_id": 2,
+ "max_position_embeddings": 128
+}
+```
+
+Keep full-attention head dimension 256 so the retained partial rotary width is 64, exactly
+`2*sum([11,11,10])`; do not shrink the head without also redesigning the MRoPE fixture.
+Preserve RMS epsilon, offset normalization, attention output gating and dense MLP.
+For local configuration validation also remap top-level image/video/start/end token IDs to
+`240/241/242/243`, but never feed multimodal tokens or run a vision graph. Save with
+`config.save_pretrained` to an absolute local config directory. The eventual reduced test must
+build the text path only; a whole-VLM graph export is not text-only qualification.
+
+Use PyTorch seed `20261015` to construct random FP32 parameters with the pinned test's initializer
+policy, then apply the FP16 export policy while preserving FP32 recurrence. Initialize RMS offset
+weights according to the model, not arbitrary Gaussian scales. Use prompts `[1,17,23,41,5,8,9]`
+and `[1,7,11,19,31,6,13]` for B=2, no padding; positions are three identical text-only aranges.
+Compare extension-off/on with 7-token prefill, chunks `[2,1,4]`, then 16 greedy feedback tokens,
+all convolution and recurrent states, and both attention caches. At M, the decoder forwards the
+runtime `attention_mask` to `create_attention_bias`
+([`qwen35.py:282-286,492-496`](https://github.com/onnxruntime/mobius/blob/b1e9c8e9dec1dfa206183f957e77d07cfb734df2/src/mobius/models/qwen35.py#L492-L496)).
+The helper uses mask contents to construct the additive bias
+([`_common.py:194-263`](https://github.com/onnxruntime/mobius/blob/b1e9c8e9dec1dfa206183f957e77d07cfb734df2/src/mobius/components/_common.py#L194-L263)).
+Under section 6.1's V1 contract, assert that this decoder's Attention remains unfused, including B=1.
+All-one feeds, no-padding prompts, and prefill/decode shape bounds do not prove structural mask
+redundancy. Do not remove or ignore the canonical mask to force a positive result. Decoder GQA
+qualification needs a separately approved mask contract; numerical equality alone is not coverage.
+
+Use the mask-free small graphs in section 9.2 for GQA-positive tests. Run chunked positives at B=1;
+for B=2 use separate prefill/decode envelopes or require GQA to remain unfused in a general dynamic
+session. Verify accepted/rejected regions in the graph trace rather than infer fusion from outputs.
+
+Public prior art for local config construction and random-weight binding is
+[`qwen38_standard_test.py:27-50,333-371`](https://github.com/onnxruntime/mobius/blob/b1e9c8e9dec1dfa206183f957e77d07cfb734df2/src/mobius/models/qwen38_standard_test.py#L333-L371).
+Its `_from_config`, preprocessing and binding calls are **pinned test internals**, not a stable
+public random-weight CLI. Reuse their construction in the authorized qualification fixture and
+record Transformers/ONNX/NumPy/PyTorch dependency pins from M's approved test environment.
+Missing dependencies require separate authorization, not installation by this design task.
+The new `RuntimeFusionCudaTest.TinyQwen35Generation` target must provide this bound-weight fixture;
+a `--no-weights` graph by itself cannot provide numerical evidence.
+
+An existing public weight-free structural route is
+[`python -m mobius build --config ... --ep onnx-standard --dtype f16 --no-weights`](https://github.com/onnxruntime/mobius/blob/b1e9c8e9dec1dfa206183f957e77d07cfb734df2/src/mobius/integrations/ort_genai/qwen38_standard_test.py#L73-L111).
+The text-only feature must be selected for the text campaign. It does not provide extension
+registration or qualify native GenAI loading. Inspect the resulting standard graph, inline
+functions through the standard-only export path, and verify no contrib nodes remain.
+
+### 9.3 Numerical acceptance gates
+
+Three comparisons are required: canonical extension-off versus on; each versus an independently
+constructed semantic reference; and, where applicable, the pinned ModelBuilder contrib path.
+Two graphs calling the same kernel are **not** independent proof. Check intermediate tensors,
+every state at every chunk/token boundary, and full logits, not just top-1 text.
+Use elementwise `abs(actual-reference) <= atol + rtol*abs(reference)` and require all values finite.
+Structural, length, mask decisions and copied cache/history values must match exactly.
+
+| Boundary dtype | Intermediate/output atol, rtol | Recurrent state atol, rtol | Reduced full logits atol, rtol | Status |
+|---|---|---|---|---|
+| FP32 | `1e-5,1e-4` | `1e-5,1e-4` | `1e-4,1e-4` | Canonical/reference control; GQA remains a non-match. |
+| FP16 | `2e-2,2e-2` | `2e-2,2e-2` | `5e-2,2e-2` | First qualified activation type; state gates apply to actual stored state dtype, including FP32 state. |
+| BF16 | `5e-2,5e-2` | `5e-2,5e-2` | `1e-1,5e-2` | Reserved gate for separate future qualification, not V1 support. |
+
+These are proposed gates to ratify **before** running, not measured accuracies. Do not loosen
+them after inspecting failures without an explicit reviewed contract change. Additionally report
+max absolute/relative error, RMS error, first divergence and logit margin. Greedy feedback tests
+use lowest-token-ID tie breaking and feed each path's generated token back to that path; require
+identical 16-token sequences for the tiny fixture as well as the tensor gates. A small logit margin
+does not excuse a changed token. Full official-model HF numerical limitations are separate known
+qualification concerns, not erased by a tiny pass or by loose FP16 bit-identity claims.
+
+### 9.4 Proposed commands and artifact update experiment
+
+Commands below require separately authorized execution and **already-built**, absolute-path
+targets. New suites and their fixtures must first be implemented. No build/install command is
+part of this assignment. Fixture tests must fail if required CUDA/DSO prerequisites are absent,
+and record the full runtime/EP/package/device pins and test count.
+
+```bash
+# Planned only. Set ORT_BUILD_DIR to an absolute directory containing authorized existing binaries.
+timeout 120s "${ORT_BUILD_DIR:?}/onnxruntime_test_all" \
+ --gtest_filter='RuntimeFusionTests.*:InferenceSessionTests.TestRegisterTransformers'
+timeout 120s "${ORT_BUILD_DIR:?}/onnxruntime_shared_lib_test" \
+ --gtest_filter='RuntimeFusionApiTest.*:ExperimentalCApiTest.*'
+timeout 300s "${ORT_BUILD_DIR:?}/onnxruntime_provider_test" \
+ --gtest_filter='RuntimeFusionCudaTest.*'
+# Planned structural export only, using M's existing CLI in an authorized pinned environment.
+python -m mobius build --config /tmp/ort-runtime-fusion/tiny-config \
+ /tmp/ort-runtime-fusion/tiny-standard --ep onnx-standard --dtype f16 \
+ --features text-only --no-weights
+```
+
+Run the smallest relevant fixture first using an exact test filter, then the family. Existing
+attention/stateful/resource test names must be enumerated from the chosen binary before scheduling;
+do not interpret an empty GoogleTest filter result as success. Sanitizer-enabled lifetime tests,
+CUDA memory checks, and CPU/WebGPU implementation campaigns require their own authorized builds.
+
+Package-update proof: serialize the tiny canonical graphs once, hash the model and every external
+data file with SHA256, then initialize with no package, explicit disable, package revision A and B.
+A recognizes one approved post-L1 topology. B additionally recognizes a proven redundant
+reshape/identity adapter form already in the **same saved artifact**, without changing the target
+op or numerical envelope. Record increased eligible coverage, graph trace and the same hashes.
+Ensure L1 actually preserves this distinguishing form; if not, the topology-survival probe blocks
+this experiment until an appropriate form is chosen. Do not re-export between A and B.
+Unavailable optional package leaves ordinary execution; required-package policy is an application
+load error. Reverting B to A must not require editing the canonical files.
+
+### 9.5 Performance envelopes
+
+Before an authorized run, record runtime/EP/package/compiler/device/driver pins, clocks/power policy,
+CPU affinity, streams, warm-up, repetitions and synchronization. No per-session microbenchmark
+chooses a fusion. Offline package qualification decides whether to ship the eligible contract.
+
+For tiny fixtures, require bounded matching storage (at most configured proposal/descriptor caps),
+and record time per graph/node rather than claiming a meaningful throughput speedup.
+For the reduced decoder and later text-only 27B campaign, pre-register these proposed envelopes:
+batch `{1,2}`, prompt `{128,512,2048}`, decode `128` generated tokens, dense FP16, nonpaged state,
+no speculative decoding, one session/stream initially; `max_total_length=2176`.
+Use a B=1 dynamic session or separate batched prefill/decode sessions where required by GQA.
+The reduced fixture uses its own context limit 128 and 16 generated tokens, not the full
+27B lengths listed here.
+Run extension-off/on with identical weights, inputs, generation settings and hardware.
+
+| Metric | Proposed release gate, ratified before collection |
+|---|---|
+| Matching/validation wall time | At most `max(50 ms, 5% of extension-off initialization)` on the same machine; separately report graph copying/Resolve. |
+| Initialization and PrePack | Report cold/warm init and persistent/scratch bytes; no hidden exclusion of load, upload or packing costs. More than 10% total-init regression needs explicit acceptance. |
+| TTFT | No greater than 5% regression in any qualified envelope, including initialization-inclusive and steady-state figures separately. |
+| Decode | Target at least 10% throughput gain in the declared pilot workload; publish per-case results and confidence intervals, not a universal gain claim. |
+| Transfers and peaks | Record H2D/D2H/D2D volume, host/pinned/device live and retained peaks, cache/state growth and initialization overlap. Baseline permits no hard-budget claim; budgeted mode must stay within the exact user-admitted limits. |
+
+Use 3 warm-ups and 10 measured repetitions per case, report median and p95 and all raw run identifiers.
+If the speed target is not met, report feasibility without claiming a performance success; it does
+not authorize changing ONNX semantics, runtime speed competition, or the user's memory budget.
+
+### 9.6 Official checkpoint campaign admission
+
+M5 requires a fresh export, not an assumed cached ONNX. Before authorization, approve a local
+read-only HF snapshot/weight revision and file manifest, verify the official raw config hash in
+section 6.4, pin M and its complete exporter dependency environment, and document the deliberate
+BF16-to-FP16 conversion policy. Export text-only, standard-only, dynamic/nonpaged cache, FP32
+recurrent state, no static-cache feature, no speculative state window. Use the pinned public CLI's
+local source route; a proposed invocation after authorization is:
+
+```bash
+# Planned only; approved-hf-snapshot must already have independently authorized local assets.
+python -m mobius build --model /tmp/ort-runtime-fusion/approved-hf-snapshot \
+ /tmp/ort-runtime-fusion/official-text-standard --ep onnx-standard --dtype f16 \
+ --features text-only
+```
+
+Check no contrib nodes or unexpanded local functions remain. Pin the ModelBuilder comparator at G
+with `linear_attn_op=linear_attention`, nonpaged mode and `state_window=0`, recording its actual
+exporter options, output state layout and all file hashes; do not invent an unverified GenAI CLI
+switch or claim canonical/contrib equivalence before comparison. Compare to HF independently,
+retaining historical numerical differences as separate failures/limitations.
+
+The authorization must name GPU/device count, CUDA/ORT build, host/device peak and storage caps,
+transfer allowance, maximum run time, number of sessions, checkpoint/artifact retention period,
+and who deletes intermediates. Reject the campaign until those concrete limits are approved;
+this document allocates nothing. Retain small digests/configs/logs and required numerical evidence,
+not redundant weight/intermediate caches. Vision/full-VLM needs separate configs, budgets and
+correctness gates and is not covered by a text-only success.
+
+## 10. Approval ledger and smallest blocked probes
+
+The October 26 design direction is EP-side registration and EP-owned pattern order. This is not
+maintainer approval of the adapters or their implementation. The named proposed approval coordinator
+for D1-D6 is
+**@titaiwangms**, the issue author/assignee. This is a routing proposal, not a claim of delegated
+authority or consent. The coordinator must confirm an individually named responsible ORT maintainer
+for each row, with subsystem/API co-review as listed. Source inspection supplies no authoritative
+individual delegation. Until ratified, **named delegated approval ownership is an attributed
+governance blocker**, and this draft does not satisfy maintainer acceptance. Do not invent reviewers
+or mark an issue-author acknowledgment as public-API approval.
+
+| ID | Concrete proposed decision and rationale | Named routing owner; required approval | Status / smallest evidence gate |
+|---|---|---|---|
+| D1 | Register/order patterns on the EP side; discover through `GetCapability` and reuse capability application. ORT resolves conflicts by EP order, not package load order or speed competition. | @titaiwangms to confirm partitioner/EP maintainer. | Revised direction; P1 must validate ordering, repeated discovery and multi-node replacement support before M1 acceptance. |
+| D2 | Reuse existing graph inspection; define only missing capability/recipe registration adapters. Do not freeze the former parallel C table. | @titaiwangms to confirm C/EP API maintainer; request `@microsoft/onnxruntime-api` review before any public additions. | Public ABI deferred; P2 gates external delivery, not an internal M1 prototype. |
+| D3 | Describe/select/apply separately; validate combined effects and boundaries; update capabilities/accounting and reconcile before final commitment. No baseline live rollback. | @titaiwangms to confirm graph/partitioner and kernel-registry maintainers. | Proposed, unapproved; P3 blocks M1 application acceptance and whole-Scan M3. |
+| D4 | Qualify FP16 GQA/convolution and separately whole-Scan LinearAttention in the model demo. Grama's engine is optional, not a system-design dependency. | @titaiwangms to confirm CUDA/stateful-kernel maintainers, with Mobius/GenAI co-review. | Demo-stage gates P4/P5; matcher fitness does not block the architecture decision. |
+| D5 | Baseline ordered selection first; target estimates add feasibility before selection in budgeted mode. Unknown is not zero; admit before materialization. | @titaiwangms to confirm resource-accountant/session/CUDA recipe maintainers. | Proposed, unapproved; P6 gates budgeted M4, not baseline M1/M2. |
+| D6 | EP owns registry/order; ORT owns selection/validated application; optional package owns matcher/adapter release. One recipe model for bundled and external patterns. | @titaiwangms to confirm ORT/EP ownership and Mobius package/release owner. | Proposed, unapproved; P2/P7 plus support ownership gate independent M4 delivery. |
+
+The coordinator must replace routing placeholders with confirmed named approvers/status and linked
+review decisions before calling the design accepted. An accepted design can retain an explicitly
+owned runtime blocker; it cannot pretend the blocked implementation contract was approved.
+
+### Separately authorized probes
+
+None has run. Public source inspection alone cannot resolve these questions. Use existing binaries
+and tiny local random data after authorization; if missing, request a bounded build scope separately.
+
+| Probe | Smallest work / observable result | Exact milestone blocked |
+|---|---|---|
+| P1: capability integration | Tiny flat fixture with two EP-ordered overlapping patterns and a multi-node replacement; repeat discovery after a revision/layout change. Trace returned/selected capabilities, registered target kernels, no duplicate application or speculative charges. Use Attention/convolution exports for M2, no official weights. | M1 ordering/application adapters; M2 topology survival and package-update choice. |
+| P2: independent ABI | After adapter approval, independently built fusion-only DSO and runtime; EP-side registration, order unchanged by load order, clones, version mismatch, failing registration, concurrent init and unload trace. | M4 external delivery; not a gate for an internal M1 prototype. |
+| P3: application closure | Two overlapping flat recipes and parent Scan/body conflict with an escaping capture; assert selected identity, boundary preservation and parent capability updates. Inject stale revision and Resolve failure; assert no dangling graph/reservation IDs or continuation of partial initialization. | M1 application contract; M3 recursive destruction. |
+| P4: semantics | GQA positives `(B,P,S)=(2,0,2),(2,3,1),(1,3,2)`, D=64; negative `(2,3,2)` and ragged/runtime masks; convolution `K=4,S=1,2`, nonzero history. Compare every output with independent reference at declared gates. | M2 first CUDA mapping acceptance. |
+| P5: recurrence | One standard-only lowered Scan `B=1,Hk=2,Hv=4,Dk=32,Dv=64,S=7`, nonsymmetric FP32 state and FP32 recurrence inputs; trace CUDA parent/body and exact cast/head order; compare whole/chunks/single-token references. | M3 recurrence mapping **before implementation delivery**; mixed dtype/ownership is a blocker, not permission to change casts/reassign. |
+| P6: preventive budget | High-priority recipe exceeding budget and lower-priority fitting recipe; unknown estimate, repeat queries, partial CSV, shared storage and initialization scratch. Verify selection/charges, final reconciliation and admission before upload; compare target declarations/allocations. | M4 budgeted integration; each uncovered target recipe remains blocked. |
+| P7: unchanged artifact | Fixed tiny canonical digest manifest, EP-bundled and external A/B matcher registration with the same explicit EP order, disable/unavailable cases. Record coverage, numerics and load behavior without model changes. | M4 independent-package value proof; W suitability evaluated here, not required. |
+
+No actual full-model export or cache package is assumed available. If P1 shows shape scaffolding
+unavailable to the target EP, shrink the boundary or report non-match; do not consume CPU-committed
+compute to force eligibility. An unassigned node and a node already committed to another EP are
+different cases.
+
+### D6 proposed ownership and support matrix
+
+Place generic selection/application adapters and conformance fixtures in ORT (`core/framework`,
+`core/optimizer`, and framework/shared-library tests). EPs own registration, priority and support
+validation; patterns may reuse shared optimizer implementations without duplicating machinery.
+Keep model-specific pattern knowledge out of the generic planner. Propose the optional, separately
+versioned performance
+distribution **`mobius-ort-cuda-fusion`** under the Mobius project's ownership; this is a name/owner
+proposal, not an existing package, new repository, or published artifact. External delivery must use
+the approved EP registration/recipe boundary, not private runtime symbols or exporter/model weights.
+Existing contrib schemas/kernels stay owned by their existing ORT maintainers. Package owners must
+qualify their supported ORT/EP/build/device tuple and respond to topology drift.
+
+Initial candidate platform scope: full ONNX-enabled Linux x64 and Windows x64 CUDA builds, matching
+native pointer width/calling convention, explicit static or DSO registration, inference only.
+The first implementation may claim only the tuples actually qualified by P2/P4/P5; this proposal
+does not certify either platform. Reduced/pruned builds decline missing kernels, and builds without
+native loading use static registration. Web/WASM cannot load native DSOs; a future native WebGPU or
+CPU adapter needs its own registered kernels, numerical/resource recipes and conformance campaign.
+No automatic package discovery and no universal cross-provider support claim.
+
+## 11. Roadmap and design acceptance
+
+These are planned milestones, **not created implementation children or assigned tasks**. Nothing
+is approved to implement by this draft. Create separate children only after their relevant
+decisions are accepted; attach actual predecessor PRs, approved API/contract revision, bounded
+affected-file scope, acceptance gates and exact verification targets. Dependencies are not schedules.
+
+| Milestone | Scope / acceptance | Predecessor and blockers |
+|---|---|---|
+| M1 capability mechanism | EP-owned ordered registry, repeatable discovery, prepared recipes and validated capability application adapters. Internal prototype permitted; no frozen external ABI or model-specific matcher dependency. `RuntimeFusionTests.*` and existing capability regressions. | D1/D3 and generic D6 approval; P1/P3. D2 external API and P2 are not internal prototype gates. |
+| M2 flat CUDA demo | Register GQA/convolution patterns in the EP; verify kernels, topology, negative masks/dtypes and all outputs/state. Evaluate W on small exported fixtures without making adoption mandatory. | Actual M1 PR and D4 flat-region contract; P1/P4. No recurrence approval implied. |
+| M3 stateful/whole-Scan | Recursive read-only inspection, core-derived closure, body accounting retirement, narrow LinearAttention sequence replacement and asymmetric-state tests. Convolution state qualification extended here. | Actual M1/M2 PRs, D3/D4 whole-Scan acceptance and P5 before sequence implementation; no wider GatedDeltaNet approval. |
+| M4 reduced generation/package/resources | Random-weight text decoder, feedback generation, independent EP-side package registration and A/B artifact proof. Approve/qualify minimal external adapter; separately connect recipe estimators and budget feasibility. | Actual M1-M3 PRs; D2/D6 and P2/P7 for external release. Budgeted part needs D5/P6 and each target recipe. No general planner prerequisite. |
+| M5 official 27B campaign | New pinned text-only export, independent numerical comparison, coverage/performance/resource report with retained provenance. Vision/full-VLM is a separate qualification. | Actual M4 PR/report; explicit weight/export/device/build/retention budget authorization. Historical Mobius #758 is context, not a substitute. |
+
+General placement search remains outside M1-M5. After the shared recipe contracts are accepted,
+P's planner can consume the same EP-ordered candidates, not a second fusion mechanism. Qualify
+query isolation, tentative-ownership/revision invalidation, and an unfused-offloads/fused-fits case;
+verify the selected graph and execution/backing plan before final commitment/materialization.
+Neither that planner nor model execution is authorized by this documentation revision.
+
+### Acceptance record
+
+- This PR supplies a reviewable design, not a qualification receipt or a maintainer approval.
+- Required for design acceptance: responsible ORT maintainer review, confirmed named D1-D6 approval
+ routing/status with rationale, and explicitly attributed blockers/probes tied to milestones.
+- Runtime tests are **not** required to close the design issue, but their absence must remain
+ explicit. A later implementation cannot substitute a planned command, skip, source registration
+ or historical report for executed evidence.
+- Design acceptance does not authorize unresolved API/lifetime, whole-Scan transaction,
+ resource-owner contracts, production coding, new implementation issues or full-model execution.
diff --git a/docs/design/onnxruntime_web_jsep_to_webgpu_ep_migration.md b/docs/design/onnxruntime_web_jsep_to_webgpu_ep_migration.md
index 5c5e4a51c7ca3..18f9a8a7098c6 100644
--- a/docs/design/onnxruntime_web_jsep_to_webgpu_ep_migration.md
+++ b/docs/design/onnxruntime_web_jsep_to_webgpu_ep_migration.md
@@ -250,11 +250,9 @@ Per-session typed options are already a superset of JSEP's, confirmed by source
covering the `.jsep.wasm` → `.asyncify.wasm` filename change, `/jsep` pinning, and the int64 `extra` opt-in.
This is the first user-facing announcement; capture the JSEP usage baseline (§11) before it ships.
-**Release gate.** The flip is gated on both: (a) a **blocking** CI job running the native WebGPU EP — the current
-native-WebGPU-EP legs are advisory (`continue-on-error` in `windows-web-ci-workflow.yml`, `continueOnError` in
-`win-web-ci.yml`) and must be promoted to blocking. The default `.` bundle and `./webgpu` are the same build, so a
-green `/webgpu` run covers the flip only if it exercises proxy/IO-binding/WebNN — which the current op-parity leg
-does not; and (b) the §8 items resolved with targeted coverage (op-parity tests don't exercise those paths).
+**Release gate.** Before switching the default to the native WebGPU EP, resolve the §8 gaps and add blocking CI
+coverage for proxy mode, GPU IO-binding, and WebNN initialization. Passing the existing operator tests alone is
+insufficient.
---
diff --git a/onnxruntime/contrib_ops/cuda/bert/hyper_connection_low_rank_mix_impl.cu b/onnxruntime/contrib_ops/cuda/bert/hyper_connection_low_rank_mix_impl.cu
index 59c9626b0035b..cb7d3ef86fd90 100644
--- a/onnxruntime/contrib_ops/cuda/bert/hyper_connection_low_rank_mix_impl.cu
+++ b/onnxruntime/contrib_ops/cuda/bert/hyper_connection_low_rank_mix_impl.cu
@@ -51,10 +51,12 @@ constexpr int kMaxFusedHidden = kMaxVecIters * kThreads * 8;
constexpr int kMaxFusedRank = 512;
constexpr int kMaxFusedBranches = 8;
+// expf(-|v|) is the argument both branches of the textbook form pass to expf (-v for v >= 0, v for v < 0), so one expf and
+// one division give the bit-identical result (verified over all 2^32 fp32 bit patterns) without the divergent two-way
+// branch, which executed both paths in nearly every warp of the prefill pre-mix.
__device__ __forceinline__ float SigmoidF32(float v) {
- if (v >= 0.0f) return 1.0f / (1.0f + expf(-v));
- const float e = expf(v);
- return e / (1.0f + e);
+ const float e = expf(-fabsf(v));
+ return (v >= 0.0f ? 1.0f : e) / (1.0f + e);
}
template
@@ -227,6 +229,16 @@ __global__ void __launch_bounds__(kThreads, ROWS >= 4 ? 4 : 5)
}
CpCommit();
+ uint4 w_reg8 = make_uint4(0u, 0u, 0u, 0u);
+ uint4 w_reg9 = make_uint4(0u, 0u, 0u, 0u);
+ if (!is_inject && col < p.rank) {
+ const T* base = w_down + k0 * p.rank + col;
+ const int row8 = rp + 128 * 8;
+ if (row8 < len) w_reg8 = __ldcg(reinterpret_cast(base + static_cast(row8) * p.rank));
+ const int row9 = rp + 128 * 9;
+ if (row9 < len) w_reg9 = __ldcg(reinterpret_cast(base + static_cast(row9) * p.rank));
+ }
+
// Norm weights for the 8 consecutive features of the slice this thread normalizes below.
const bool norm_active = tid * 8 < len;
uint4 wv = make_uint4(0u, 0u, 0u, 0u);
@@ -330,7 +342,7 @@ __global__ void __launch_bounds__(kThreads, ROWS >= 4 ? 4 : 5)
float wf[8];
const uint4 w_vec = (i < kLocalAsyncWeightSlots)
? wsm[i * kThreads + tid]
- : __ldcg(reinterpret_cast(base + static_cast(row) * p.rank));
+ : (i == 8 ? w_reg8 : w_reg9);
Unpack8(w_vec, wf);
#pragma unroll
for (int r = 0; r < ROWS; ++r) {
@@ -418,11 +430,11 @@ __global__ void __launch_bounds__(kThreads, ROWS >= 4 ? 4 : 5)
// CTA (hb, 0, z) owns features [16 hb, 16 hb + 16) of every branch for ROWS rows.
template
-__global__ void __launch_bounds__(kThreads, 5)
+__global__ void __launch_bounds__(kThreads, (BRANCHES == 4 && RANK == 320 && ROWS >= 4) ? 3 : 5)
HcUpPreKernel(const T* __restrict__ streams, const T* __restrict__ norm_w, const T* __restrict__ w_up,
const float* __restrict__ partial, const float* __restrict__ inv_rms_in,
T* __restrict__ pre_out, T* __restrict__ gate_out, UpParams p) {
- static_assert(ROWS <= 2, "the cross-warp gate reduction below assumes ROWS * 16 * branches <= kThreads");
+ static_assert(ROWS <= 4, "the cross-warp gate reduction below assumes ROWS * 16 * branches <= kThreads");
constexpr int kLocalWeightSlots = kWeightSlots;
const int hb = blockIdx.x;
const int64_t row0 = static_cast(blockIdx.z) * ROWS;
@@ -445,7 +457,8 @@ __global__ void __launch_bounds__(kThreads, 5)
__shared__ __align__(16) uint4 wsm[kLocalWeightSlots * kThreads];
__shared__ __align__(16) float sf[ROWS][RANK == 0 ? kMaxFusedRank : RANK];
- __shared__ float sg[ROWS][kBranchStorage * kUpTile];
+ __shared__ float sg[(BRANCHES == 4 && RANK == 320) ? 1 : ROWS]
+ [(BRANCHES == 4 && RANK == 320) ? 1 : kBranchStorage * kUpTile];
// Prefetch the first chunk of up weights before waiting for the down kernel.
const T* wbase = w_up + static_cast(c) * hidden + col;
@@ -458,10 +471,55 @@ __global__ void __launch_bounds__(kThreads, 5)
}
CpCommit();
+ constexpr bool kFastEp4x320 = (BRANCHES == 4 && RANK == 320);
+ float nw_pre[kFastEp4x320 ? 1 : kBranchStorage];
+#pragma unroll
+ for (int b = 0; b < (kFastEp4x320 ? 1 : kBranchStorage); ++b) nw_pre[b] = 0.0f;
+ const int r_ep = kFastEp4x320 ? (tid >> 6) : (tid / kUpTile);
+ const int hl_ep = kFastEp4x320 ? ((tid >> 2) & 15) : (tid % kUpTile);
+ const int b_ep = tid & 3;
+ const bool ep_valid = (tid < (kFastEp4x320 ? (ROWS * 64) : (ROWS * kUpTile))) &&
+ ((h0 + hl_ep) < hidden) && RowOk(row0, r_ep, p.rows);
+ if (ep_valid) {
+ if constexpr (kFastEp4x320) {
+ nw_pre[0] = to_float(__ldg(norm_w + static_cast(b_ep) * hidden + h0 + hl_ep));
+ } else {
+#pragma unroll
+ for (int b = 0; b < kBranchStorage; ++b) {
+ if (b < branches) nw_pre[b] = to_float(__ldg(norm_w + static_cast(b) * hidden + h0 + hl_ep));
+ }
+ }
+ }
+
// The next block's down kernel may start (and prefetch its weights) while this kernel waits.
GridTrigger();
GridWait();
+ float xv_pre[kFastEp4x320 ? 1 : kBranchStorage];
+ float inv_pre[kFastEp4x320 ? 1 : kBranchStorage];
+#pragma unroll
+ for (int b = 0; b < (kFastEp4x320 ? 1 : kBranchStorage); ++b) {
+ xv_pre[b] = 0.0f;
+ inv_pre[b] = 0.0f;
+ }
+ if (ep_valid) {
+ const int64_t r_idx_ep = RowIdx(row0, r_ep, p.rows);
+ if constexpr (kFastEp4x320) {
+ xv_pre[0] = BitsToFloat(__ldcg(reinterpret_cast(
+ streams + r_idx_ep * K + static_cast(b_ep) * hidden + h0 + hl_ep)));
+ inv_pre[0] = __ldcg(inv_rms_in + r_idx_ep * 4 + b_ep);
+ } else {
+#pragma unroll
+ for (int b = 0; b < kBranchStorage; ++b) {
+ if (b < branches) {
+ xv_pre[b] = BitsToFloat(__ldcg(reinterpret_cast(
+ streams + r_idx_ep * K + static_cast(b) * hidden + h0 + hl_ep)));
+ inv_pre[b] = __ldcg(inv_rms_in + r_idx_ep * branches + b);
+ }
+ }
+ }
+ }
+
// Fixed-order reduction of the K-split partial sums, then ScaledSiLU; four columns per thread.
const int nq = rank / 4;
for (int idx = tid; idx < ROWS * nq; idx += kThreads) {
@@ -469,16 +527,27 @@ __global__ void __launch_bounds__(kThreads, 5)
const int cq = (idx % nq) * 4;
const int64_t r_idx = RowIdx(row0, r, p.rows);
float4 va = make_float4(0.f, 0.f, 0.f, 0.f);
- for (int sp0 = 0; sp0 < p.splits; sp0 += 8) {
+ const float* p_base = partial + r_idx * p.splits * p.partial_stride + cq;
+ if (p.splits == 8) {
#pragma unroll
for (int u = 0; u < 8; ++u) {
- if (sp0 + u < p.splits) {
- const float4 q = __ldcg(reinterpret_cast(
- partial + (r_idx * p.splits + sp0 + u) * p.partial_stride + cq));
- va.x += q.x;
- va.y += q.y;
- va.z += q.z;
- va.w += q.w;
+ const float4 q = __ldcg(reinterpret_cast(p_base + u * p.partial_stride));
+ va.x += q.x;
+ va.y += q.y;
+ va.z += q.z;
+ va.w += q.w;
+ }
+ } else {
+ for (int sp0 = 0; sp0 < p.splits; sp0 += 8) {
+#pragma unroll
+ for (int u = 0; u < 8; ++u) {
+ if (sp0 + u < p.splits) {
+ const float4 q = __ldcg(reinterpret_cast(p_base + (sp0 + u) * p.partial_stride));
+ va.x += q.x;
+ va.y += q.y;
+ va.z += q.z;
+ va.w += q.w;
+ }
}
}
}
@@ -491,13 +560,23 @@ __global__ void __launch_bounds__(kThreads, 5)
}
*reinterpret_cast(&sf[r][cq]) = make_float4(out[0], out[1], out[2], out[3]);
}
- if (hb == 0 && tid < branches * ROWS) {
- const int r = tid / branches;
- const int t = tid % branches;
+ const int gt_base = (ROWS * nq > kThreads) ? (ROWS * nq - kThreads) : 0;
+ const int gt = tid - gt_base;
+ if (hb == 0 && gt >= 0 && gt < branches * ROWS) {
+ const int r = gt / branches;
+ const int t = gt % branches;
const int64_t r_idx = RowIdx(row0, r, p.rows);
float v = 0.0f;
- for (int sp = 0; sp < p.splits; ++sp) {
- v += __ldcg(partial + (r_idx * p.splits + sp) * p.partial_stride + rank + t);
+ const float* p_base = partial + r_idx * p.splits * p.partial_stride + rank + t;
+ if (p.splits == 8) {
+#pragma unroll
+ for (int sp = 0; sp < 8; ++sp) {
+ v += __ldcg(p_base + sp * p.partial_stride);
+ }
+ } else {
+ for (int sp = 0; sp < p.splits; ++sp) {
+ v += __ldcg(p_base + sp * p.partial_stride);
+ }
}
const float inject = Rnd(v);
const float u = Rnd(SigmoidF32(Rnd(inject / p.gate_div)));
@@ -556,37 +635,51 @@ __global__ void __launch_bounds__(kThreads, 5)
}
}
__syncthreads();
- // The up kernel runs with at most 2 rows, so ROWS * groups * 8 = 16 * ROWS * branches <= kThreads (branches <= 8).
+ // With ROWS <= 4, ROWS * groups * 8 = 16 * ROWS * branches <= kThreads (branches <= 4 for ROWS=4, <= 8 for ROWS<=2).
const int g8 = groups * 8;
- if (tid < ROWS * g8) {
- const int r = tid / g8;
- const int idx = tid % g8;
- float v = 0.0f;
+ if constexpr (kFastEp4x320) {
+ if (tid < ROWS * g8) {
+ const int idx = (b_ep * 2 + (hl_ep >> 3)) * 8 + (hl_ep & 7);
+ float v = 0.0f;
#pragma unroll
- for (int q = 0; q < 8; ++q) v += red[(q * ROWS + r) * g8 + idx];
- const int gg = idx >> 3;
- const int local = (gg & 1) * 8 + (idx & 7);
- sg[r][(gg >> 1) * kUpTile + local] = Rnd(SigmoidF32(Rnd(v)));
- }
- __syncthreads();
- if (tid < ROWS * kUpTile) {
- const int r = tid / kUpTile;
- const int hl = tid % kUpTile;
- const int64_t r_idx = RowIdx(row0, r, p.rows);
- if ((h0 + hl) < hidden && RowOk(row0, r, p.rows)) {
+ for (int q = 0; q < 8; ++q) v += red[(q * ROWS + r_ep) * g8 + idx];
+ const float my_s = Rnd(SigmoidF32(Rnd(v)));
+ const float my_n = Rnd(xv_pre[0] * inv_pre[0] * nw_pre[0]);
+ float sum = 0.0f;
+#pragma unroll
+ for (int b = 0; b < 4; ++b) {
+ const float nb = __shfl_sync(0xffffffffu, my_n, b, 4);
+ const float sb = __shfl_sync(0xffffffffu, my_s, b, 4);
+ sum = fmaf(nb, sb, sum);
+ }
+ if (b_ep == 0 && ep_valid) {
+ const int64_t r_idx = RowIdx(row0, r_ep, p.rows);
+ pre_out[r_idx * hidden + h0 + hl_ep] = from_float(sum * p.reduction_scale);
+ }
+ }
+ } else {
+ if (tid < ROWS * g8) {
+ const int r = tid / g8;
+ const int idx = tid % g8;
+ float v = 0.0f;
+#pragma unroll
+ for (int q = 0; q < 8; ++q) v += red[(q * ROWS + r) * g8 + idx];
+ const int gg = idx >> 3;
+ const int local = (gg & 1) * 8 + (idx & 7);
+ sg[r][(gg >> 1) * kUpTile + local] = Rnd(SigmoidF32(Rnd(v)));
+ }
+ __syncthreads();
+ if (ep_valid) {
+ const int64_t r_idx = RowIdx(row0, r_ep, p.rows);
float sum = 0.0f;
#pragma unroll
for (int b = 0; b < kBranchStorage; ++b) {
if (b < branches) {
- const float nw = to_float(norm_w[static_cast(b) * hidden + h0 + hl]);
- const float xv = BitsToFloat(__ldcg(reinterpret_cast(
- streams + r_idx * K + static_cast(b) * hidden + h0 + hl)));
- const float inv = __ldcg(inv_rms_in + r_idx * branches + b);
- const float n = Rnd(xv * inv * nw);
- sum = fmaf(n, sg[r][b * kUpTile + hl], sum);
+ const float n = Rnd(xv_pre[b] * inv_pre[b] * nw_pre[b]);
+ sum = fmaf(n, sg[r_ep][b * kUpTile + hl_ep], sum);
}
}
- pre_out[r_idx * hidden + h0 + hl] = from_float(sum * p.reduction_scale);
+ pre_out[r_idx * hidden + h0 + hl_ep] = from_float(sum * p.reduction_scale);
}
}
}
@@ -747,6 +840,78 @@ __global__ void PreMixKernel(const T* __restrict__ hn, const T* __restrict__ up,
pre[i] = from_float(sum * reduction_scale);
}
+template
+__global__ void __launch_bounds__(320, 4)
+ FusedPreMixInjectPostGateB4H2560Kernel(const T* __restrict__ hn, const T* __restrict__ up,
+ const T* __restrict__ w_inject, T* __restrict__ pre_out,
+ T* __restrict__ gate_out, int64_t rows, float reduction_scale,
+ float gate_div, float gate_mul) {
+ const int64_t row = blockIdx.x;
+ if (row >= rows) return;
+ const int tid = threadIdx.x;
+ constexpr int kHidden = 2560;
+ constexpr int64_t K = 4 * kHidden;
+ const int h = tid * 8;
+ const T* hn_row = hn + row * K;
+ const T* up_row = up + row * K;
+
+ float sum[8] = {0.f, 0.f, 0.f, 0.f, 0.f, 0.f, 0.f, 0.f};
+ float acc[4] = {0.f, 0.f, 0.f, 0.f};
+
+#pragma unroll
+ for (int b = 0; b < 4; ++b) {
+ const int k_base = b * kHidden + h;
+ float hn_v[8], up_v[8];
+ Unpack8(__ldg(reinterpret_cast(hn_row + k_base)), hn_v);
+ Unpack8(__ldg(reinterpret_cast(up_row + k_base)), up_v);
+ const uint4* w_ptr = reinterpret_cast(w_inject + k_base * 4);
+#pragma unroll
+ for (int m = 0; m < 8; ++m) {
+ const float gate = Rnd(SigmoidF32(up_v[m]));
+ sum[m] = fmaf(hn_v[m], gate, sum[m]);
+ }
+#pragma unroll
+ for (int p = 0; p < 4; ++p) {
+ float wf[8];
+ Unpack8(__ldg(w_ptr + p), wf);
+ const float xv0 = hn_v[2 * p];
+ const float xv1 = hn_v[2 * p + 1];
+ acc[0] = fmaf(xv0, wf[0], acc[0]);
+ acc[1] = fmaf(xv0, wf[1], acc[1]);
+ acc[2] = fmaf(xv0, wf[2], acc[2]);
+ acc[3] = fmaf(xv0, wf[3], acc[3]);
+ acc[0] = fmaf(xv1, wf[4], acc[0]);
+ acc[1] = fmaf(xv1, wf[5], acc[1]);
+ acc[2] = fmaf(xv1, wf[6], acc[2]);
+ acc[3] = fmaf(xv1, wf[7], acc[3]);
+ }
+ }
+
+ float out[8];
+#pragma unroll
+ for (int m = 0; m < 8; ++m) {
+ out[m] = sum[m] * reduction_scale;
+ }
+ *reinterpret_cast(pre_out + row * kHidden + h) = Pack8(out);
+
+ __shared__ float smem[10][4];
+#pragma unroll
+ for (int b = 0; b < 4; ++b) {
+#pragma unroll
+ for (int off = 16; off > 0; off >>= 1) acc[b] += __shfl_down_sync(0xffffffffu, acc[b], off);
+ if ((tid & 31) == 0) smem[tid >> 5][b] = acc[b];
+ }
+ __syncthreads();
+ if (tid < 4) {
+ float v = 0.0f;
+#pragma unroll
+ for (int w = 0; w < 10; ++w) v += smem[w][tid];
+ const float t = Rnd(Rnd(v) / gate_div);
+ const float u = Rnd(SigmoidF32(t));
+ gate_out[row * 4 + tid] = from_float(u * gate_mul);
+ }
+}
+
template
__global__ void FusedInjectPostGateB4Kernel(const T* __restrict__ hn, const T* __restrict__ w_inject,
T* __restrict__ gate_out, int64_t rows, int64_t K,
@@ -976,8 +1141,10 @@ bool IsProgrammaticLaunchAvailable() {
};
available = available && image_has_griddepcontrol(HcDownKernel) &&
image_has_griddepcontrol(HcDownKernel) &&
- image_has_griddepcontrol(HcUpPreKernel) &&
image_has_griddepcontrol(HcUpPreKernel);
+ if constexpr (UP_ROWS <= 2) {
+ available = available && image_has_griddepcontrol(HcUpPreKernel);
+ }
(void)cudaGetLastError();
cached_device = device;
cached_available = available;
@@ -1002,9 +1169,11 @@ Status LaunchFusedRows(bool use_pdl, cudaStream_t stream, const HyperConnectionL
(void)cudaFuncSetAttribute(HcUpPreKernel,
cudaFuncAttributePreferredSharedMemoryCarveout,
cudaSharedmemCarveoutMaxShared);
- (void)cudaFuncSetAttribute(HcUpPreKernel,
- cudaFuncAttributePreferredSharedMemoryCarveout,
- cudaSharedmemCarveoutMaxShared);
+ if constexpr (UP_ROWS <= 2) {
+ (void)cudaFuncSetAttribute(HcUpPreKernel,
+ cudaFuncAttributePreferredSharedMemoryCarveout,
+ cudaSharedmemCarveoutMaxShared);
+ }
(void)cudaGetLastError();
return true;
}();
@@ -1027,18 +1196,26 @@ Status LaunchFusedRows(bool use_pdl, cudaStream_t stream, const HyperConnectionL
}
const T* up_streams = tensors.prev_block_output != nullptr ? tensors.streams_out : tensors.streams;
- if (config.branches == 4 && config.rank == 320) {
+ if constexpr (UP_ROWS > 2) {
return LaunchWithProgrammaticSerialization(use_pdl, HcUpPreKernel, up_grid, dim3(kThreads),
stream, up_streams, tensors.norm_scale, tensors.up_weight,
static_cast(partial),
static_cast(inv_rms), tensors.pre_output,
tensors.post_gate, up);
+ } else {
+ if (config.branches == 4 && config.rank == 320) {
+ return LaunchWithProgrammaticSerialization(use_pdl, HcUpPreKernel, up_grid, dim3(kThreads),
+ stream, up_streams, tensors.norm_scale, tensors.up_weight,
+ static_cast(partial),
+ static_cast(inv_rms), tensors.pre_output,
+ tensors.post_gate, up);
+ }
+ return LaunchWithProgrammaticSerialization(use_pdl, HcUpPreKernel, up_grid, dim3(kThreads), stream,
+ up_streams, tensors.norm_scale, tensors.up_weight,
+ static_cast(partial),
+ static_cast(inv_rms), tensors.pre_output,
+ tensors.post_gate, up);
}
- return LaunchWithProgrammaticSerialization(use_pdl, HcUpPreKernel, up_grid, dim3(kThreads), stream,
- up_streams, tensors.norm_scale, tensors.up_weight,
- static_cast(partial),
- static_cast(inv_rms), tensors.pre_output,
- tensors.post_gate, up);
}
} // namespace
@@ -1096,7 +1273,7 @@ Status LaunchHyperConnectionLowRankMixFused(const cudaDeviceProp& device_prop, c
if (config.rows == 1) {
return LaunchFusedRows(use_pdl, stream, config, tensors, g, partial, inv_rms, down, up);
}
- // Fused row-group dispatch: rows == 1 -> 1x1, rows 2..7 -> 2x2, rows == 8 -> 4x2 (down x up rows per CTA).
+ // Fused row-group dispatch: rows == 1 -> 1x1, rows 2..7 -> 2x2, rows == 8 -> 4x2.
if (config.rows >= 8) {
return LaunchFusedRows(use_pdl, stream, config, tensors, g, partial, inv_rms, down, up);
}
@@ -1170,6 +1347,19 @@ Status LaunchHyperConnectionLowRankMixUnfused(cublasHandle_t cublas, cudaStream_
ORT_RETURN_IF_ERROR(CUDA_CALL(cudaGetLastError()));
ORT_RETURN_IF_ERROR(Gemm(cublas, rows, K, config.rank, down, tensors.up_weight, up));
+ if constexpr (sizeof(T) == 2) {
+ if (config.branches == 4 && config.hidden == 2560 && rows <= std::numeric_limits::max() &&
+ (reinterpret_cast(hn) & 15u) == 0 &&
+ (reinterpret_cast(up) & 15u) == 0 &&
+ (reinterpret_cast(tensors.pre_output) & 15u) == 0 &&
+ (reinterpret_cast(tensors.inject_weight) & 15u) == 0) {
+ FusedPreMixInjectPostGateB4H2560Kernel<<(rows), 320, 0, stream>>>(
+ hn, up, tensors.inject_weight, tensors.pre_output, tensors.post_gate, rows,
+ config.reduction_scale, config.gate_div, config.gate_mul);
+ return CUDA_CALL(cudaGetLastError());
+ }
+ }
+
const int64_t pre_count = rows * config.hidden;
if constexpr (sizeof(T) == 2) {
const int64_t vec_count = pre_count / 8;
diff --git a/onnxruntime/contrib_ops/cuda/cuda_contrib_kernels.cc b/onnxruntime/contrib_ops/cuda/cuda_contrib_kernels.cc
index fc4c1633baac4..da2c42ae9458d 100644
--- a/onnxruntime/contrib_ops/cuda/cuda_contrib_kernels.cc
+++ b/onnxruntime/contrib_ops/cuda/cuda_contrib_kernels.cc
@@ -161,6 +161,9 @@ class CUDA_ONNX_OP_TYPED_CLASS_NAME(1, MLFloat16, ParametricSoftplus);
class CUDA_MS_OP_TYPED_CLASS_NAME(1, float, RotaryEmbedding);
class CUDA_MS_OP_TYPED_CLASS_NAME(1, MLFloat16, RotaryEmbedding);
class CUDA_MS_OP_TYPED_CLASS_NAME(1, BFloat16, RotaryEmbedding);
+class CUDA_MS_OP_TYPED_CLASS_NAME(1, float, SwiGLU);
+class CUDA_MS_OP_TYPED_CLASS_NAME(1, MLFloat16, SwiGLU);
+class CUDA_MS_OP_TYPED_CLASS_NAME(1, BFloat16, SwiGLU);
class CUDA_MS_OP_TYPED_CLASS_NAME(1, float, MRotaryEmbedding);
class CUDA_MS_OP_TYPED_CLASS_NAME(1, MLFloat16, MRotaryEmbedding);
class CUDA_MS_OP_TYPED_CLASS_NAME(1, BFloat16, MRotaryEmbedding);
@@ -531,6 +534,9 @@ Status RegisterCudaContribKernels(KernelRegistry& kernel_registry) {
BuildKernelCreateInfo,
BuildKernelCreateInfo,
BuildKernelCreateInfo,
+ BuildKernelCreateInfo,
+ BuildKernelCreateInfo,
+ BuildKernelCreateInfo,
BuildKernelCreateInfo,
BuildKernelCreateInfo,
BuildKernelCreateInfo,
diff --git a/onnxruntime/contrib_ops/cuda/math/swiglu.cc b/onnxruntime/contrib_ops/cuda/math/swiglu.cc
new file mode 100644
index 0000000000000..abf1ac358c4ca
--- /dev/null
+++ b/onnxruntime/contrib_ops/cuda/math/swiglu.cc
@@ -0,0 +1,77 @@
+// Copyright (c) Microsoft Corporation. All rights reserved.
+// Licensed under the MIT License.
+
+#include "contrib_ops/cuda/math/swiglu.h"
+
+#include "contrib_ops/cuda/math/swiglu_impl.h"
+#include "core/providers/cuda/cuda_common.h"
+
+using namespace onnxruntime::cuda;
+
+namespace onnxruntime {
+namespace contrib {
+namespace cuda {
+
+#define REGISTER_KERNEL_TYPED(T) \
+ ONNX_OPERATOR_TYPED_KERNEL_EX( \
+ SwiGLU, \
+ kMSDomain, \
+ 1, \
+ T, \
+ kCudaExecutionProvider, \
+ (*KernelDefBuilder::Create()) \
+ .TypeConstraint("T", DataTypeImpl::GetTensorType()), \
+ SwiGLU);
+
+REGISTER_KERNEL_TYPED(float)
+REGISTER_KERNEL_TYPED(MLFloat16)
+REGISTER_KERNEL_TYPED(BFloat16)
+
+template
+SwiGLU::SwiGLU(const OpKernelInfo& info) : CudaKernel(info) {
+ limit_ = info.GetAttrOrDefault("limit", 0.0f);
+ alpha_ = info.GetAttrOrDefault("activation_alpha", 1.0f);
+ beta_ = info.GetAttrOrDefault("activation_beta", 0.0f);
+ round_activation_ = info.GetAttrOrDefault("round_activation", 0) != 0;
+}
+
+template
+Status SwiGLU::ComputeInternal(OpKernelContext* context) const {
+ const Tensor* gate = context->Input(0);
+ const Tensor* up = context->Input(1);
+
+ if (up != nullptr) {
+ if (gate->Shape() != up->Shape()) {
+ return ORT_MAKE_STATUS(ONNXRUNTIME, INVALID_ARGUMENT, "gate is ", gate->Shape(),
+ " but up is ", up->Shape(), "; they must match.");
+ }
+ Tensor* output = context->Output(0, gate->Shape());
+ const int64_t count = gate->Shape().Size();
+ if (count == 0) return Status::OK();
+ return LaunchSwiGLU(Stream(context), count, 0, limit_, alpha_, beta_, round_activation_, gate->Data(),
+ up->Data(), output->MutableData());
+ }
+
+ // One projection produced both halves; slice it here rather than paying for a Split node.
+ const TensorShape& in_shape = gate->Shape();
+ const size_t rank = in_shape.NumDimensions();
+ if (rank == 0 || in_shape[rank - 1] % 2 != 0) {
+ return ORT_MAKE_STATUS(ONNXRUNTIME, INVALID_ARGUMENT, "with `up` omitted, gate is ", in_shape,
+ "; it must have a last dimension, and that dimension must be even.");
+ }
+ const int64_t cols = in_shape[rank - 1] / 2;
+
+ TensorShapeVector dims = in_shape.AsShapeVector();
+ dims[rank - 1] = cols;
+ Tensor* output = context->Output(0, TensorShape(dims));
+ const int64_t count = output->Shape().Size();
+ if (count == 0) return Status::OK();
+
+ const T* data = gate->Data();
+ return LaunchSwiGLU(Stream(context), count, cols, limit_, alpha_, beta_, round_activation_, data, data + cols,
+ output->MutableData());
+}
+
+} // namespace cuda
+} // namespace contrib
+} // namespace onnxruntime
diff --git a/onnxruntime/contrib_ops/cuda/math/swiglu.h b/onnxruntime/contrib_ops/cuda/math/swiglu.h
new file mode 100644
index 0000000000000..b048516555331
--- /dev/null
+++ b/onnxruntime/contrib_ops/cuda/math/swiglu.h
@@ -0,0 +1,31 @@
+// Copyright (c) Microsoft Corporation. All rights reserved.
+// Licensed under the MIT License.
+
+#pragma once
+
+#include "core/providers/cuda/cuda_kernel.h"
+
+namespace onnxruntime {
+namespace contrib {
+namespace cuda {
+
+using namespace onnxruntime::cuda;
+
+// Clamped SwiGLU on a gate/up pair given either as two tensors or as the two halves of one
+// `[.., 2 * inter]` projection. Same alpha/beta formula as MoE/QMoE, but non-positive limit disables clamping.
+template
+class SwiGLU final : public CudaKernel {
+ public:
+ explicit SwiGLU(const OpKernelInfo& info);
+ Status ComputeInternal(OpKernelContext* context) const override;
+
+ private:
+ float limit_;
+ float alpha_;
+ float beta_;
+ bool round_activation_;
+};
+
+} // namespace cuda
+} // namespace contrib
+} // namespace onnxruntime
diff --git a/onnxruntime/contrib_ops/cuda/math/swiglu_activation.cuh b/onnxruntime/contrib_ops/cuda/math/swiglu_activation.cuh
new file mode 100644
index 0000000000000..4c24f5e1dffde
--- /dev/null
+++ b/onnxruntime/contrib_ops/cuda/math/swiglu_activation.cuh
@@ -0,0 +1,55 @@
+// Copyright (c) Microsoft Corporation. All rights reserved.
+// Licensed under the MIT License.
+
+#pragma once
+
+#include
+
+namespace onnxruntime {
+namespace contrib {
+namespace cuda {
+
+// The clamped SwiGLU of one element, shared by every kernel that fuses it:
+// G = min(gate, limit), L = clamp(up, -limit, limit) (clamps off when limit <= 0)
+// result = G * sigmoid(alpha * G) * (L + beta)
+// Evaluated in float, widened to double for the product when alpha * G is below -87 so that a negative tail is not
+// lost to float exp underflow. With `round_activation`, G * sigmoid(alpha * G) is narrowed to CudaT before the
+// multiply by (L + beta), exactly as a QuickGelu node followed by a Mul in CudaT would.
+// Keep the arithmetic identical to the SwiGLU operator.
+template
+__device__ __forceinline__ float SwiGluActivation(float gate, float up, float limit, float alpha, float beta,
+ bool round_activation = false) {
+ // `l + beta` with beta == +0.0 turns l == -0.0 into +0.0 and can flip the sign of a zero result, which a
+ // QuickGelu followed by a Mul keeps. -0.0 is the additive identity for every l, so a zero beta adds that.
+ const float bias = beta == 0.0f ? -0.0f : beta;
+ float g = gate;
+ float l = up;
+ if (limit > 0.0f) {
+ // Unlike fmin and fmax, these comparisons preserve NaNs.
+ if (g > limit) g = limit;
+ if (l < -limit) l = -limit;
+ if (l > limit) l = limit;
+ }
+ const float scaled = alpha * g;
+ if (scaled < -87.0f) {
+ // Float exp enters its subnormal range near -87 and underflows to zero by -104, which would zero a tail that
+ // the following multiply by (l + beta) can bring back into range. Keep the whole product wide.
+ const double wide_g = static_cast(g);
+ const double e = exp(wide_g * static_cast(alpha));
+ const double activated = wide_g * e / (1.0 + e);
+ if (round_activation) {
+ // QuickGelu narrows double -> float -> CudaT.
+ return static_cast(static_cast(static_cast(activated))) * (l + bias);
+ }
+ return static_cast(activated * (static_cast(l) + static_cast(bias)));
+ }
+ // Keeps the exponent non-positive without cancelling small negative-input results.
+ const float sigmoid = scaled > 0.0f ? 1.0f / (1.0f + expf(-scaled)) : expf(scaled) / (1.0f + expf(scaled));
+ float activated = g * sigmoid;
+ if (round_activation) activated = static_cast(static_cast(activated));
+ return activated * (l + bias);
+}
+
+} // namespace cuda
+} // namespace contrib
+} // namespace onnxruntime
diff --git a/onnxruntime/contrib_ops/cuda/math/swiglu_impl.cu b/onnxruntime/contrib_ops/cuda/math/swiglu_impl.cu
new file mode 100644
index 0000000000000..4cc7148373ef3
--- /dev/null
+++ b/onnxruntime/contrib_ops/cuda/math/swiglu_impl.cu
@@ -0,0 +1,75 @@
+// Copyright (c) Microsoft Corporation. All rights reserved.
+// Licensed under the MIT License.
+
+#include "contrib_ops/cuda/math/swiglu_impl.h"
+
+#include "contrib_ops/cuda/math/swiglu_activation.cuh"
+#include "core/providers/cuda/cu_inc/common.cuh"
+
+using namespace onnxruntime::cuda;
+
+namespace onnxruntime {
+namespace contrib {
+namespace cuda {
+
+namespace {
+
+constexpr int kThreads = 256;
+
+// The gated activation of a feed-forward or MoE expert block. The graph spells this as two
+// casts, two clips, a sigmoid, two multiplies and a cast back -- eight passes over the same
+// buffer. Matches MoE/QMoE's alpha/beta formula and positive-limit clamping rule.
+// With `Fused`, gate and up are the two halves of one `[.., 2 * cols]` row, so the source
+// index skips the half that the other operand owns.
+// With `round_activation`, G * Sigmoid(alpha * G) is narrowed to T before the multiply by (L + beta),
+// exactly as a QuickGelu node followed by a Mul in T would, so SwiGluFusion does not change results.
+template
+__global__ void SwiGLUKernel(int64_t count, int64_t cols, float limit, float alpha, float beta,
+ bool round_activation, const CudaT* __restrict__ gate,
+ const CudaT* __restrict__ up, CudaT* __restrict__ output) {
+ const int64_t stride = static_cast(gridDim.x) * blockDim.x;
+ for (int64_t i = static_cast(blockIdx.x) * blockDim.x + threadIdx.x; i < count;
+ i += stride) {
+ const int64_t src = Fused ? i + (i / cols) * cols : i;
+ const float g = static_cast(gate[src]);
+ const float l = static_cast(up[src]);
+ output[i] = static_cast(SwiGluActivation(g, l, limit, alpha, beta, round_activation));
+ }
+}
+
+} // namespace
+
+template
+Status LaunchSwiGLU(cudaStream_t stream, int64_t count, int64_t cols, float limit,
+ float alpha, float beta, bool round_activation,
+ const T* gate, const T* up, T* output) {
+ typedef typename ToCudaType::MappedType CudaT;
+ const int64_t want = (count + kThreads - 1) / kThreads;
+ const int blocks = static_cast(want < 65535 ? want : 65535);
+ auto* g = reinterpret_cast(gate);
+ auto* u = reinterpret_cast(up);
+ auto* o = reinterpret_cast(output);
+ if (cols > 0) {
+ SwiGLUKernel<<>>(count, cols, limit, alpha, beta,
+ round_activation, g, u, o);
+ } else {
+ SwiGLUKernel<<>>(count, 0, limit, alpha, beta,
+ round_activation, g, u, o);
+ }
+ CUDA_RETURN_IF_ERROR(cudaGetLastError());
+ return Status::OK();
+}
+
+#define INSTANTIATE(T) \
+ template Status LaunchSwiGLU(cudaStream_t, int64_t, int64_t, float, float, float, bool, \
+ const T*, const T*, T*)
+
+INSTANTIATE(float);
+INSTANTIATE(MLFloat16);
+INSTANTIATE(BFloat16);
+
+#undef INSTANTIATE
+
+} // namespace cuda
+} // namespace contrib
+} // namespace onnxruntime
diff --git a/onnxruntime/contrib_ops/cuda/math/swiglu_impl.h b/onnxruntime/contrib_ops/cuda/math/swiglu_impl.h
new file mode 100644
index 0000000000000..0275ed08edf12
--- /dev/null
+++ b/onnxruntime/contrib_ops/cuda/math/swiglu_impl.h
@@ -0,0 +1,22 @@
+// Copyright (c) Microsoft Corporation. All rights reserved.
+// Licensed under the MIT License.
+
+#pragma once
+
+#include "core/providers/cuda/cuda_common.h"
+
+namespace onnxruntime {
+namespace contrib {
+namespace cuda {
+
+// `cols` is the width of one half when gate and up are the two halves of a single
+// `[.., 2 * cols]` row; pass 0 when they are separate, densely packed tensors.
+// `round_activation` narrows gate * Sigmoid(alpha * gate) to T before it is multiplied by the up half.
+template
+Status LaunchSwiGLU(cudaStream_t stream, int64_t count, int64_t cols, float limit,
+ float alpha, float beta, bool round_activation,
+ const T* gate, const T* up, T* output);
+
+} // namespace cuda
+} // namespace contrib
+} // namespace onnxruntime
diff --git a/onnxruntime/core/graph/contrib_ops/contrib_defs.cc b/onnxruntime/core/graph/contrib_ops/contrib_defs.cc
index fe19adb003419..75c8a286a3cd1 100644
--- a/onnxruntime/core/graph/contrib_ops/contrib_defs.cc
+++ b/onnxruntime/core/graph/contrib_ops/contrib_defs.cc
@@ -874,6 +874,67 @@ ONNX_MS_OPERATOR_SET_SCHEMA(BiasSoftmax, 1,
"Constrain input and output types to float tensors.")
.TypeAndShapeInferenceFunction(ONNX_NAMESPACE::propagateShapeAndTypeFromFirstInput));
+ONNX_MS_OPERATOR_SET_SCHEMA(
+ SwiGLU, 1,
+ OpSchema()
+ .SetDoc(
+ "Clamped SwiGLU, the gated activation of a feed-forward or MoE expert block:\n"
+ " G = clamp(gate, max=limit)\n"
+ " L = clamp(up, min=-limit, max=limit)\n"
+ " output = G * Sigmoid(activation_alpha * G) * (L + activation_beta)\n"
+ "with the arithmetic done in float regardless of T, widened to double for the product "
+ "when `activation_alpha * G` is below -87 so that a negative tail is not lost to "
+ "float exp underflow. A `limit` of zero or less disables both clamps.\n"
+ "\n"
+ "The result is rounded to T once, after the multiply by (L + activation_beta). With "
+ "`round_activation` set, G * Sigmoid(activation_alpha * G) is first rounded to T, as a "
+ "QuickGelu node followed by a Mul in T would, which keeps a fused QuickGelu + Mul "
+ "bit-identical to the graph it replaces.\n"
+ "\n"
+ "This uses the same activation_alpha / activation_beta formula and positive-limit "
+ "clamping rule as MoE and QMoE. The no-clamp sentinel differs: here any `limit` "
+ "of zero or less disables clamping, while MoE/QMoE use infinity when `swiglu_limit` "
+ "is omitted and clamp finite values, including an explicit zero.\n"
+ "With `up` supplied and the default `limit` and `activation_beta`, it "
+ "equals the ONNX SwiGLU operator (opset 28) with `alpha` = `activation_alpha`.\n"
+ "\n"
+ "`up` is optional. When it is omitted, `gate` carries both halves of one "
+ "`[.., 2 * inter]` projection -- gate first, then up -- and the split is done "
+ "internally, so a fused sibling GEMM needs no `Split` node.")
+ .Attr("limit", "Clamp applied to both halves; zero or less disables it.",
+ AttributeProto::FLOAT, 0.0f)
+ .Attr("activation_alpha", "Multiplier inside the sigmoid.", AttributeProto::FLOAT, 1.0f)
+ .Attr("activation_beta",
+ "Value added to the linear half. A zero value leaves it unchanged, including a negative zero.",
+ AttributeProto::FLOAT, 0.0f)
+ .Attr("round_activation",
+ "If non-zero, round G * Sigmoid(activation_alpha * G) to T before multiplying by "
+ "(L + activation_beta). Zero rounds only the final result.",
+ AttributeProto::INT, static_cast(0))
+ .Input(0, "gate", "Gate half of the projection, or both halves concatenated.", "T")
+ .Input(1, "up", "Up half of the projection, same shape as gate.", "T", OpSchema::Optional)
+ .Output(0, "output", "Activated tensor, shaped like one half of the projection.", "T")
+ .TypeConstraint("T", {"tensor(float16)", "tensor(bfloat16)", "tensor(float)"},
+ "Constrain the activation type to float tensors.")
+ .TypeAndShapeInferenceFunction([](ONNX_NAMESPACE::InferenceContext& ctx) {
+ propagateElemTypeFromInputToOutput(ctx, 0, 0);
+ if (!hasInputShape(ctx, 0)) return;
+ if (ctx.getNumInputs() > 1 && ctx.getInputType(1) != nullptr) {
+ propagateShapeFromInputToOutput(ctx, 0, 0);
+ return;
+ }
+ const auto& in = getInputShape(ctx, 0);
+ if (in.dim_size() == 0) fail_shape_inference("gate must not be a scalar.");
+ auto* out = ctx.getOutputType(0)->mutable_tensor_type()->mutable_shape();
+ for (int i = 0; i + 1 < in.dim_size(); ++i) *out->add_dim() = in.dim(i);
+ const auto& last = in.dim(in.dim_size() - 1);
+ auto* d = out->add_dim();
+ if (last.has_dim_value()) {
+ if (last.dim_value() % 2 != 0) fail_shape_inference("gate's last dim must be even.");
+ d->set_dim_value(last.dim_value() / 2);
+ }
+ }));
+
ONNX_MS_OPERATOR_SET_SCHEMA(BiasDropout, 1,
OpSchema()
.SetDoc(
diff --git a/onnxruntime/core/graph/contrib_ops/ms_opset.h b/onnxruntime/core/graph/contrib_ops/ms_opset.h
index 5b24454c085cf..7bfbc49c762c0 100644
--- a/onnxruntime/core/graph/contrib_ops/ms_opset.h
+++ b/onnxruntime/core/graph/contrib_ops/ms_opset.h
@@ -122,6 +122,7 @@ class ONNX_OPERATOR_SET_SCHEMA_CLASS_NAME(Microsoft, 1, MRotaryEmbedding);
class ONNX_OPERATOR_SET_SCHEMA_CLASS_NAME(Microsoft, 1, GemmaRotaryEmbedding);
class ONNX_OPERATOR_SET_SCHEMA_CLASS_NAME(Microsoft, 1, SampleOp);
class ONNX_OPERATOR_SET_SCHEMA_CLASS_NAME(Microsoft, 1, Sampling);
+class ONNX_OPERATOR_SET_SCHEMA_CLASS_NAME(Microsoft, 1, SwiGLU);
class ONNX_OPERATOR_SET_SCHEMA_CLASS_NAME(Microsoft, 1, SkipGroupNorm);
class ONNX_OPERATOR_SET_SCHEMA_CLASS_NAME(Microsoft, 1, SkipLayerNormalization);
class ONNX_OPERATOR_SET_SCHEMA_CLASS_NAME(Microsoft, 1, SkipSimplifiedLayerNormalization);
@@ -262,6 +263,7 @@ class OpSet_Microsoft_ver1 {
fn(GetOpSchema());
fn(GetOpSchema());
fn(GetOpSchema());
+ fn(GetOpSchema());
fn(GetOpSchema());
fn(GetOpSchema());
fn(GetOpSchema());
diff --git a/onnxruntime/core/optimizer/graph_transformer_utils.cc b/onnxruntime/core/optimizer/graph_transformer_utils.cc
index cc2c49dbc49e3..64372c52e1a20 100644
--- a/onnxruntime/core/optimizer/graph_transformer_utils.cc
+++ b/onnxruntime/core/optimizer/graph_transformer_utils.cc
@@ -88,6 +88,7 @@
#include "core/optimizer/skip_layer_norm_fusion.h"
#include "core/optimizer/slice_elimination.h"
#include "core/optimizer/slice_concat_to_space_to_depth_fusion.h"
+#include "core/optimizer/swiglu_fusion.h"
#include "core/optimizer/transpose_optimizer.h"
#include "core/optimizer/unsqueeze_elimination.h"
#ifdef ENABLE_TRAINING
@@ -428,6 +429,7 @@ InlinedVector> GenerateTransformers(
// Expose pre-opset-20 tanh GELU to ConvActivationFusion on WebGPU.
transformers.emplace_back(std::make_unique(cpu_cuda_dml_webgpu_eps));
transformers.emplace_back(std::make_unique(cpu_acl_cuda_dml_js_webgpu_eps));
+ transformers.emplace_back(std::make_unique(cuda_eps));
// GeluApproximation has side effects which may change results. It needs to be manually enabled,
// or alternatively the model can be updated offline using a model conversion script
diff --git a/onnxruntime/core/optimizer/pattern_rewriter.cc b/onnxruntime/core/optimizer/pattern_rewriter.cc
index 23339047a6bd2..9e73722bb2031 100644
--- a/onnxruntime/core/optimizer/pattern_rewriter.cc
+++ b/onnxruntime/core/optimizer/pattern_rewriter.cc
@@ -2598,8 +2598,22 @@ PatternRewriter::PatternRewriter(Rules rules, Options options)
}
Status PatternRewriter::Run(Graph& graph, Statistics& statistics) const {
+ return RunImpl(graph, statistics, nullptr);
+}
+
+Status PatternRewriter::Run(Graph& graph, Statistics& statistics, InlinedHashSet& region) const {
+ return RunImpl(graph, statistics, ®ion);
+}
+
+Status PatternRewriter::RunImpl(Graph& graph, Statistics& statistics, InlinedHashSet* region) const {
statistics = {};
ORT_RETURN_IF(graph.GraphResolveNeeded(), "PatternRewriter requires an already resolved graph.");
+ if (region) {
+ for (auto index : *region)
+ ORT_RETURN_IF(index >= static_cast(graph.MaxNodeIndex()) || graph.GetNode(index) == nullptr,
+ "Rewrite region contains an invalid node index: ", index);
+ if (region->empty()) return Status::OK();
+ }
auto report = [&](size_t id, NodeIndex root, FailureStage stage, std::string_view reason) {
if (options_.on_failure)
options_.on_failure({rules_[id].name, root, stage, std::string(reason)});
@@ -2634,8 +2648,10 @@ Status PatternRewriter::Run(Graph& graph, Statistics& statistics) const {
GraphViewer viewer(graph);
for (NodeIndex index : viewer.GetNodesInTopologicalOrder()) {
const Node& node = *graph.GetNode(index);
- worklist.push_back(index);
- queued.insert(index);
+ if (!region || region->count(index)) {
+ worklist.push_back(index);
+ queued.insert(index);
+ }
for (const NodeArg* output : node.OutputDefs()) {
if (output->Exists()) producers[output] = index;
}
@@ -2663,6 +2679,13 @@ Status PatternRewriter::Run(Graph& graph, Statistics& statistics) const {
++statistics.structural_checks;
candidate.matched = MatchPattern(pattern, graph, *root, producers, candidate.match,
candidate.stage, candidate.reason);
+ if (candidate.matched && region &&
+ std::any_of(candidate.match.nodes.begin(), candidate.match.nodes.end(),
+ [&](const Node* node) { return !region->count(node->Index()); })) {
+ candidate.matched = false;
+ candidate.stage = FailureStage::Boundary;
+ candidate.reason = "Match includes nodes outside the rewrite region.";
+ }
if (candidate.matched && !CanRewritePartition(candidate.match, options_)) {
candidate.matched = false;
candidate.stage = FailureStage::ExecutionProvider;
@@ -2739,11 +2762,18 @@ Status PatternRewriter::Run(Graph& graph, Statistics& statistics) const {
ORT_RETURN_IF(!declared(node.domain), "Replacement domain has no rule opset requirement: ", node.domain);
return Replace(graph, *winning_pattern, winning_match, recipe, producers, added);
};
+ InlinedVector removed;
+ if (region)
+ for (const Node* node : winning_match.nodes) removed.push_back(node->Index());
auto status = replace();
if (!status.IsOK()) {
report(winner, root->Index(), FailureStage::Replacement, status.ErrorMessage());
return status;
}
+ if (region) {
+ for (auto index : removed) region->erase(index);
+ region->insert(added.begin(), added.end());
+ }
++statistics.rewrites;
affected.insert(affected.end(), added.begin(), added.end());
// A changed producer or use count can affect a deeper pattern root, not just a direct consumer.
@@ -2755,7 +2785,7 @@ Status PatternRewriter::Run(Graph& graph, Statistics& statistics) const {
if (!seen.insert(index).second) continue;
const Node* node = graph.GetNode(index);
if (!node) continue;
- if (queued.insert(index).second) next_worklist.push_back(index);
+ if ((!region || region->count(index)) && queued.insert(index).second) next_worklist.push_back(index);
for (auto it = node->OutputEdgesBegin(); it != node->OutputEdgesEnd(); ++it) {
next.push_back(it->GetNode().Index());
}
diff --git a/onnxruntime/core/optimizer/pattern_rewriter.h b/onnxruntime/core/optimizer/pattern_rewriter.h
index 51d8d6e1daae0..ac0d896905601 100644
--- a/onnxruntime/core/optimizer/pattern_rewriter.h
+++ b/onnxruntime/core/optimizer/pattern_rewriter.h
@@ -335,7 +335,7 @@ enum class FailureStage { Opset,
TypeShape,
Constant,
Attribute,
- Boundary,
+ Boundary, // Unsafe deletion or a match containing nodes outside the region.
ExecutionProvider,
CheckAndBind,
Replacement };
@@ -365,12 +365,19 @@ class PatternRewriter final : public GraphTransformer {
// Requires an already resolved graph. Resolve again after mutation before general graph use.
// A later failure does not roll back earlier rewrites; inspect the returned status before Resolve/use.
Status Run(Graph& graph, Statistics& statistics) const;
+ // Only matches wholly inside region are eligible. Successful rewrites replace
+ // their source IDs with the new IDs in region, including across iterations.
+ // Boundary values are preserved; their edges to outside nodes may be rewired.
+ // IDs must initially be live in graph. Empty regions do no work. As with Run,
+ // failures do not roll back earlier rewrites; region tracks completed rewrites.
+ Status Run(Graph& graph, Statistics& statistics, InlinedHashSet& region) const;
private:
struct Group {
Pattern pattern;
InlinedVector rules;
};
+ Status RunImpl(Graph& graph, Statistics& statistics, InlinedHashSet* region) const;
Status ApplyImpl(Graph& graph, bool& modified, int graph_level, const logging::Logger& logger) const override;
Rules rules_;
InlinedVector groups_;
diff --git a/onnxruntime/core/optimizer/qdq_transformer/selectors_actions/qdq_selectors.cc b/onnxruntime/core/optimizer/qdq_transformer/selectors_actions/qdq_selectors.cc
index e374b5ec6563d..96d4530cfdfb8 100644
--- a/onnxruntime/core/optimizer/qdq_transformer/selectors_actions/qdq_selectors.cc
+++ b/onnxruntime/core/optimizer/qdq_transformer/selectors_actions/qdq_selectors.cc
@@ -46,6 +46,42 @@ constexpr bool IsNBitsIntType(int32_t data_type) {
return Is2BitIntType(data_type) || Is4BitIntType(data_type) || Is8BitIntType(data_type);
}
+constexpr bool IsFloatQuantType(int32_t data_type) {
+ return (data_type == ONNX_NAMESPACE::TensorProto_DataType::TensorProto_DataType_FLOAT8E4M3FN) ||
+ (data_type == ONNX_NAMESPACE::TensorProto_DataType::TensorProto_DataType_FLOAT8E4M3FNUZ) ||
+ (data_type == ONNX_NAMESPACE::TensorProto_DataType::TensorProto_DataType_FLOAT8E5M2) ||
+ (data_type == ONNX_NAMESPACE::TensorProto_DataType::TensorProto_DataType_FLOAT8E5M2FNUZ) ||
+ (data_type == ONNX_NAMESPACE::TensorProto_DataType::TensorProto_DataType_FLOAT8E8M0) ||
+ (data_type == ONNX_NAMESPACE::TensorProto_DataType::TensorProto_DataType_FLOAT4E2M1);
+}
+
+// The QLinear*/QGemm replacement ops only accept integer quantized types, and dropping the Q/DQ nodes feeds the
+// quantized tensor to the target op directly, which most ops do not accept for float8/float4. Either rewrite would
+// produce an invalid graph, so node groups with float quantized types are left unchanged.
+bool HasFloatQuantType(const GraphViewer& graph_viewer, const NodeGroup& node_group) {
+ const auto is_float_quant_arg = [](const NodeArg* arg) {
+ const auto* type_proto = arg != nullptr ? arg->TypeAsProto() : nullptr;
+ return type_proto != nullptr && type_proto->has_tensor_type() &&
+ IsFloatQuantType(type_proto->tensor_type().elem_type());
+ };
+
+ for (NodeIndex dq_idx : node_group.dq_nodes) {
+ const Node* dq_node = graph_viewer.GetNode(dq_idx);
+ if (dq_node != nullptr && is_float_quant_arg(dq_node->InputDefs()[0])) {
+ return true;
+ }
+ }
+
+ for (NodeIndex q_idx : node_group.q_nodes) {
+ const Node* q_node = graph_viewer.GetNode(q_idx);
+ if (q_node != nullptr && is_float_quant_arg(q_node->OutputDefs()[0])) {
+ return true;
+ }
+ }
+
+ return false;
+}
+
// adjust for an optional input/output that has an entry but does not exist
int NumActualValues(const Node& node, bool input) {
const auto& defs = input ? node.InputDefs() : node.OutputDefs();
@@ -188,6 +224,10 @@ std::optional BaseSelector::Select(const GraphViewer& gr
return std::nullopt;
}
+ if (HasFloatQuantType(graph_viewer, *qdq_group)) {
+ return std::nullopt;
+ }
+
NodesToOptimizeIndicesBuilder builder;
// TODO(edgchen1) update NodeGroup to use InlinedVector
builder.input_nodes.assign(qdq_group->dq_nodes.begin(), qdq_group->dq_nodes.end());
diff --git a/onnxruntime/core/optimizer/swiglu_fusion.cc b/onnxruntime/core/optimizer/swiglu_fusion.cc
new file mode 100644
index 0000000000000..0bb842a13502e
--- /dev/null
+++ b/onnxruntime/core/optimizer/swiglu_fusion.cc
@@ -0,0 +1,110 @@
+// Copyright (c) Microsoft Corporation. All rights reserved.
+// Licensed under the MIT License.
+
+#include "core/optimizer/swiglu_fusion.h"
+
+#include
+
+#include "core/framework/tensorprotoutils.h"
+#include "core/graph/graph_utils.h"
+#include "core/optimizer/utils.h"
+
+namespace onnxruntime {
+
+namespace {
+
+// QuickGelu's schema default; QuickGeluFusion always writes the attribute explicitly.
+constexpr float kQuickGeluDefaultAlpha = 1.702f;
+
+// Mul broadcasts but SwiGLU does not, so both operands must provably share one shape.
+bool HaveSameShape(const NodeArg& a, const NodeArg& b) {
+ const auto* shape_a = a.Shape();
+ const auto* shape_b = b.Shape();
+ if (shape_a == nullptr || shape_b == nullptr || shape_a->dim_size() == 0 ||
+ shape_a->dim_size() != shape_b->dim_size()) {
+ return false;
+ }
+
+ for (int i = 0; i < shape_a->dim_size(); ++i) {
+ const auto& dim_a = shape_a->dim(i);
+ const auto& dim_b = shape_b->dim(i);
+ if (utils::HasDimValue(dim_a) && utils::HasDimValue(dim_b)) {
+ if (dim_a.dim_value() != dim_b.dim_value()) return false;
+ } else if (utils::HasDimParam(dim_a) && utils::HasDimParam(dim_b)) {
+ if (dim_a.dim_param() != dim_b.dim_param()) return false;
+ } else {
+ return false;
+ }
+ }
+ return true;
+}
+
+bool IsSupportedType(const NodeArg& arg) {
+ const auto* type = arg.Type();
+ return type != nullptr &&
+ (*type == "tensor(float)" || *type == "tensor(float16)" || *type == "tensor(bfloat16)");
+}
+
+} // namespace
+
+Status SwiGluFusion::ApplyImpl(Graph& graph, bool& modified, int graph_level, const logging::Logger& logger) const {
+ GraphViewer graph_viewer(graph);
+ const auto& node_topology_list = graph_viewer.GetNodesInTopologicalOrder();
+ for (auto node_index : node_topology_list) {
+ auto* p_node = graph.GetNode(node_index);
+ if (p_node == nullptr) continue;
+
+ Node& mul_node = *p_node;
+ ORT_RETURN_IF_ERROR(Recurse(mul_node, modified, graph_level, logger));
+
+ if (!graph_utils::IsSupportedOptypeVersionAndDomain(mul_node, "Mul", {7, 13, 14}) ||
+ !graph_utils::IsSupportedProvider(mul_node, GetCompatibleExecutionProviders())) {
+ continue;
+ }
+
+ int gate_index = -1;
+ for (int i = 0; i < 2; ++i) {
+ const Node* input_node = graph_utils::GetInputNode(mul_node, i);
+ if (input_node != nullptr &&
+ graph_utils::IsSupportedOptypeVersionAndDomain(*input_node, "QuickGelu", {1}, kMSDomain) &&
+ input_node->GetExecutionProviderType() == mul_node.GetExecutionProviderType() &&
+ optimizer_utils::CheckOutputEdges(graph, *input_node, 1)) {
+ gate_index = i;
+ break;
+ }
+ }
+ if (gate_index < 0) continue;
+
+ const int up_index = 1 - gate_index;
+ Node& quick_gelu_node = *graph.GetNode(graph_utils::GetInputNode(mul_node, gate_index)->Index());
+ NodeArg* gate_arg = quick_gelu_node.MutableInputDefs()[0];
+ NodeArg* up_arg = mul_node.MutableInputDefs()[up_index];
+ if (!IsSupportedType(*gate_arg) || !HaveSameShape(*gate_arg, *up_arg)) continue;
+
+ const auto* alpha_attr = graph_utils::GetNodeAttribute(quick_gelu_node, "alpha");
+ const float alpha = alpha_attr != nullptr ? alpha_attr->f() : kQuickGeluDefaultAlpha;
+
+ // FinalizeNodeFusion only moves input edges of the first fused node, so `up` is re-wired below.
+ const Node::EdgeEnd* up_edge = graph_utils::GetInputEdge(mul_node, up_index);
+ const NodeIndex up_src_node = up_edge != nullptr ? up_edge->GetNode().Index() : 0;
+ const int up_src_arg_index = up_edge != nullptr ? up_edge->GetSrcArgIndex() : -1;
+
+ Node& swiglu_node = graph.AddNode(graph.GenerateNodeName(mul_node.Name() + "/SwiGluFusion/"), "SwiGLU",
+ "fused QuickGelu and Mul", std::array{gate_arg, up_arg},
+ std::array{mul_node.MutableOutputDefs()[0]}, nullptr, kMSDomain);
+ swiglu_node.AddAttribute("activation_alpha", alpha);
+ // QuickGelu rounds to T before the Mul; keep that rounding so the fusion does not change results.
+ swiglu_node.AddAttribute("round_activation", static_cast(1));
+ swiglu_node.SetExecutionProviderType(mul_node.GetExecutionProviderType());
+
+ graph_utils::FinalizeNodeFusion(graph, {quick_gelu_node, mul_node}, swiglu_node);
+ if (up_src_arg_index >= 0) {
+ graph.AddEdge(up_src_node, swiglu_node.Index(), up_src_arg_index, 1);
+ }
+ modified = true;
+ }
+
+ return Status::OK();
+}
+
+} // namespace onnxruntime
diff --git a/onnxruntime/core/optimizer/swiglu_fusion.h b/onnxruntime/core/optimizer/swiglu_fusion.h
new file mode 100644
index 0000000000000..6ab794942c282
--- /dev/null
+++ b/onnxruntime/core/optimizer/swiglu_fusion.h
@@ -0,0 +1,27 @@
+// Copyright (c) Microsoft Corporation. All rights reserved.
+// Licensed under the MIT License.
+
+#pragma once
+
+#include "core/optimizer/graph_transformer.h"
+
+namespace onnxruntime {
+
+/**
+ * @brief Rewrite QuickGelu(gate) * up to com.microsoft SwiGLU(gate, up).
+ *
+ * QuickGeluFusion first turns gate * Sigmoid(alpha * gate) into QuickGelu, so running after it
+ * covers the decomposed SiLU gated MLP that exporters emit today.
+ *
+ * The fused node sets round_activation so QuickGelu's output is still rounded to the tensor type
+ * before it meets `up`, which keeps the result identical to the unfused pair.
+ */
+class SwiGluFusion : public GraphTransformer {
+ public:
+ explicit SwiGluFusion(const InlinedHashSet& compatible_execution_providers = {}) noexcept
+ : GraphTransformer("SwiGluFusion", compatible_execution_providers) {}
+
+ Status ApplyImpl(Graph& graph, bool& modified, int graph_level, const logging::Logger& logger) const override;
+};
+
+} // namespace onnxruntime
diff --git a/onnxruntime/core/providers/cuda/cu_inc/common.cuh b/onnxruntime/core/providers/cuda/cu_inc/common.cuh
index 93257da22e269..595ce5d5ae4d2 100644
--- a/onnxruntime/core/providers/cuda/cu_inc/common.cuh
+++ b/onnxruntime/core/providers/cuda/cu_inc/common.cuh
@@ -311,6 +311,13 @@ __device__ __inline__ double _Log(double a) { return log(a); }
template <>
__device__ __inline__ half _Log(half a) { return half(logf((float)a)); }
+// Keeps the exponent non-positive without cancelling small negative-input results.
+__device__ __inline__ float _Sigmoid(float a) {
+ if (a > 0.0f) return 1.0f / (1.0f + expf(-a));
+ const float exp_a = expf(a);
+ return exp_a / (1.0f + exp_a);
+}
+
template
__device__ __inline T _Tanh(T a);
diff --git a/onnxruntime/core/providers/cuda/cuda_execution_provider.cc b/onnxruntime/core/providers/cuda/cuda_execution_provider.cc
index 1c1ca9088b7a5..e434df5b1456d 100755
--- a/onnxruntime/core/providers/cuda/cuda_execution_provider.cc
+++ b/onnxruntime/core/providers/cuda/cuda_execution_provider.cc
@@ -721,6 +721,14 @@ class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kO
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 7, bool, And);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 7, bool, Or);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 7, bool, Xor);
+class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 18, uint8_t, BitwiseXor);
+class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 18, uint16_t, BitwiseXor);
+class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 18, uint32_t, BitwiseXor);
+class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 18, uint64_t, BitwiseXor);
+class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 18, int8_t, BitwiseXor);
+class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 18, int16_t, BitwiseXor);
+class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 18, int32_t, BitwiseXor);
+class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 18, int64_t, BitwiseXor);
class ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 6, 7, Sum);
class ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 8, 12, Sum);
class ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 6, 11, Max);
@@ -1384,6 +1392,18 @@ class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain,
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 13, BFloat16, ReduceSum);
class ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 13, 27, Mod);
class ONNX_OPERATOR_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 28, Mod);
+class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, 27, uint8_t, BitShift);
+class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, 27, uint16_t, BitShift);
+class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, 27, uint32_t, BitShift);
+class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, 27, uint64_t, BitShift);
+class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 28, uint8_t, BitShift);
+class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 28, uint16_t, BitShift);
+class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 28, uint32_t, BitShift);
+class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 28, uint64_t, BitShift);
+class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 28, int8_t, BitShift);
+class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 28, int16_t, BitShift);
+class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 28, int32_t, BitShift);
+class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 28, int64_t, BitShift);
class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 13, 18, int8_t, QuantizeLinear);
class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 13, 18, uint8_t, QuantizeLinear);
class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 13, 18, int8_t, DequantizeLinear);
@@ -2052,6 +2072,14 @@ static Status RegisterCudaKernels(KernelRegistry& kernel_registry) {
BuildKernelCreateInfo,
BuildKernelCreateInfo,
BuildKernelCreateInfo,
+ BuildKernelCreateInfo,
+ BuildKernelCreateInfo,
+ BuildKernelCreateInfo,
+ BuildKernelCreateInfo,
+ BuildKernelCreateInfo,
+ BuildKernelCreateInfo,
+ BuildKernelCreateInfo,
+ BuildKernelCreateInfo,
BuildKernelCreateInfo,
BuildKernelCreateInfo,
BuildKernelCreateInfo,
@@ -2716,6 +2744,18 @@ static Status RegisterCudaKernels(KernelRegistry& kernel_registry) {
BuildKernelCreateInfo,
BuildKernelCreateInfo,
BuildKernelCreateInfo,
+ BuildKernelCreateInfo,
+ BuildKernelCreateInfo,
+ BuildKernelCreateInfo,
+ BuildKernelCreateInfo,
+ BuildKernelCreateInfo,
+ BuildKernelCreateInfo,
+ BuildKernelCreateInfo,
+ BuildKernelCreateInfo,
+ BuildKernelCreateInfo,
+ BuildKernelCreateInfo,
+ BuildKernelCreateInfo,
+ BuildKernelCreateInfo,
BuildKernelCreateInfo,
BuildKernelCreateInfo,
BuildKernelCreateInfo,
diff --git a/onnxruntime/core/providers/cuda/math/binary_elementwise_ops.cc b/onnxruntime/core/providers/cuda/math/binary_elementwise_ops.cc
index 0340f3b4c7e8b..04f1d5ad42a38 100644
--- a/onnxruntime/core/providers/cuda/math/binary_elementwise_ops.cc
+++ b/onnxruntime/core/providers/cuda/math/binary_elementwise_ops.cc
@@ -292,6 +292,14 @@ BINARY_OP_REGISTER_VERSIONED_CLASS_HFD(Pow, Pow_7, 7, 11)
BINARY_LOGICALOP_TYPED(And, 7, bool)
BINARY_LOGICALOP_TYPED(Or, 7, bool)
BINARY_LOGICALOP_TYPED(Xor, 7, bool)
+BINARY_OP_TYPED(BitwiseXor, 18, uint8_t)
+BINARY_OP_TYPED(BitwiseXor, 18, uint16_t)
+BINARY_OP_TYPED(BitwiseXor, 18, uint32_t)
+BINARY_OP_TYPED(BitwiseXor, 18, uint64_t)
+BINARY_OP_TYPED(BitwiseXor, 18, int8_t)
+BINARY_OP_TYPED(BitwiseXor, 18, int16_t)
+BINARY_OP_TYPED(BitwiseXor, 18, int32_t)
+BINARY_OP_TYPED(BitwiseXor, 18, int64_t)
BINARY_OP_VERSIONED_HFD(PRelu, 7, 8)
BINARY_OP_VERSIONED_HFD(PRelu, 9, 15)
// Opset-16 adds BFloat16 to allowed types for the PRelu operator
diff --git a/onnxruntime/core/providers/cuda/math/binary_elementwise_ops.h b/onnxruntime/core/providers/cuda/math/binary_elementwise_ops.h
index 337c7909f46e9..18020cd5ed43d 100644
--- a/onnxruntime/core/providers/cuda/math/binary_elementwise_ops.h
+++ b/onnxruntime/core/providers/cuda/math/binary_elementwise_ops.h
@@ -200,6 +200,13 @@ class Xor final : public BinaryElementwise {
Status ComputeInternal(OpKernelContext* context) const override;
};
+template
+class BitwiseXor final : public BinaryElementwise {
+ public:
+ BitwiseXor(const OpKernelInfo& info) : BinaryElementwise(info) {}
+ Status ComputeInternal(OpKernelContext* context) const override;
+};
+
// PRelu is activation function, but it's closer to binary elementwise ops in implementation
template
class PRelu final : public BinaryElementwise {
diff --git a/onnxruntime/core/providers/cuda/math/binary_elementwise_ops_impl.cu b/onnxruntime/core/providers/cuda/math/binary_elementwise_ops_impl.cu
index 2b08614b1e9c4..283f459002b83 100644
--- a/onnxruntime/core/providers/cuda/math/binary_elementwise_ops_impl.cu
+++ b/onnxruntime/core/providers/cuda/math/binary_elementwise_ops_impl.cu
@@ -159,6 +159,14 @@ SPECIALIZED_BINARY_ELEMENTWISE_IMPL_HFD(Pow_7)
SPECIALIZED_BINARY_ELEMENTWISE_IMPL(And, bool)
SPECIALIZED_BINARY_ELEMENTWISE_IMPL(Or, bool)
SPECIALIZED_BINARY_ELEMENTWISE_IMPL(Xor, bool)
+SPECIALIZED_BINARY_ELEMENTWISE_IMPL(BitwiseXor, uint8_t)
+SPECIALIZED_BINARY_ELEMENTWISE_IMPL(BitwiseXor, uint16_t)
+SPECIALIZED_BINARY_ELEMENTWISE_IMPL(BitwiseXor, uint32_t)
+SPECIALIZED_BINARY_ELEMENTWISE_IMPL(BitwiseXor, uint64_t)
+SPECIALIZED_BINARY_ELEMENTWISE_IMPL(BitwiseXor, int8_t)
+SPECIALIZED_BINARY_ELEMENTWISE_IMPL(BitwiseXor, int16_t)
+SPECIALIZED_BINARY_ELEMENTWISE_IMPL(BitwiseXor, int32_t)
+SPECIALIZED_BINARY_ELEMENTWISE_IMPL(BitwiseXor, int64_t)
SPECIALIZED_BINARY_ELEMENTWISE_IMPL_HFD(PRelu)
SPECIALIZED_BINARY_ELEMENTWISE_IMPL_UZILHFD(Max)
SPECIALIZED_BINARY_ELEMENTWISE_IMPL_UZILHFD(Min)
diff --git a/onnxruntime/core/providers/cuda/math/binary_elementwise_ops_impl.h b/onnxruntime/core/providers/cuda/math/binary_elementwise_ops_impl.h
index 5311f73f01a78..dd778ccb15be2 100644
--- a/onnxruntime/core/providers/cuda/math/binary_elementwise_ops_impl.h
+++ b/onnxruntime/core/providers/cuda/math/binary_elementwise_ops_impl.h
@@ -24,6 +24,7 @@ namespace cuda {
BINARY_OP_NAME_EXPR(And, (a & b)) \
BINARY_OP_NAME_EXPR(Or, (a | b)) \
BINARY_OP_NAME_EXPR(Xor, (a ^ b)) \
+ BINARY_OP_NAME_EXPR(BitwiseXor, (a ^ b)) \
BINARY_OP_NAME_EXPR(PRelu, (a > (T)0 ? a : a * b)) \
BINARY_OP_NAME_EXPR(Max, _Max(a, b)) \
BINARY_OP_NAME_EXPR(Min, _Min(a, b)) \
diff --git a/onnxruntime/core/providers/cuda/math/bitshift.cc b/onnxruntime/core/providers/cuda/math/bitshift.cc
new file mode 100644
index 0000000000000..86d6858f299c8
--- /dev/null
+++ b/onnxruntime/core/providers/cuda/math/bitshift.cc
@@ -0,0 +1,75 @@
+// Copyright (c) Microsoft Corporation. All rights reserved.
+// Licensed under the MIT License.
+
+#include "core/providers/shared_library/provider_api.h"
+#include "core/providers/cuda/math/binary_elementwise_ops.h"
+#include "core/providers/cuda/math/bitshift_impl.h"
+
+namespace onnxruntime {
+namespace cuda {
+
+template
+class BitShift final : public BinaryElementwise {
+ public:
+ explicit BitShift(const OpKernelInfo& info) : BinaryElementwise(info) {
+ std::string direction;
+ ORT_THROW_IF_ERROR(info.GetAttr("direction", &direction));
+ ORT_ENFORCE(direction == "LEFT" || direction == "RIGHT",
+ "Invalid direction value of '", direction, "'. Valid values are 'LEFT' or 'RIGHT'.");
+ shift_left_ = direction == "LEFT";
+ }
+
+ Status ComputeInternal(OpKernelContext* context) const override {
+ BinaryElementwisePreparation prepare;
+ ORT_RETURN_IF_ERROR(Prepare(context, &prepare));
+ Impl_BitShift(
+ Stream(context),
+ prepare.output_rank_or_simple_broadcast,
+ &prepare.lhs_padded_strides,
+ prepare.lhs_tensor->Data(),
+ &prepare.rhs_padded_strides,
+ prepare.rhs_tensor->Data(),
+ &prepare.fdm_output_strides,
+ prepare.fdm_H,
+ prepare.fdm_C,
+ prepare.output_tensor->MutableData(),
+ prepare.output_tensor->Shape().Size(),
+ shift_left_);
+ return Status::OK();
+ }
+
+ private:
+ bool shift_left_;
+};
+
+#define REGISTER_BITSHIFT_VERSIONED(T) \
+ ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_EX( \
+ BitShift, kOnnxDomain, 11, 27, T, kCudaExecutionProvider, \
+ (*KernelDefBuilder::Create()).TypeConstraint("T", DataTypeImpl::GetTensorType()), \
+ BitShift);
+
+#define REGISTER_BITSHIFT(T) \
+ ONNX_OPERATOR_TYPED_KERNEL_EX( \
+ BitShift, kOnnxDomain, 28, T, kCudaExecutionProvider, \
+ (*KernelDefBuilder::Create()).TypeConstraint("T", DataTypeImpl::GetTensorType()), \
+ BitShift);
+
+REGISTER_BITSHIFT_VERSIONED(uint8_t)
+REGISTER_BITSHIFT_VERSIONED(uint16_t)
+REGISTER_BITSHIFT_VERSIONED(uint32_t)
+REGISTER_BITSHIFT_VERSIONED(uint64_t)
+
+REGISTER_BITSHIFT(uint8_t)
+REGISTER_BITSHIFT(uint16_t)
+REGISTER_BITSHIFT(uint32_t)
+REGISTER_BITSHIFT(uint64_t)
+REGISTER_BITSHIFT(int8_t)
+REGISTER_BITSHIFT(int16_t)
+REGISTER_BITSHIFT(int32_t)
+REGISTER_BITSHIFT(int64_t)
+
+#undef REGISTER_BITSHIFT_VERSIONED
+#undef REGISTER_BITSHIFT
+
+} // namespace cuda
+} // namespace onnxruntime
diff --git a/onnxruntime/core/providers/cuda/math/bitshift_impl.cu b/onnxruntime/core/providers/cuda/math/bitshift_impl.cu
new file mode 100644
index 0000000000000..bd03b213b61e5
--- /dev/null
+++ b/onnxruntime/core/providers/cuda/math/bitshift_impl.cu
@@ -0,0 +1,96 @@
+// Copyright (c) Microsoft Corporation. All rights reserved.
+// Licensed under the MIT License.
+
+#include
+#include
+
+#include "core/providers/cuda/math/bitshift_impl.h"
+#include "core/providers/cuda/cu_inc/common.cuh"
+#include "core/providers/cuda/cu_inc/binary_elementwise_impl.cuh"
+
+namespace onnxruntime {
+namespace cuda {
+
+template
+struct BitShiftFunctor {
+ __device__ __inline__ T operator()(T value, T shift) const {
+ using UnsignedT = std::make_unsigned_t;
+ // Keep narrow integers unsigned after integer promotion as well.
+ using WorkT = std::conditional_t<(sizeof(T) < sizeof(uint32_t)), uint32_t, UnsignedT>;
+ constexpr unsigned bit_width = sizeof(T) * 8;
+ const auto amount = static_cast(shift);
+ if (amount >= bit_width) {
+ if constexpr (!ShiftLeft && std::is_signed_v) {
+ if (value < 0) {
+ return T{-1};
+ }
+ }
+ return T{0};
+ }
+
+ WorkT bits = static_cast(value);
+ if constexpr (ShiftLeft) {
+ bits <<= amount;
+ } else {
+ bits >>= amount;
+ if constexpr (std::is_signed_v) {
+ if (value < 0 && amount != 0) {
+ bits |= ~WorkT{0} << (bit_width - amount);
+ }
+ }
+ }
+
+ // Preserve the two's-complement result without an out-of-range signed conversion.
+ const auto result_bits = static_cast(bits);
+ T result;
+ memcpy(&result, &result_bits, sizeof(result));
+ return result;
+ }
+};
+
+template
+void Impl_BitShift(
+ cudaStream_t stream,
+ int32_t output_rank_or_simple_broadcast,
+ const TArray* lhs_padded_strides,
+ const T* lhs_data,
+ const TArray* rhs_padded_strides,
+ const T* rhs_data,
+ const TArray* fdm_output_strides,
+ const fast_divmod& fdm_H,
+ const fast_divmod& fdm_C,
+ T* output_data,
+ size_t count,
+ bool shift_left) {
+ if (shift_left) {
+ BinaryElementWiseImpl(stream, output_rank_or_simple_broadcast,
+ lhs_padded_strides, lhs_data, rhs_padded_strides, rhs_data,
+ fdm_output_strides, fdm_H, fdm_C, output_data,
+ BitShiftFunctor{}, count);
+ } else {
+ BinaryElementWiseImpl(stream, output_rank_or_simple_broadcast,
+ lhs_padded_strides, lhs_data, rhs_padded_strides, rhs_data,
+ fdm_output_strides, fdm_H, fdm_C, output_data,
+ BitShiftFunctor{}, count);
+ }
+}
+
+#define INSTANTIATE_BITSHIFT(T) \
+ template void Impl_BitShift( \
+ cudaStream_t, int32_t, const TArray*, const T*, \
+ const TArray*, const T*, const TArray*, \
+ const fast_divmod&, const fast_divmod&, T*, size_t, bool);
+
+INSTANTIATE_BITSHIFT(uint8_t)
+INSTANTIATE_BITSHIFT(uint16_t)
+INSTANTIATE_BITSHIFT(uint32_t)
+INSTANTIATE_BITSHIFT(uint64_t)
+INSTANTIATE_BITSHIFT(int8_t)
+INSTANTIATE_BITSHIFT(int16_t)
+INSTANTIATE_BITSHIFT(int32_t)
+INSTANTIATE_BITSHIFT(int64_t)
+
+#undef INSTANTIATE_BITSHIFT
+
+} // namespace cuda
+} // namespace onnxruntime
diff --git a/onnxruntime/core/providers/cuda/math/bitshift_impl.h b/onnxruntime/core/providers/cuda/math/bitshift_impl.h
new file mode 100644
index 0000000000000..994d7603a5e51
--- /dev/null
+++ b/onnxruntime/core/providers/cuda/math/bitshift_impl.h
@@ -0,0 +1,27 @@
+// Copyright (c) Microsoft Corporation. All rights reserved.
+// Licensed under the MIT License.
+
+#pragma once
+
+#include "core/providers/cuda/math/binary_elementwise_ops_impl.h"
+
+namespace onnxruntime {
+namespace cuda {
+
+template
+void Impl_BitShift(
+ cudaStream_t stream,
+ int32_t output_rank_or_simple_broadcast,
+ const TArray* lhs_padded_strides,
+ const T* lhs_data,
+ const TArray* rhs_padded_strides,
+ const T* rhs_data,
+ const TArray* fdm_output_strides,
+ const fast_divmod& fdm_H,
+ const fast_divmod& fdm_C,
+ T* output_data,
+ size_t count,
+ bool shift_left);
+
+} // namespace cuda
+} // namespace onnxruntime
diff --git a/onnxruntime/core/providers/webgpu/ep/ep.cc b/onnxruntime/core/providers/webgpu/ep/ep.cc
index de4bac1f91415..6397634704afb 100644
--- a/onnxruntime/core/providers/webgpu/ep/ep.cc
+++ b/onnxruntime/core/providers/webgpu/ep/ep.cc
@@ -12,6 +12,7 @@
#include "core/session/onnxruntime_run_options_config_keys.h"
#include "core/session/plugin_ep/ep_kernel_registration.h"
#include "core/providers/webgpu/webgpu_execution_provider.h"
+#include "core/providers/webgpu/quantization/quantize_linear.h"
#include "ep/get_capability_utils.h"
@@ -149,6 +150,21 @@ OrtStatus* ORT_API_CALL Ep::GetCapabilityImpl(OrtEp* this_ptr, const OrtGraph* g
}
}
+ if (node.GetOperatorType() == "QuantizeLinear" && node.GetDomain() == kOnnxDomain) {
+ bool precision_supported = true;
+ for (const auto& attr : node.GetAttributes()) {
+ if (attr.GetName() == "precision") {
+ int64_t precision = 0;
+ RETURN_IF_ERROR(attr.GetValue(precision));
+ precision_supported = IsQuantizeLinearPrecisionSupported(precision);
+ break;
+ }
+ }
+ if (!precision_supported) {
+ continue;
+ }
+ }
+
candidate_nodes.push_back(node);
tentative_candidate_nodes.push_back(node);
}
diff --git a/onnxruntime/core/providers/webgpu/math/binary_elementwise_ops.cc b/onnxruntime/core/providers/webgpu/math/binary_elementwise_ops.cc
index 8e0d20031571e..d2fdc042d449e 100644
--- a/onnxruntime/core/providers/webgpu/math/binary_elementwise_ops.cc
+++ b/onnxruntime/core/providers/webgpu/math/binary_elementwise_ops.cc
@@ -435,8 +435,16 @@ static std::string GetMinImpl(int lhs_element_type, int /* rhs_element_type */)
WEBGPU_VARIADIC_IMPL(Max, "max_v(vec4(a), vec4(b))", GetMaxImpl)
WEBGPU_VARIADIC_IMPL(Min, "min_v(vec4(a), vec4(b))", GetMinImpl)
-std::string GetPowImpl(int lhs_element_type, int /* rhs_element_type */) {
+std::string GetPowImpl(int lhs_element_type, int rhs_element_type) {
SS(s, 1024);
+ // Take an integer exponent's parity from the integer itself. Converting it to f32 first rounds odd
+ // values above 2^24 to even ones (16777217 -> 16777216), which flips the sign for a negative base.
+ // int64 exponents are read as their low 32 bits, which keep the parity.
+ const bool is_integer_exponent = rhs_element_type == ONNX_TENSOR_ELEMENT_DATA_TYPE_INT32 ||
+ rhs_element_type == ONNX_TENSOR_ELEMENT_DATA_TYPE_UINT32 ||
+ rhs_element_type == ONNX_TENSOR_ELEMENT_DATA_TYPE_INT64;
+ const char* exponent_is_odd = is_integer_exponent ? "(b & input_b_element_t(1)) != input_b_element_t(0)"
+ : "round(abs(f32(b)) % 2.0) == 1.0";
std::string round_str;
if (lhs_element_type == ONNX_TENSOR_ELEMENT_DATA_TYPE_INT32) {
round_str = "round";
@@ -453,17 +461,24 @@ std::string GetPowImpl(int lhs_element_type, int /* rhs_element_type */) {
" }\n";
}
- s << "fn pow_custom(a : input_a_element_t, b : f32) -> input_a_element_t {\n"
+ s << "fn pow_exponent_is_odd(b : input_b_element_t) -> bool {\n"
+ " return "
+ << exponent_is_odd << ";\n"
+ << "}\n"
+ "fn pow_custom(a : input_a_element_t, b : f32, b_is_odd : bool) -> input_a_element_t {\n"
" if (b == 0.0) {\n"
" return input_a_element_t(1.0);\n"
" } else if (a < input_a_element_t(0.0) && b != floor(b)) {\n"
" return input_a_element_t(pow(f32(a), b)); // NaN\n"
" }\n"
<< use_pow_shortcut
- << " return select(sign(a), input_a_element_t(1.0), round(abs(b) % 2.0) != 1.0) * input_a_element_t(" << round_str << "(pow(f32(abs(a)), b)));\n"
+ << " return select(input_a_element_t(1.0), sign(a), b_is_odd) * input_a_element_t(" << round_str << "(pow(f32(abs(a)), b)));\n"
<< "}\n"
"fn pow_v(a : vec4