From 548aa4fba61b2b81e62ec3e19dfa0ca862238d9a Mon Sep 17 00:00:00 2001 From: Jiajia Qin Date: Wed, 16 Sep 2026 11:18:56 +0800 Subject: [PATCH 01/25] docs: design WebGPU MatMul algorithm scheduler --- ...ebgpu-matmul-algorithm-scheduler-design.md | 63 +++++++++++++++++++ 1 file changed, 63 insertions(+) create mode 100644 docs/superpowers/specs/2026-09-16-webgpu-matmul-algorithm-scheduler-design.md diff --git a/docs/superpowers/specs/2026-09-16-webgpu-matmul-algorithm-scheduler-design.md b/docs/superpowers/specs/2026-09-16-webgpu-matmul-algorithm-scheduler-design.md new file mode 100644 index 0000000000000..b1b3778ebd2f7 --- /dev/null +++ b/docs/superpowers/specs/2026-09-16-webgpu-matmul-algorithm-scheduler-design.md @@ -0,0 +1,63 @@ +# WebGPU MatMul Algorithm Scheduler Design + +## Goal + +Make every `ComputeMatMul` implementation path explicit and independently testable without changing the default runtime behavior. Selection policy must be separated from execution, preserve the current rule order, and allow vendor-specific policy to extend the common rules. + +## Algorithms + +Introduce `MatMulAlgorithm` with five concrete values: + +- `SubgroupMatrix`: the common subgroup-matrix implementation in `subgroup_matrix_matmul.cc`. +- `Naive`: `MatMulNaiveProgram`. +- `IntelSubgroup`: the Intel subgroup implementation in `vendor/intel/math/matmul.cc`. +- `Packed`: the generic packed `MatMulProgram` without Split-K. +- `PackedSplitK`: the same packed program with Split-K initialization and atomic accumulation. + +There is no `Auto` algorithm value. Automatic versus forced selection is represented by an optional forced algorithm, which prevents an unset policy state from being confused with an executable implementation. + +## Selection Architecture + +Add a `MatMulAlgorithmScheduler` base class. Its common rule order mirrors the existing `ComputeMatMul` condition order: + +1. Select `SubgroupMatrix` when its implementation reports that it can handle the problem. +2. Select `Naive` when `N < 8 && K < 8`. +3. Ask a virtual vendor-policy hook for a vendor algorithm. +4. Select `PackedSplitK` when the existing `SplitKConfig::UseSplitK` rule succeeds. +5. Fall back to `Packed`. + +An Intel-derived scheduler implements the vendor hook and selects `IntelSubgroup` under the current Intel subgroup rule. Other vendors use the base scheduler unchanged. Future vendor policies can derive from the base scheduler without adding vendor conditionals to `ComputeMatMul`. + +The scheduler accepts already-computed selection facts rather than owning tensor execution. This keeps it deterministic and unit-testable without a WebGPU device. Vendor policy may inspect dimensions and device-derived capability facts, but it must return only a `MatMulAlgorithm`. + +## Forced Test Selection + +Add the internal WebGPU session configuration key `ep.webgpuexecutionprovider.forceMatmulAlgorithm`. Accepted values are `subgroup_matrix`, `naive`, `intel_subgroup`, `packed`, and `packed_split_k`. The option is parsed when the WebGPU EP is created, stored as `std::optional`, and exposed read-only through `ComputeContextBase`. + +When set, the scheduler returns the requested enum before applying heuristic rules. The dispatcher then validates the algorithm's hard prerequisites. Unsupported device features, data types, layouts, deterministic-compute settings, or other correctness constraints produce a descriptive failure naming the forced algorithm; forced mode never silently falls back. + +Heuristic thresholds are not hard prerequisites. For example, forcing Intel subgroup bypasses its current `M/N/K` performance thresholds while still requiring an Intel adapter with subgroup support. Forcing Split-K bypasses performance thresholds while still requiring a usable Split-K configuration, non-deterministic compute, compatible packing/activation, and supported bias layout. + +Invalid option strings fail during WebGPU provider creation and list accepted values. + +## Dispatch and Implementation Boundaries + +`ComputeMatMul` computes shared shape facts once, asks the scheduler for exactly one enum, and switches directly to the matching implementation. Algorithm bodies are extracted into focused helpers where necessary. The generic packed helper takes an explicit Split-K mode; it does not re-run the selection heuristic. + +The subgroup-matrix optional implementation gains a non-mutating applicability query and an execution method that no longer communicates selection through a `handled` output. This removes trial execution as a dispatch mechanism. + +The existing per-kernel cache continues to own device-dependent subgroup-matrix state and the scheduler, so MatMul, pointwise Conv, and Attention callers retain their current caching and behavior. + +## Compatibility + +With no forcing option, the selected algorithm and precedence remain equivalent to the current code. Existing call sites keep using `ComputeMatMul`; the refactor does not change the operator API or model semantics. + +The option is intentionally internal and test-only: it is declared with WebGPU provider options for configuration plumbing but is not added to public user documentation. + +## Testing + +- Add device-independent scheduler unit tests covering every common branch, precedence, Intel override, default fallback, and forced override. +- Add parser/configuration tests for every accepted value and invalid input. +- Add WebGPU MatMul tests that choose shapes which normally select a different path, force a compatible algorithm, and verify numerical output. Hardware-specific forced algorithms are tested only when their hard capabilities are present; strict-failure tests cover unsupported forced choices. +- Build the WebGPU provider and provider-test target on Windows. Run scheduler/parser tests locally. Run hardware-backed MatMul tests where the available adapter supports them; macOS-arm64 Metal CI remains the correctness source of truth for WebGPU MatMul execution because lavapipe cannot execute MatMul reliably. + From 21ecabe9e7a2673b06666eab060253f5f9b0a515 Mon Sep 17 00:00:00 2001 From: Jiajia Qin Date: Wed, 16 Sep 2026 11:21:32 +0800 Subject: [PATCH 02/25] docs: plan WebGPU MatMul scheduler refactor --- ...09-16-webgpu-matmul-algorithm-scheduler.md | 223 ++++++++++++++++++ ...ebgpu-matmul-algorithm-scheduler-design.md | 5 +- 2 files changed, 226 insertions(+), 2 deletions(-) create mode 100644 docs/superpowers/plans/2026-09-16-webgpu-matmul-algorithm-scheduler.md diff --git a/docs/superpowers/plans/2026-09-16-webgpu-matmul-algorithm-scheduler.md b/docs/superpowers/plans/2026-09-16-webgpu-matmul-algorithm-scheduler.md new file mode 100644 index 0000000000000..9e8be4e90db7a --- /dev/null +++ b/docs/superpowers/plans/2026-09-16-webgpu-matmul-algorithm-scheduler.md @@ -0,0 +1,223 @@ +# WebGPU MatMul Algorithm Scheduler Implementation Plan + +> **For agentic workers:** REQUIRED SUB-SKILL: Use superpowers:subagent-driven-development (recommended) or superpowers:executing-plans to implement this plan task-by-task. Steps use checkbox (`- [ ]`) syntax for tracking. + +**Goal:** Replace conditional trial dispatch in `ComputeMatMul` with an explicit, vendor-extensible algorithm scheduler and a test-only forced-algorithm session option. + +**Architecture:** A lightweight algorithm enum is shared by WebGPU provider configuration and MatMul. A pure scheduler preserves current precedence and delegates vendor rules to an Intel subclass; `ComputeMatMul` dispatches the chosen enum through one switch and validates hard prerequisites when a test forces a path. + +**Tech Stack:** C++17, ONNX Runtime WebGPU EP, GoogleTest, CMake/Visual Studio on Windows + +**Spec:** `docs/superpowers/specs/2026-09-16-webgpu-matmul-algorithm-scheduler-design.md` + +## Global Constraints + +- Automatic selection must preserve current behavior and rule precedence. +- Forced selection bypasses heuristics but never bypasses correctness or hardware prerequisites. +- Unsupported forced algorithms fail explicitly and never silently fall back. +- `PackedSplitK` is independently selectable from `Packed`. +- The option remains internal/test-only and is not added to public documentation. +- Existing MatMul, pointwise Conv, and Attention callers continue through `ComputeMatMul`. +- The Windows validation build enables Dawn Vulkan, disables Dawn D3D12, and tests explicitly select `dawnBackendType=Vulkan`. + +--- + +### Task 1: Algorithm Identity and Configuration Plumbing + +**Files:** +- Create: `onnxruntime/core/providers/webgpu/math/matmul_algorithm.h` +- Modify: `onnxruntime/core/providers/webgpu/webgpu_provider_options.h` +- Modify: `onnxruntime/core/providers/webgpu/webgpu_execution_provider.h` +- Modify: `onnxruntime/core/providers/webgpu/webgpu_execution_provider.cc` +- Modify: `onnxruntime/core/providers/webgpu/webgpu_provider_factory.cc` +- Modify: `onnxruntime/core/providers/webgpu/compute_context.h` +- Test: `onnxruntime/test/providers/webgpu/matmul_algorithm_scheduler_test.cc` + +**Interfaces:** +- Produces: `enum class MatMulAlgorithm { SubgroupMatrix, Naive, IntelSubgroup, Packed, PackedSplitK }`. +- Produces: `std::optional ParseMatMulAlgorithm(std::string_view)` and `std::string_view MatMulAlgorithmName(MatMulAlgorithm)`. +- Produces: `ComputeContextBase::ForcedMatMulAlgorithm() const` returning `std::optional`. + +- [ ] **Step 1: Write the failing enum parsing tests** + +```cpp +EXPECT_EQ(ParseMatMulAlgorithm("packed_split_k"), MatMulAlgorithm::PackedSplitK); +EXPECT_EQ(MatMulAlgorithmName(MatMulAlgorithm::PackedSplitK), "packed_split_k"); +EXPECT_EQ(ParseMatMulAlgorithm("unknown"), std::nullopt); +``` + +- [ ] **Step 2: Build to verify the new tests fail** + +```powershell +python tools/ci_build/build.py --config RelWithDebInfo --build_dir .\build\WGPU-Vulkan --use_webgpu --update --build --parallel --target onnxruntime_provider_test --cmake_extra_defines onnxruntime_ENABLE_DAWN_BACKEND_VULKAN=ON onnxruntime_ENABLE_DAWN_BACKEND_D3D12=OFF +``` + +Expected: compilation fails because the algorithm interface is missing. + +- [ ] **Step 3: Implement enum conversion and option plumbing** + +Use exact case-sensitive names `subgroup_matrix`, `naive`, `intel_subgroup`, `packed`, and `packed_split_k`. Declare `kForceMatMulAlgorithm = "ep.webgpuexecutionprovider.forceMatmulAlgorithm"`, parse it in `ParseEpConfig`, reject invalid values with the accepted-value list, store an `std::optional` in `WebGpuExecutionProvider`, and expose it through `ComputeContextBase`. + +- [ ] **Step 4: Build and run parser tests** + +Build with Step 2, locate `onnxruntime_provider_test.exe`, and run `--gtest_filter="MatMulAlgorithmParsingTest.*"` from its directory. Confirm a non-zero test count and zero failures. + +- [ ] **Step 5: Commit** + +```powershell +git add onnxruntime/core/providers/webgpu onnxruntime/test/providers/webgpu/matmul_algorithm_scheduler_test.cc +git commit -m "webgpu: add forced MatMul algorithm option" +``` + +### Task 2: Common and Intel Algorithm Schedulers + +**Files:** +- Modify: `onnxruntime/core/providers/webgpu/math/matmul.h` +- Modify: `onnxruntime/core/providers/webgpu/math/matmul.cc` +- Modify: `onnxruntime/core/providers/webgpu/vendor/intel/math/matmul.h` +- Modify: `onnxruntime/core/providers/webgpu/vendor/intel/math/matmul.cc` +- Test: `onnxruntime/test/providers/webgpu/matmul_algorithm_scheduler_test.cc` + +**Interfaces:** +- Consumes: `MatMulAlgorithm` and the optional forced value from Task 1. +- Produces: `MatMulAlgorithmSelectionParams` with common rule facts. +- Produces: `MatMulAlgorithmScheduler::Select(const MatMulAlgorithmSelectionParams&, std::optional) const`. +- Produces: an Intel scheduler overriding the vendor rule hook. + +- [ ] **Step 1: Write failing scheduler tests** + +Cover forced precedence, subgroup-matrix precedence, exact `N < 8 && K < 8` boundaries, Intel selection, Split-K selection, and packed fallback. + +```cpp +EXPECT_EQ(base.Select(small, std::nullopt), MatMulAlgorithm::Naive); +EXPECT_EQ(intel.Select(intel_problem, std::nullopt), MatMulAlgorithm::IntelSubgroup); +EXPECT_EQ(base.Select(split_k_problem, std::nullopt), MatMulAlgorithm::PackedSplitK); +EXPECT_EQ(base.Select({}, MatMulAlgorithm::Packed), MatMulAlgorithm::Packed); +``` + +- [ ] **Step 2: Build and confirm compilation fails for missing scheduler types** + +Run the Task 1 build command and retain the compiler failure as the red TDD result. + +- [ ] **Step 3: Implement the common scheduler** + +Implement a non-virtual `Select` that checks forced selection, subgroup matrix, naive, a protected virtual vendor hook, Split-K, and packed in that order. The base vendor hook returns `std::nullopt`. + +- [ ] **Step 4: Implement the Intel scheduler** + +The Intel override returns `IntelSubgroup` only for the supplied current-rule fact. Add a scheduler factory selected by `context.AdapterInfo().vendor`. + +- [ ] **Step 5: Build, run `--gtest_filter="MatMulAlgorithmSchedulerTest.*"`, and commit** + +```powershell +git add onnxruntime/core/providers/webgpu/math/matmul.h onnxruntime/core/providers/webgpu/math/matmul.cc onnxruntime/core/providers/webgpu/vendor/intel/math/matmul.h onnxruntime/core/providers/webgpu/vendor/intel/math/matmul.cc onnxruntime/test/providers/webgpu/matmul_algorithm_scheduler_test.cc +git commit -m "refactor: add WebGPU MatMul algorithm scheduler" +``` + +### Task 3: Direct Enum Dispatch + +**Files:** +- Modify: `onnxruntime/core/providers/webgpu/math/matmul.h` +- Modify: `onnxruntime/core/providers/webgpu/math/matmul.cc` +- Modify: `onnxruntime/core/providers/webgpu/math/subgroup_matrix_matmul.cc` +- Modify: `onnxruntime/core/providers/webgpu/vendor/intel/math/gemm_subgroup.h` +- Modify: `onnxruntime/core/providers/webgpu/vendor/intel/math/gemm_subgroup.cc` +- Modify: `onnxruntime/core/providers/webgpu/vendor/intel/math/matmul.h` +- Modify: `onnxruntime/core/providers/webgpu/vendor/intel/math/matmul.cc` +- Test: `onnxruntime/test/providers/webgpu/matmul_algorithm_scheduler_test.cc` + +**Interfaces:** +- Produces: `MatMulOptImpl::CanApply(...) const` and `Compute(...)` without a `handled` output. +- Produces: separate Intel hard-capability and automatic-heuristic checks. +- Produces: one `switch (algorithm)` in `ComputeMatMul` and one helper per implementation. + +- [ ] **Step 1: Write failing prerequisite tests** + +Cover forced Split-K rejection for deterministic compute and incompatible packing/activation/layout. Cover Intel hard capability separately from its `M >= 64 && N >= 512 && K >= 32` automatic heuristic. + +- [ ] **Step 2: Run focused tests and preserve the red result** + +Run `onnxruntime_provider_test.exe --gtest_filter="MatMulAlgorithm*"`; expect missing validation interfaces or failing assertions. + +- [ ] **Step 3: Split subgroup-matrix applicability from execution** + +Move non-mutating early-decline checks into `CanApply`, make `Compute` error if called inapplicably, remove `handled`, and preserve odd-N cached padding. + +- [ ] **Step 4: Separate Intel capability from heuristic policy** + +Hard capability is `vendor == intel && Subgroups`. Automatic policy adds the existing `M >= 64 && N >= 512 && K >= 32` thresholds. Forced mode uses only hard capability. + +- [ ] **Step 5: Extract naive and packed helpers** + +Make the packed helper accept explicit `bool use_split_k`. False never runs the Split-K heuristic; true validates a configured split size, non-deterministic compute, vec4 packing, no fused activation, and compatible bias layout before atomic accumulation. + +- [ ] **Step 6: Replace conditional execution with one switch** + +```cpp +switch (algorithm) { + case MatMulAlgorithm::SubgroupMatrix: return subgroup_impl->Compute(...); + case MatMulAlgorithm::Naive: return ApplyMatMulNaive(...); + case MatMulAlgorithm::IntelSubgroup: return intel::ApplyMatMulIntel(...); + case MatMulAlgorithm::Packed: return ApplyMatMulPacked(..., false); + case MatMulAlgorithm::PackedSplitK: return ApplyMatMulPacked(..., true); +} +``` + +Every forced-prerequisite error includes `MatMulAlgorithmName(algorithm)`. + +- [ ] **Step 7: Build, run `--gtest_filter="MatMulAlgorithm*"`, and commit** + +```powershell +git add onnxruntime/core/providers/webgpu/math onnxruntime/core/providers/webgpu/vendor/intel/math onnxruntime/test/providers/webgpu/matmul_algorithm_scheduler_test.cc +git commit -m "refactor: dispatch WebGPU MatMul by algorithm" +``` + +### Task 4: Forced-Path Integration and Verification + +**Files:** +- Modify: `onnxruntime/test/providers/webgpu/matmul_large_test.cc` +- Modify: `onnxruntime/test/providers/webgpu/matmul_algorithm_scheduler_test.cc` + +**Interfaces:** +- Consumes: `kForceMatMulAlgorithm` and `WebGpuExecutionProviderWithOptions`. +- Produces: end-to-end proof that forcing changes dispatch and preserves output. + +- [ ] **Step 1: Write failing integration tests** + +Create each EP with `kDawnBackendType=Vulkan`. Use a non-small shape forced to `Naive`, a small shape forced to `Packed`, a compatible large shape forced to `PackedSplitK`, and a compatible f16 shape forced to `SubgroupMatrix`. Add an expected-failure case with deliberately unmet hard prerequisites. + +- [ ] **Step 2: Run `--gtest_filter="WebGpuMatMulAlgorithmTest.*"` and preserve the failing result** + +Confirm a non-zero selected test count and at least one expected pre-implementation failure. + +- [ ] **Step 3: Complete only wiring or diagnostics exposed by the tests** + +Ensure test EP construction passes the option through `ConfigOptions`, the same path used by WebGPU session creation. + +- [ ] **Step 4: Run focused and regression tests** + +```powershell +.\onnxruntime_provider_test.exe --gtest_filter="MatMulAlgorithm*:WebGpuMatMulAlgorithmTest.*" +.\onnxruntime_provider_test.exe --gtest_filter="MathOpTest.MatMulFloatType:MathOpTest.MatMul_Float16:WebGpuMatMulLargeTest.*" +``` + +Confirm non-zero counts and zero failures. The subgroup-matrix case must verify that the Vulkan adapter advertises the required subgroup-matrix configuration and subgroup-size control; otherwise it may skip only with the exact missing capability in its reason. + +- [ ] **Step 5: Run formatting and diff checks** + +Run the repository C++ formatter/linter for changed files, then: + +```powershell +git diff --check main...HEAD +git status --short +git diff --stat main...HEAD +``` + +- [ ] **Step 6: Commit and report limitations** + +```powershell +git add onnxruntime/test/providers/webgpu/matmul_large_test.cc onnxruntime/test/providers/webgpu/matmul_algorithm_scheduler_test.cc +git commit -m "test: force WebGPU MatMul algorithms" +``` + +Report exact commands, counts, Vulkan adapter identity, and observed subgroup-matrix capabilities. State that macOS-arm64 Metal CI remains additional cross-backend validation and that lavapipe cannot execute this family reliably. diff --git a/docs/superpowers/specs/2026-09-16-webgpu-matmul-algorithm-scheduler-design.md b/docs/superpowers/specs/2026-09-16-webgpu-matmul-algorithm-scheduler-design.md index b1b3778ebd2f7..c70ce892c2970 100644 --- a/docs/superpowers/specs/2026-09-16-webgpu-matmul-algorithm-scheduler-design.md +++ b/docs/superpowers/specs/2026-09-16-webgpu-matmul-algorithm-scheduler-design.md @@ -59,5 +59,6 @@ The option is intentionally internal and test-only: it is declared with WebGPU p - Add device-independent scheduler unit tests covering every common branch, precedence, Intel override, default fallback, and forced override. - Add parser/configuration tests for every accepted value and invalid input. - Add WebGPU MatMul tests that choose shapes which normally select a different path, force a compatible algorithm, and verify numerical output. Hardware-specific forced algorithms are tested only when their hard capabilities are present; strict-failure tests cover unsupported forced choices. -- Build the WebGPU provider and provider-test target on Windows. Run scheduler/parser tests locally. Run hardware-backed MatMul tests where the available adapter supports them; macOS-arm64 Metal CI remains the correctness source of truth for WebGPU MatMul execution because lavapipe cannot execute MatMul reliably. - +- Build Dawn and the WebGPU provider on Windows with the Vulkan backend enabled and D3D12 disabled. Tests explicitly request `dawnBackendType=Vulkan`. +- Verify the selected Vulkan adapter exposes subgroup size control, f16, and the cooperative/subgroup-matrix configuration required by the 8x16x16 kernel before claiming subgroup-matrix execution coverage. +- Run scheduler/parser tests and hardware-backed MatMul tests on the local Intel Arc Vulkan adapter. macOS-arm64 Metal CI remains additional cross-backend coverage; lavapipe is not used because it cannot execute MatMul reliably. From dce8a2d98d039ebf6e13698e8cb286799692c242 Mon Sep 17 00:00:00 2001 From: Jiajia Qin Date: Wed, 16 Sep 2026 12:19:11 +0800 Subject: [PATCH 03/25] build: skip DXC for Vulkan-only WebGPU --- cmake/external/onnxruntime_external_deps.cmake | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/cmake/external/onnxruntime_external_deps.cmake b/cmake/external/onnxruntime_external_deps.cmake index 7e56d910c3fb2..cdd98b34c358c 100644 --- a/cmake/external/onnxruntime_external_deps.cmake +++ b/cmake/external/onnxruntime_external_deps.cmake @@ -751,10 +751,6 @@ if (onnxruntime_USE_WEBGPU) set(DAWN_ENABLE_SPIRV_VALIDATION OFF CACHE BOOL "" FORCE) if (WIN32) - # building this requires the HLSL writer to be enabled in Tint. TBD if that we need either of these to be ON. - set(DAWN_USE_BUILT_DXC ON CACHE BOOL "" FORCE) - set(TINT_BUILD_HLSL_WRITER ON CACHE BOOL "" FORCE) - if ((NOT onnxruntime_ENABLE_DAWN_BACKEND_VULKAN) AND (NOT onnxruntime_ENABLE_DAWN_BACKEND_D3D12)) message(FATAL_ERROR "At least one of onnxruntime_ENABLE_DAWN_BACKEND_VULKAN or onnxruntime_ENABLE_DAWN_BACKEND_D3D12 must be enabled when using Dawn on Windows.") endif() @@ -765,8 +761,12 @@ if (onnxruntime_USE_WEBGPU) set(DAWN_ENABLE_VULKAN OFF CACHE BOOL "" FORCE) endif() if (onnxruntime_ENABLE_DAWN_BACKEND_D3D12) + set(DAWN_USE_BUILT_DXC ON CACHE BOOL "" FORCE) + set(TINT_BUILD_HLSL_WRITER ON CACHE BOOL "" FORCE) set(DAWN_ENABLE_D3D12 ON CACHE BOOL "" FORCE) else() + set(DAWN_USE_BUILT_DXC OFF CACHE BOOL "" FORCE) + set(TINT_BUILD_HLSL_WRITER OFF CACHE BOOL "" FORCE) set(DAWN_ENABLE_D3D12 OFF CACHE BOOL "" FORCE) endif() # We are currently always using the D3D12 backend. From 05a892e1bbb43ac7658b532747ac04b18ae40d08 Mon Sep 17 00:00:00 2001 From: Jiajia Qin Date: Wed, 16 Sep 2026 12:19:11 +0800 Subject: [PATCH 04/25] webgpu: add forced MatMul algorithm option --- .../core/providers/webgpu/compute_context.h | 4 ++ .../providers/webgpu/math/matmul_algorithm.h | 56 +++++++++++++++++++ .../webgpu/webgpu_execution_provider.cc | 1 + .../webgpu/webgpu_execution_provider.h | 5 ++ .../webgpu/webgpu_provider_factory.cc | 8 +++ .../webgpu/webgpu_provider_options.h | 2 + .../webgpu/matmul_algorithm_scheduler_test.cc | 39 +++++++++++++ 7 files changed, 115 insertions(+) create mode 100644 onnxruntime/core/providers/webgpu/math/matmul_algorithm.h create mode 100644 onnxruntime/test/providers/webgpu/matmul_algorithm_scheduler_test.cc diff --git a/onnxruntime/core/providers/webgpu/compute_context.h b/onnxruntime/core/providers/webgpu/compute_context.h index 99b2ca5cf600b..0d41fc1da9a32 100644 --- a/onnxruntime/core/providers/webgpu/compute_context.h +++ b/onnxruntime/core/providers/webgpu/compute_context.h @@ -122,6 +122,10 @@ class ComputeContextBase { return ep_.EnableMatmulFp32Accumulation(); } + inline std::optional ForcedMatMulAlgorithm() const { + return ep_.ForcedMatMulAlgorithm(); + } + // // Get the logger. // diff --git a/onnxruntime/core/providers/webgpu/math/matmul_algorithm.h b/onnxruntime/core/providers/webgpu/math/matmul_algorithm.h new file mode 100644 index 0000000000000..9740bfd451f8d --- /dev/null +++ b/onnxruntime/core/providers/webgpu/math/matmul_algorithm.h @@ -0,0 +1,56 @@ +// Copyright (c) Microsoft Corporation. All rights reserved. +// Licensed under the MIT License. + +#pragma once + +#include +#include + +namespace onnxruntime { +namespace webgpu { + +enum class MatMulAlgorithm { + SubgroupMatrix, + Naive, + IntelSubgroup, + Packed, + PackedSplitK, +}; + +inline std::optional ParseMatMulAlgorithm(std::string_view name) { + if (name == "subgroup_matrix") { + return MatMulAlgorithm::SubgroupMatrix; + } + if (name == "naive") { + return MatMulAlgorithm::Naive; + } + if (name == "intel_subgroup") { + return MatMulAlgorithm::IntelSubgroup; + } + if (name == "packed") { + return MatMulAlgorithm::Packed; + } + if (name == "packed_split_k") { + return MatMulAlgorithm::PackedSplitK; + } + return std::nullopt; +} + +inline std::string_view MatMulAlgorithmName(MatMulAlgorithm algorithm) { + switch (algorithm) { + case MatMulAlgorithm::SubgroupMatrix: + return "subgroup_matrix"; + case MatMulAlgorithm::Naive: + return "naive"; + case MatMulAlgorithm::IntelSubgroup: + return "intel_subgroup"; + case MatMulAlgorithm::Packed: + return "packed"; + case MatMulAlgorithm::PackedSplitK: + return "packed_split_k"; + } + return "unknown"; +} + +} // namespace webgpu +} // namespace onnxruntime diff --git a/onnxruntime/core/providers/webgpu/webgpu_execution_provider.cc b/onnxruntime/core/providers/webgpu/webgpu_execution_provider.cc index 0db9ca708a3f4..5960714e1e828 100644 --- a/onnxruntime/core/providers/webgpu/webgpu_execution_provider.cc +++ b/onnxruntime/core/providers/webgpu/webgpu_execution_provider.cc @@ -605,6 +605,7 @@ WebGpuExecutionProvider::WebGpuExecutionProvider(int context_id, multi_rotary_cache_concat_offset_{config.multi_rotary_cache_concat_offset}, kv_cache_quantization_bits_{config.kv_cache_quantization_bits}, enable_matmul_fp32_accumulation_{config.enable_matmul_fp32_accumulation}, + forced_matmul_algorithm_{config.forced_matmul_algorithm}, prepack_allocator_{CreateWebGpuAllocator( /*device_free=*/!context.HasDevice(), [this]() -> const webgpu::BufferManager& { return context_.InitializerBufferManager(); }, false)} { diff --git a/onnxruntime/core/providers/webgpu/webgpu_execution_provider.h b/onnxruntime/core/providers/webgpu/webgpu_execution_provider.h index 73ae3562b5231..bd3a4e7cd44ed 100644 --- a/onnxruntime/core/providers/webgpu/webgpu_execution_provider.h +++ b/onnxruntime/core/providers/webgpu/webgpu_execution_provider.h @@ -6,6 +6,7 @@ #include #include +#include #include #include #include @@ -17,6 +18,7 @@ #include "core/graph/constants.h" #include "core/providers/providers.h" #include "core/providers/webgpu/buffer_manager.h" +#include "core/providers/webgpu/math/matmul_algorithm.h" #include "core/providers/webgpu/session_buffer_pool.h" #if defined(ENABLE_PIX_FOR_WEBGPU_EP) @@ -57,6 +59,7 @@ struct WebGpuExecutionProviderConfig { // This is the single line that decides the shipped default for the // "enableMatmulFp32Accumulation" provider option. bool enable_matmul_fp32_accumulation{false}; + std::optional forced_matmul_algorithm; std::vector force_cpu_node_names{}; }; @@ -122,6 +125,7 @@ class WebGpuExecutionProvider : public IExecutionProvider { uint32_t KvCacheQuantizationBits() const { return kv_cache_quantization_bits_; } bool KvCacheQuantizationEnabled() const { return kv_cache_quantization_bits_ != 0; } bool EnableMatmulFp32Accumulation() const { return enable_matmul_fp32_accumulation_; } + std::optional ForcedMatMulAlgorithm() const { return forced_matmul_algorithm_; } #if defined(ORT_USE_EP_API_ADAPTERS) inline onnxruntime::ep::adapter::Logger& GetEpLogger() const { @@ -148,6 +152,7 @@ class WebGpuExecutionProvider : public IExecutionProvider { uint32_t multi_rotary_cache_concat_offset_ = 0; uint32_t kv_cache_quantization_bits_ = 0; bool enable_matmul_fp32_accumulation_ = false; + std::optional forced_matmul_algorithm_; std::unordered_map graph_id_to_run_count_; // Required regular runs before graph capture for any necessary allocations. const int min_num_runs_before_graph_capture_ = 0; diff --git a/onnxruntime/core/providers/webgpu/webgpu_provider_factory.cc b/onnxruntime/core/providers/webgpu/webgpu_provider_factory.cc index 27a8dd6c1e424..b4d8811372d6a 100644 --- a/onnxruntime/core/providers/webgpu/webgpu_provider_factory.cc +++ b/onnxruntime/core/providers/webgpu/webgpu_provider_factory.cc @@ -124,6 +124,14 @@ WebGpuExecutionProviderConfig ParseEpConfig(const ConfigOptions& config_options) } } + if (std::string forced_matmul_algorithm_str; + config_options.TryGetConfigEntry(kForceMatMulAlgorithm, forced_matmul_algorithm_str)) { + webgpu_ep_config.forced_matmul_algorithm = ParseMatMulAlgorithm(forced_matmul_algorithm_str); + ORT_ENFORCE(webgpu_ep_config.forced_matmul_algorithm.has_value(), + "Invalid forced MatMul algorithm: ", forced_matmul_algorithm_str, + ". Must be one of: subgroup_matrix, naive, intel_subgroup, packed, packed_split_k."); + } + // parse force CPU node names // The force CPU node names are separated by EOL (\n or \r\n) in the config entry. // each line is a node name that will be forced to run on CPU. diff --git a/onnxruntime/core/providers/webgpu/webgpu_provider_options.h b/onnxruntime/core/providers/webgpu/webgpu_provider_options.h index ace6f4ce5177d..4ee0b534e935f 100644 --- a/onnxruntime/core/providers/webgpu/webgpu_provider_options.h +++ b/onnxruntime/core/providers/webgpu/webgpu_provider_options.h @@ -32,6 +32,8 @@ constexpr const char* kKvCacheQuantizationBits = "ep.webgpuexecutionprovider.kvC // Today this covers MatMulNBits and its fused variants; the unquantized MatMul family is planned // as follow-up work under the same option. constexpr const char* kEnableMatmulFp32Accumulation = "ep.webgpuexecutionprovider.enableMatmulFp32Accumulation"; +// Internal test option for selecting a concrete unquantized MatMul implementation. +constexpr const char* kForceMatMulAlgorithm = "ep.webgpuexecutionprovider.forceMatmulAlgorithm"; constexpr const char* kDawnProcTable = "ep.webgpuexecutionprovider.dawnProcTable"; diff --git a/onnxruntime/test/providers/webgpu/matmul_algorithm_scheduler_test.cc b/onnxruntime/test/providers/webgpu/matmul_algorithm_scheduler_test.cc new file mode 100644 index 0000000000000..c8c68821588b8 --- /dev/null +++ b/onnxruntime/test/providers/webgpu/matmul_algorithm_scheduler_test.cc @@ -0,0 +1,39 @@ +// Copyright (c) Microsoft Corporation. All rights reserved. +// Licensed under the MIT License. + +#include "gtest/gtest.h" + +#include "core/providers/webgpu/math/matmul_algorithm.h" + +namespace onnxruntime { +namespace webgpu { +namespace test { + +TEST(MatMulAlgorithmParsingTest, RoundTripsEveryAlgorithmName) { + struct TestCase { + std::string_view name; + MatMulAlgorithm algorithm; + }; + + constexpr TestCase test_cases[] = { + {"subgroup_matrix", MatMulAlgorithm::SubgroupMatrix}, + {"naive", MatMulAlgorithm::Naive}, + {"intel_subgroup", MatMulAlgorithm::IntelSubgroup}, + {"packed", MatMulAlgorithm::Packed}, + {"packed_split_k", MatMulAlgorithm::PackedSplitK}, + }; + + for (const auto& test_case : test_cases) { + SCOPED_TRACE(test_case.name); + EXPECT_EQ(ParseMatMulAlgorithm(test_case.name), test_case.algorithm); + EXPECT_EQ(MatMulAlgorithmName(test_case.algorithm), test_case.name); + } +} + +TEST(MatMulAlgorithmParsingTest, RejectsUnknownAlgorithmName) { + EXPECT_EQ(ParseMatMulAlgorithm("unknown"), std::nullopt); +} + +} // namespace test +} // namespace webgpu +} // namespace onnxruntime From 2cf47049348746dfc21730ea6af49adadd8fd1c5 Mon Sep 17 00:00:00 2001 From: Jiajia Qin Date: Wed, 16 Sep 2026 12:41:00 +0800 Subject: [PATCH 05/25] refactor: add WebGPU MatMul algorithm scheduler --- .../webgpu/math/matmul_algorithm_scheduler.h | 56 ++++++++++++++++ .../intel/math/matmul_algorithm_scheduler.h | 26 ++++++++ .../webgpu/matmul_algorithm_scheduler_test.cc | 64 +++++++++++++++++++ 3 files changed, 146 insertions(+) create mode 100644 onnxruntime/core/providers/webgpu/math/matmul_algorithm_scheduler.h create mode 100644 onnxruntime/core/providers/webgpu/vendor/intel/math/matmul_algorithm_scheduler.h diff --git a/onnxruntime/core/providers/webgpu/math/matmul_algorithm_scheduler.h b/onnxruntime/core/providers/webgpu/math/matmul_algorithm_scheduler.h new file mode 100644 index 0000000000000..1c134d7553d50 --- /dev/null +++ b/onnxruntime/core/providers/webgpu/math/matmul_algorithm_scheduler.h @@ -0,0 +1,56 @@ +// Copyright (c) Microsoft Corporation. All rights reserved. +// Licensed under the MIT License. + +#pragma once + +#include +#include + +#include "core/providers/webgpu/math/matmul_algorithm.h" + +namespace onnxruntime { +namespace webgpu { + +struct MatMulAlgorithmSelectionParams { + int64_t m = 0; + int64_t n = 0; + int64_t k = 0; + bool can_use_subgroup_matrix = false; + bool has_intel_subgroup_capability = false; + bool use_split_k = false; +}; + +class MatMulAlgorithmScheduler { + public: + virtual ~MatMulAlgorithmScheduler() = default; + + MatMulAlgorithm Select( + const MatMulAlgorithmSelectionParams& params, + std::optional forced_algorithm = std::nullopt) const { + if (forced_algorithm.has_value()) { + return *forced_algorithm; + } + if (params.can_use_subgroup_matrix) { + return MatMulAlgorithm::SubgroupMatrix; + } + if (params.n < 8 && params.k < 8) { + return MatMulAlgorithm::Naive; + } + if (const auto vendor_algorithm = SelectVendorAlgorithm(params); vendor_algorithm.has_value()) { + return *vendor_algorithm; + } + if (params.use_split_k) { + return MatMulAlgorithm::PackedSplitK; + } + return MatMulAlgorithm::Packed; + } + + protected: + virtual std::optional SelectVendorAlgorithm( + const MatMulAlgorithmSelectionParams& /*params*/) const { + return std::nullopt; + } +}; + +} // namespace webgpu +} // namespace onnxruntime diff --git a/onnxruntime/core/providers/webgpu/vendor/intel/math/matmul_algorithm_scheduler.h b/onnxruntime/core/providers/webgpu/vendor/intel/math/matmul_algorithm_scheduler.h new file mode 100644 index 0000000000000..37bd97332fcb2 --- /dev/null +++ b/onnxruntime/core/providers/webgpu/vendor/intel/math/matmul_algorithm_scheduler.h @@ -0,0 +1,26 @@ +// Copyright (c) Microsoft Corporation. All rights reserved. +// Licensed under the MIT License. + +#pragma once + +#include "core/providers/webgpu/math/matmul_algorithm_scheduler.h" + +namespace onnxruntime { +namespace webgpu { +namespace intel { + +class IntelMatMulAlgorithmScheduler final : public MatMulAlgorithmScheduler { + protected: + std::optional SelectVendorAlgorithm( + const MatMulAlgorithmSelectionParams& params) const override { + if (params.has_intel_subgroup_capability && + params.m >= 64 && params.n >= 512 && params.k >= 32) { + return MatMulAlgorithm::IntelSubgroup; + } + return std::nullopt; + } +}; + +} // namespace intel +} // namespace webgpu +} // namespace onnxruntime diff --git a/onnxruntime/test/providers/webgpu/matmul_algorithm_scheduler_test.cc b/onnxruntime/test/providers/webgpu/matmul_algorithm_scheduler_test.cc index c8c68821588b8..6f0df83ccc54b 100644 --- a/onnxruntime/test/providers/webgpu/matmul_algorithm_scheduler_test.cc +++ b/onnxruntime/test/providers/webgpu/matmul_algorithm_scheduler_test.cc @@ -4,6 +4,8 @@ #include "gtest/gtest.h" #include "core/providers/webgpu/math/matmul_algorithm.h" +#include "core/providers/webgpu/math/matmul_algorithm_scheduler.h" +#include "core/providers/webgpu/vendor/intel/math/matmul_algorithm_scheduler.h" namespace onnxruntime { namespace webgpu { @@ -34,6 +36,68 @@ TEST(MatMulAlgorithmParsingTest, RejectsUnknownAlgorithmName) { EXPECT_EQ(ParseMatMulAlgorithm("unknown"), std::nullopt); } +TEST(MatMulAlgorithmSchedulerTest, ForcedAlgorithmTakesPrecedence) { + MatMulAlgorithmScheduler scheduler; + MatMulAlgorithmSelectionParams params{}; + params.can_use_subgroup_matrix = true; + + EXPECT_EQ(scheduler.Select(params, MatMulAlgorithm::Packed), MatMulAlgorithm::Packed); +} + +TEST(MatMulAlgorithmSchedulerTest, SubgroupMatrixTakesAutomaticPrecedence) { + MatMulAlgorithmScheduler scheduler; + MatMulAlgorithmSelectionParams params{}; + params.n = 4; + params.k = 4; + params.can_use_subgroup_matrix = true; + params.use_split_k = true; + + EXPECT_EQ(scheduler.Select(params), MatMulAlgorithm::SubgroupMatrix); +} + +TEST(MatMulAlgorithmSchedulerTest, NaiveUsesStrictSmallDimensionBoundaries) { + MatMulAlgorithmScheduler scheduler; + + MatMulAlgorithmSelectionParams small_params{}; + small_params.n = 7; + small_params.k = 7; + EXPECT_EQ(scheduler.Select(small_params), MatMulAlgorithm::Naive); + + MatMulAlgorithmSelectionParams n_boundary = small_params; + n_boundary.n = 8; + EXPECT_EQ(scheduler.Select(n_boundary), MatMulAlgorithm::Packed); + + MatMulAlgorithmSelectionParams k_boundary = small_params; + k_boundary.k = 8; + EXPECT_EQ(scheduler.Select(k_boundary), MatMulAlgorithm::Packed); +} + +TEST(MatMulAlgorithmSchedulerTest, IntelSchedulerAppliesCurrentVendorRule) { + intel::IntelMatMulAlgorithmScheduler scheduler; + MatMulAlgorithmSelectionParams params{}; + params.m = 64; + params.n = 512; + params.k = 32; + params.has_intel_subgroup_capability = true; + + EXPECT_EQ(scheduler.Select(params), MatMulAlgorithm::IntelSubgroup); + + params.n = 511; + EXPECT_EQ(scheduler.Select(params), MatMulAlgorithm::Packed); +} + +TEST(MatMulAlgorithmSchedulerTest, SplitKPrecedesPackedFallback) { + MatMulAlgorithmScheduler scheduler; + MatMulAlgorithmSelectionParams params{}; + params.n = 64; + params.k = 1024; + params.use_split_k = true; + EXPECT_EQ(scheduler.Select(params), MatMulAlgorithm::PackedSplitK); + + params.use_split_k = false; + EXPECT_EQ(scheduler.Select(params), MatMulAlgorithm::Packed); +} + } // namespace test } // namespace webgpu } // namespace onnxruntime From 3354ebf53f72b2a62cef0e3e96d2153244629e57 Mon Sep 17 00:00:00 2001 From: Jiajia Qin Date: Wed, 16 Sep 2026 13:20:49 +0800 Subject: [PATCH 06/25] refactor: dispatch WebGPU MatMul by algorithm --- .../core/providers/webgpu/math/matmul.cc | 382 +++++++++++------- .../core/providers/webgpu/math/matmul.h | 12 +- .../webgpu/math/matmul_algorithm_scheduler.h | 31 ++ .../webgpu/math/subgroup_matrix_matmul.cc | 97 +++-- .../webgpu/vendor/intel/math/matmul.cc | 5 +- .../webgpu/vendor/intel/math/matmul.h | 2 +- .../webgpu/matmul_algorithm_scheduler_test.cc | 37 ++ 7 files changed, 388 insertions(+), 178 deletions(-) diff --git a/onnxruntime/core/providers/webgpu/math/matmul.cc b/onnxruntime/core/providers/webgpu/math/matmul.cc index 65157575fbb3b..d4f5a066d988f 100644 --- a/onnxruntime/core/providers/webgpu/math/matmul.cc +++ b/onnxruntime/core/providers/webgpu/math/matmul.cc @@ -11,6 +11,7 @@ #include "core/providers/webgpu/webgpu_supported_types.h" #include "core/providers/webgpu/nn/fuse_utils.h" #include "core/providers/webgpu/data_transfer.h" +#include "core/providers/webgpu/vendor/intel/math/matmul_algorithm_scheduler.h" #include "core/providers/webgpu/vendor/intel/math/matmul.h" #include "core/providers/webgpu/webgpu_utils.h" @@ -26,6 +27,17 @@ MatMulOptImpl* MatMulOptImplCache::GetOrCreate(const ComputeContextBase& context return subgroup_impl_.get(); } +const MatMulAlgorithmScheduler& MatMulOptImplCache::GetOrCreateScheduler(const ComputeContextBase& context) { + std::call_once(scheduler_init_flag_, [&]() { + if (context.AdapterInfo().vendor == std::string_view{"intel"}) { + scheduler_ = std::make_unique(); + } else { + scheduler_ = std::make_unique(); + } + }); + return *scheduler_; +} + ONNX_OPERATOR_VERSIONED_KERNEL_EX( MatMul, kOnnxDomain, @@ -156,191 +168,273 @@ Status MatMul::ComputeInternal(ComputeContext& context) const { /*is_channels_last=*/true, compute_cache_, b_is_constant_); } -Status ComputeMatMul(ComputeContext* context, - const Activation& activation, std::vector& inputs, Tensor* output_tensor, bool is_channels_last, - MatMulOptImplCache& cache, - bool b_is_constant) { +static Status ApplyMatMulNaive(ComputeContext& context, + const Activation& activation, + const std::vector& inputs, + Tensor* output_tensor, + bool is_channels_last, + const MatMulComputeHelper& helper) { const auto* a = inputs[0]; const auto* b = inputs[1]; - bool has_bias = inputs.size() > 2; - const TensorShape& logical_a_shape = a->Shape(); - const TensorShape& logical_b_shape = b->Shape(); - ORT_RETURN_IF_NOT(logical_a_shape.NumDimensions() >= 2 && logical_b_shape.NumDimensions() >= 2, - "ComputeMatMul expects matrix or batched-matrix inputs."); - - MatMulComputeHelper helper; - ORT_THROW_IF_ERROR(helper.Compute(logical_a_shape, logical_b_shape)); - - MatMulOptImpl* subgroup_impl = cache.GetOrCreate(*context); - if (subgroup_impl != nullptr) { - bool handled = false; - ORT_RETURN_IF_ERROR(subgroup_impl->Compute( - *context, inputs, output_tensor, - activation, is_channels_last, b_is_constant, handled)); - if (handled) { - return Status::OK(); - } - } - if (helper.N() < 8 && helper.K() < 8) { - const uint32_t m = narrow(helper.M()); - const uint32_t n = narrow(helper.N()); - const uint32_t k = narrow(helper.K()); - const int components = GetMaxComponents(n); - const int a_components = GetMaxComponents(k); - const int64_t output_number = GetMaxComponents(m); - const TensorShape& logical_output_shape = helper.OutputShape(); - const size_t output_rank = logical_output_shape.NumDimensions(); - const TensorShape outer_dims = - output_rank > 2 ? logical_output_shape.Slice(0, output_rank - 2) : TensorShape({}); - const int64_t output_rows = logical_a_shape[logical_a_shape.NumDimensions() - 2]; - const TensorShape output_program_shape{ - outer_dims.Size(), output_rows, n / components}; - const uint32_t output_size = - narrow(logical_output_shape.Size() / components / output_number); - - MatMulNaiveProgram program{activation, output_rank, output_number, has_bias, is_channels_last}; - program - .CacheHint(activation.CacheKey(), std::to_string(components), - std::to_string(a_components), std::to_string(output_number), - std::to_string(is_channels_last)) - .AddInputs({{a, ProgramTensorMetadataDependency::TypeAndRank, a_components}, - {b, ProgramTensorMetadataDependency::TypeAndRank, components}}); - if (has_bias) { - const int bias_components = is_channels_last ? components : 1; - program.AddInput({inputs[2], ProgramTensorMetadataDependency::Rank, bias_components}); - } - program - .AddOutputs({{output_tensor, ProgramTensorMetadataDependency::None, - output_program_shape, components}}) - .SetDispatchGroupSize(CeilDiv(output_size, 64u)) - .AddIndices(outer_dims) - .AddUniformVariables({{output_size}, {m}, {n}, {k}}); - AppendActivationUniformsData(activation, program); - return context->RunProgram(program); + const bool has_bias = inputs.size() > 2; + const uint32_t m = narrow(helper.M()); + const uint32_t n = narrow(helper.N()); + const uint32_t k = narrow(helper.K()); + const int components = GetMaxComponents(n); + const int a_components = GetMaxComponents(k); + const int64_t output_number = GetMaxComponents(m); + const TensorShape& logical_output_shape = helper.OutputShape(); + const size_t output_rank = logical_output_shape.NumDimensions(); + const TensorShape outer_dims = + output_rank > 2 ? logical_output_shape.Slice(0, output_rank - 2) : TensorShape({}); + const int64_t output_rows = a->Shape()[a->Shape().NumDimensions() - 2]; + const TensorShape output_program_shape{ + outer_dims.Size(), output_rows, n / components}; + const uint32_t output_size = + narrow(logical_output_shape.Size() / components / output_number); + + MatMulNaiveProgram program{activation, output_rank, output_number, has_bias, is_channels_last}; + program + .CacheHint(activation.CacheKey(), std::to_string(components), + std::to_string(a_components), std::to_string(output_number), + std::to_string(is_channels_last)) + .AddInputs({{a, ProgramTensorMetadataDependency::TypeAndRank, a_components}, + {b, ProgramTensorMetadataDependency::TypeAndRank, components}}); + if (has_bias) { + const int bias_components = is_channels_last ? components : 1; + program.AddInput({inputs[2], ProgramTensorMetadataDependency::Rank, bias_components}); } + program + .AddOutputs({{output_tensor, ProgramTensorMetadataDependency::None, + output_program_shape, components}}) + .SetDispatchGroupSize(CeilDiv(output_size, 64u)) + .AddIndices(outer_dims) + .AddUniformVariables({{output_size}, {m}, {n}, {k}}); + AppendActivationUniformsData(activation, program); + return context.RunProgram(program); +} - if (intel::CanApplyMatMulIntel(*context, helper.M(), helper.N(), helper.K())) { - return intel::ApplyMatMulIntel(*context, activation, inputs, output_tensor, is_channels_last); +static bool ShouldUsePackedSplitK(ComputeContext& context, + const Activation& activation, + const std::vector& inputs, + bool is_channels_last, + const MatMulComputeHelper& helper) { + if (context.KernelContext().GetUseDeterministicCompute()) { + return false; } - TensorShape a_shape = logical_a_shape; - TensorShape b_shape = logical_b_shape; + TensorShape a_shape = inputs[0]->Shape(); + TensorShape b_shape = inputs[1]->Shape(); TensorShape output_shape = helper.OutputShape(); - const int64_t batchA = + const int64_t batch_a = a_shape.NumDimensions() > 2 ? a_shape.SizeToDimension(a_shape.NumDimensions() - 2) : 1; - const int64_t batchB = + const int64_t batch_b = b_shape.NumDimensions() > 2 ? b_shape.SizeToDimension(b_shape.NumDimensions() - 2) : 1; - - // The generic path benefits from folding A's batch dimensions into M when B - // is shared. The subgroup and Intel paths derive their own dispatch shapes - // directly from the tensor views and have already declined above. - if (batchA != 1 && batchB == 1) { - const int64_t batchAndM = a_shape.SizeToDimension(a_shape.NumDimensions() - 1); - a_shape = TensorShape({batchAndM, helper.K()}); + if (batch_a != 1 && batch_b == 1) { + const int64_t batch_and_m = a_shape.SizeToDimension(a_shape.NumDimensions() - 1); + a_shape = TensorShape({batch_and_m, helper.K()}); b_shape = TensorShape({helper.K(), helper.N()}); - output_shape = TensorShape({batchAndM, helper.N()}); + output_shape = TensorShape({batch_and_m, helper.N()}); } - // helpful dimension variables - TensorShape outer_dims_a = a_shape.NumDimensions() > 2 - ? a_shape.Slice(0, a_shape.NumDimensions() - 2) - : TensorShape({}); + const int64_t batch_size = output_shape.NumDimensions() > 2 + ? output_shape.SizeToDimension(output_shape.NumDimensions() - 2) + : 1; + const uint32_t m = narrow(a_shape[a_shape.NumDimensions() - 2]); + const uint32_t k = narrow(a_shape[a_shape.NumDimensions() - 1]); + const uint32_t n = narrow(b_shape[b_shape.NumDimensions() - 1]); + const bool is_vec4 = k % 4 == 0 && n % 4 == 0; + return context.GetSplitKConfig().UseSplitK( + is_vec4, activation.activation_kind_, batch_size, m, n, k, is_channels_last); +} - TensorShape outer_dims_b = b_shape.NumDimensions() > 2 - ? b_shape.Slice(0, b_shape.NumDimensions() - 2) - : TensorShape({}); +static Status ApplyMatMulPacked(ComputeContext& context, + const Activation& activation, + const std::vector& inputs, + Tensor* output_tensor, + bool is_channels_last, + const MatMulComputeHelper& helper, + bool use_split_k) { + const auto* a = inputs[0]; + const auto* b = inputs[1]; + const bool has_bias = inputs.size() > 2; + TensorShape a_shape = a->Shape(); + TensorShape b_shape = b->Shape(); + TensorShape output_shape = helper.OutputShape(); + const int64_t batch_a = + a_shape.NumDimensions() > 2 ? a_shape.SizeToDimension(a_shape.NumDimensions() - 2) : 1; + const int64_t batch_b = + b_shape.NumDimensions() > 2 ? b_shape.SizeToDimension(b_shape.NumDimensions() - 2) : 1; - TensorShape outer_dims = output_shape.NumDimensions() > 2 - ? output_shape.Slice(0, output_shape.NumDimensions() - 2) - : TensorShape({}); + if (batch_a != 1 && batch_b == 1) { + const int64_t batch_and_m = a_shape.SizeToDimension(a_shape.NumDimensions() - 1); + a_shape = TensorShape({batch_and_m, helper.K()}); + b_shape = TensorShape({helper.K(), helper.N()}); + output_shape = TensorShape({batch_and_m, helper.N()}); + } + const TensorShape outer_dims_a = a_shape.NumDimensions() > 2 + ? a_shape.Slice(0, a_shape.NumDimensions() - 2) + : TensorShape({}); + const TensorShape outer_dims_b = b_shape.NumDimensions() > 2 + ? b_shape.Slice(0, b_shape.NumDimensions() - 2) + : TensorShape({}); + const TensorShape outer_dims = output_shape.NumDimensions() > 2 + ? output_shape.Slice(0, output_shape.NumDimensions() - 2) + : TensorShape({}); const int64_t batch_size = outer_dims.Size(); - - // Get dimensions for matrix multiplication from TensorShape - const uint32_t dim_a_outer = narrow(a_shape[a_shape.NumDimensions() - 2]); // left matrix second dimension - const uint32_t dim_inner = narrow(a_shape[a_shape.NumDimensions() - 1]); // left matrix first dimension - const uint32_t dim_b_outer = narrow(b_shape[b_shape.NumDimensions() - 1]); // right matrix first dimension - + const uint32_t dim_a_outer = narrow(a_shape[a_shape.NumDimensions() - 2]); + const uint32_t dim_inner = narrow(a_shape[a_shape.NumDimensions() - 1]); + const uint32_t dim_b_outer = narrow(b_shape[b_shape.NumDimensions() - 1]); const bool is_vec4 = dim_inner % 4 == 0 && dim_b_outer % 4 == 0; InlinedVector elements_per_thread = dim_a_outer <= 8 ? InlinedVector({4, 1, 1}) : InlinedVector({4, 4, 1}); - - const uint32_t dispatch_x = narrow((dim_b_outer + MatMul::MATMUL_PACKED_WORKGROUP_SIZE_X * elements_per_thread[0] - 1) / - (MatMul::MATMUL_PACKED_WORKGROUP_SIZE_X * elements_per_thread[0])); - const uint32_t dispatch_y = narrow((dim_a_outer + MatMul::MATMUL_PACKED_WORKGROUP_SIZE_Y * elements_per_thread[1] - 1) / - (MatMul::MATMUL_PACKED_WORKGROUP_SIZE_Y * elements_per_thread[1])); - uint32_t dispatch_z = narrow((static_cast(batch_size) + MatMul::MATMUL_PACKED_WORKGROUP_SIZE_Z * elements_per_thread[2] - 1) / - (MatMul::MATMUL_PACKED_WORKGROUP_SIZE_Z * elements_per_thread[2])); + const uint32_t dispatch_x = narrow( + (dim_b_outer + MatMul::MATMUL_PACKED_WORKGROUP_SIZE_X * elements_per_thread[0] - 1) / + (MatMul::MATMUL_PACKED_WORKGROUP_SIZE_X * elements_per_thread[0])); + const uint32_t dispatch_y = narrow( + (dim_a_outer + MatMul::MATMUL_PACKED_WORKGROUP_SIZE_Y * elements_per_thread[1] - 1) / + (MatMul::MATMUL_PACKED_WORKGROUP_SIZE_Y * elements_per_thread[1])); + uint32_t dispatch_z = narrow( + (static_cast(batch_size) + + MatMul::MATMUL_PACKED_WORKGROUP_SIZE_Z * elements_per_thread[2] - 1) / + (MatMul::MATMUL_PACKED_WORKGROUP_SIZE_Z * elements_per_thread[2])); const int components = is_vec4 ? 4 : 1; - const TensorShape a_shape_temp = CreateMatMulIntermediateShape(outer_dims_a, dim_a_outer, dim_inner, components); - const TensorShape b_shape_temp = CreateMatMulIntermediateShape(outer_dims_b, dim_inner, dim_b_outer, components); - const TensorShape output_shape_temp = TensorShape({batch_size, dim_a_outer, dim_b_outer / components}); - + const TensorShape a_shape_temp = + CreateMatMulIntermediateShape(outer_dims_a, dim_a_outer, dim_inner, components); + const TensorShape b_shape_temp = + CreateMatMulIntermediateShape(outer_dims_b, dim_inner, dim_b_outer, components); + const TensorShape output_shape_temp{batch_size, dim_a_outer, dim_b_outer / components}; ProgramOutput output(output_tensor, ProgramTensorMetadataDependency::Rank, output_shape_temp, components); const Tensor* bias = has_bias ? inputs[2] : nullptr; bool use_bias_in_matmul = has_bias; uint32_t split_dim_inner = 1; uint32_t splits_per_batch = 1; - // Current Split-K implementation relies on atomic operations, which are not deterministic. - if (!context->KernelContext().GetUseDeterministicCompute()) { - const SplitKConfig& split_k_config = context->GetSplitKConfig(); - const bool need_split_k = split_k_config.UseSplitK( - is_vec4, activation.activation_kind_, batch_size, dim_a_outer, dim_b_outer, - dim_inner, is_channels_last); - if (need_split_k) { - ORT_ENFORCE(is_vec4, "Split-K MatMul requires vec4 packing."); - - if (has_bias) { - ORT_ENFORCE(is_channels_last, "Split-K MatMul only supports channels-last format."); - } - - // Initialize `output_tensor` with 0 or bias before MatMulProgram with Split-K enabled. - const auto fill_bias_program = CreateMatMulFillBiasOrZeroBeforeSplitKProgram(bias, output_tensor, /*is_gemm*/ false, /*beta*/ 1.0f, /*bias_components*/ 4, output_shape_temp, narrow(batch_size)); - ORT_RETURN_IF_ERROR(context->RunProgram(fill_bias_program)); - - // `bias` has been handled in the execution of `fill_bias_program` so we don't need to set - // `bias` again in `MatMulProgram`. - use_bias_in_matmul = false; - - // With Split-K, `dim_inner` will be split into multiple parts. `dispatch_z` encodes - // both the split-k index and the batch index: dispatch_z = splits_per_batch * batch_size. - split_dim_inner = split_k_config.GetSplitDimInner(); - splits_per_batch = (dim_inner + split_dim_inner - 1) / split_dim_inner; - const uint64_t dispatch_z_u64 = static_cast(batch_size) * static_cast(splits_per_batch); - ORT_ENFORCE(dispatch_z_u64 <= static_cast(std::numeric_limits::max()), - "dispatch_z exceeds uint32_t range: ", dispatch_z_u64); - dispatch_z = narrow(dispatch_z_u64); - - // The output should be declared in atomic types in `MatMulProgram` for the use of atomic - // built-in functions. - output.is_atomic = true; - } + if (use_split_k) { + ORT_RETURN_IF(context.KernelContext().GetUseDeterministicCompute(), + "MatMul algorithm packed_split_k does not support deterministic compute."); + ORT_RETURN_IF_NOT(context.GetSplitKConfig().GetSplitDimInner() != 0, + "MatMul algorithm packed_split_k is not configured for this adapter."); + ORT_RETURN_IF_NOT(is_vec4, + "MatMul algorithm packed_split_k requires vec4 packing."); + ORT_RETURN_IF_NOT(activation.activation_kind_ == ActivationKind::None, + "MatMul algorithm packed_split_k does not support a fused activation."); + ORT_RETURN_IF_NOT(!has_bias || is_channels_last, + "MatMul algorithm packed_split_k requires channels-last bias layout."); + + const auto fill_bias_program = CreateMatMulFillBiasOrZeroBeforeSplitKProgram( + bias, output_tensor, /*is_gemm=*/false, /*beta=*/1.0f, + /*output_components=*/4, output_shape_temp, narrow(batch_size)); + ORT_RETURN_IF_ERROR(context.RunProgram(fill_bias_program)); + use_bias_in_matmul = false; + split_dim_inner = context.GetSplitKConfig().GetSplitDimInner(); + splits_per_batch = (dim_inner + split_dim_inner - 1) / split_dim_inner; + const uint64_t dispatch_z_u64 = + static_cast(batch_size) * static_cast(splits_per_batch); + ORT_RETURN_IF_NOT(dispatch_z_u64 <= static_cast(std::numeric_limits::max()), + "MatMul algorithm packed_split_k dispatch_z exceeds uint32_t range: ", dispatch_z_u64); + dispatch_z = narrow(dispatch_z_u64); + output.is_atomic = true; } - MatMulProgram matmul_program{activation, use_bias_in_matmul, is_vec4, elements_per_thread, is_channels_last, split_dim_inner}; - matmul_program - .CacheHint(activation.CacheKey(), absl::StrJoin(elements_per_thread, "-"), std::to_string(is_vec4), components, is_channels_last, split_dim_inner) + MatMulProgram program{activation, use_bias_in_matmul, is_vec4, elements_per_thread, + is_channels_last, split_dim_inner}; + program + .CacheHint(activation.CacheKey(), absl::StrJoin(elements_per_thread, "-"), + std::to_string(is_vec4), components, is_channels_last, split_dim_inner) .AddInputs({{a, ProgramTensorMetadataDependency::TypeAndRank, a_shape_temp, components}, {b, ProgramTensorMetadataDependency::TypeAndRank, b_shape_temp, components}}) .AddUniformVariables({{dim_a_outer}, {dim_b_outer}, {dim_inner}, {dispatch_x}, {dispatch_y}, {dispatch_z}, {splits_per_batch}}) .AddIndices(outer_dims) .SetDispatchGroupSize(dispatch_x, dispatch_y, dispatch_z) - .SetWorkgroupSize(MatMul::MATMUL_PACKED_WORKGROUP_SIZE_X, MatMul::MATMUL_PACKED_WORKGROUP_SIZE_Y, MatMul::MATMUL_PACKED_WORKGROUP_SIZE_Z) + .SetWorkgroupSize(MatMul::MATMUL_PACKED_WORKGROUP_SIZE_X, + MatMul::MATMUL_PACKED_WORKGROUP_SIZE_Y, + MatMul::MATMUL_PACKED_WORKGROUP_SIZE_Z) .AddOutput(std::move(output)); - // Activation uniforms must remain last because definitions and values are matched by index. - AppendActivationUniformsData(activation, matmul_program); + AppendActivationUniformsData(activation, program); if (use_bias_in_matmul) { - auto bias_components = is_channels_last ? components : 1; - TensorShape reduced_bias_shape = ReduceShapeByComponents(bias->Shape(), bias_components); - matmul_program.AddInput({bias, ProgramTensorMetadataDependency::Rank, reduced_bias_shape, bias_components}); + const int bias_components = is_channels_last ? components : 1; + const TensorShape reduced_bias_shape = ReduceShapeByComponents(bias->Shape(), bias_components); + program.AddInput({bias, ProgramTensorMetadataDependency::Rank, reduced_bias_shape, bias_components}); + } + + return context.RunProgram(program); +} + +Status ComputeMatMul(ComputeContext* context, + const Activation& activation, std::vector& inputs, Tensor* output_tensor, + bool is_channels_last, MatMulOptImplCache& cache, + bool b_is_constant) { + const auto* a = inputs[0]; + const auto* b = inputs[1]; + const bool has_bias = inputs.size() > 2; + const TensorShape& logical_a_shape = a->Shape(); + const TensorShape& logical_b_shape = b->Shape(); + ORT_RETURN_IF_NOT(logical_a_shape.NumDimensions() >= 2 && logical_b_shape.NumDimensions() >= 2, + "ComputeMatMul expects matrix or batched-matrix inputs."); + + MatMulComputeHelper helper; + ORT_RETURN_IF_ERROR(helper.Compute(logical_a_shape, logical_b_shape)); + + MatMulOptImpl* subgroup_impl = cache.GetOrCreate(*context); + const bool can_use_subgroup_matrix = + subgroup_impl != nullptr && + subgroup_impl->CanApply(*context, inputs, is_channels_last, b_is_constant); + const bool has_intel_subgroup_capability = intel::HasMatMulIntelCapability(*context); + const bool use_split_k = + ShouldUsePackedSplitK(*context, activation, inputs, is_channels_last, helper); + + MatMulAlgorithmSelectionParams selection_params{}; + selection_params.m = helper.M(); + selection_params.n = helper.N(); + selection_params.k = helper.K(); + selection_params.can_use_subgroup_matrix = can_use_subgroup_matrix; + selection_params.has_intel_subgroup_capability = has_intel_subgroup_capability; + selection_params.use_split_k = use_split_k; + + const MatMulAlgorithm algorithm = + cache.GetOrCreateScheduler(*context).Select(selection_params, context->ForcedMatMulAlgorithm()); + + MatMulAlgorithmPrerequisites prerequisites{}; + prerequisites.can_use_subgroup_matrix = can_use_subgroup_matrix; + prerequisites.has_intel_subgroup_capability = has_intel_subgroup_capability; + prerequisites.split_k_configured = context->GetSplitKConfig().GetSplitDimInner() != 0; + prerequisites.deterministic_compute = context->KernelContext().GetUseDeterministicCompute(); + prerequisites.is_vec4 = helper.K() % 4 == 0 && helper.N() % 4 == 0; + prerequisites.has_fused_activation = activation.activation_kind_ != ActivationKind::None; + prerequisites.split_k_bias_layout_supported = !has_bias || is_channels_last; + ORT_RETURN_IF_NOT(MeetsMatMulAlgorithmPrerequisites(algorithm, prerequisites), + "MatMul algorithm ", MatMulAlgorithmName(algorithm), + " does not support these inputs or this device."); + + switch (algorithm) { + case MatMulAlgorithm::SubgroupMatrix: + ORT_RETURN_IF_NOT(subgroup_impl != nullptr, + "MatMul algorithm subgroup_matrix is unavailable."); + return subgroup_impl->Compute( + *context, inputs, output_tensor, activation, is_channels_last, b_is_constant); + case MatMulAlgorithm::Naive: + return ApplyMatMulNaive( + *context, activation, inputs, output_tensor, is_channels_last, helper); + case MatMulAlgorithm::IntelSubgroup: + return intel::ApplyMatMulIntel( + *context, activation, inputs, output_tensor, is_channels_last); + case MatMulAlgorithm::Packed: + return ApplyMatMulPacked( + *context, activation, inputs, output_tensor, is_channels_last, helper, + /*use_split_k=*/false); + case MatMulAlgorithm::PackedSplitK: + return ApplyMatMulPacked( + *context, activation, inputs, output_tensor, is_channels_last, helper, + /*use_split_k=*/true); } - return context->RunProgram(matmul_program); + return ORT_MAKE_STATUS(ONNXRUNTIME, FAIL, "Unknown MatMul algorithm."); } MatMulFillBiasOrZeroBeforeSplitKProgram CreateMatMulFillBiasOrZeroBeforeSplitKProgram( diff --git a/onnxruntime/core/providers/webgpu/math/matmul.h b/onnxruntime/core/providers/webgpu/math/matmul.h index 4aaca49f0edbe..fb572f0b2aee1 100644 --- a/onnxruntime/core/providers/webgpu/math/matmul.h +++ b/onnxruntime/core/providers/webgpu/math/matmul.h @@ -9,6 +9,7 @@ #include "core/providers/webgpu/webgpu_kernel.h" #include "core/providers/webgpu/program.h" #include "core/providers/cpu/math/matmul_helper.h" +#include "core/providers/webgpu/math/matmul_algorithm_scheduler.h" #include "core/providers/webgpu/math/matmul_utils.h" #include "core/providers/webgpu/math/matmul_packed.h" #include "core/providers/webgpu/webgpu_utils.h" @@ -21,13 +22,17 @@ class MatMulOptImpl { public: virtual ~MatMulOptImpl() = default; + virtual bool CanApply(const ComputeContext& context, + const std::vector& inputs, + bool is_channels_last, + bool b_is_constant) const = 0; + virtual Status Compute(ComputeContext& context, const std::vector& inputs, Tensor* output, const Activation& activation, bool is_channels_last, - bool b_is_constant, - /*out*/ bool& handled) = 0; + bool b_is_constant) = 0; }; class MatMulOptImplCache { @@ -36,10 +41,13 @@ class MatMulOptImplCache { ORT_DISALLOW_COPY_ASSIGNMENT_AND_MOVE(MatMulOptImplCache); MatMulOptImpl* GetOrCreate(const ComputeContextBase& context); + const MatMulAlgorithmScheduler& GetOrCreateScheduler(const ComputeContextBase& context); private: std::once_flag subgroup_impl_init_flag_; std::unique_ptr subgroup_impl_; + std::once_flag scheduler_init_flag_; + std::unique_ptr scheduler_; }; Status ComputeMatMul(ComputeContext* context, const Activation& activation, std::vector& inputs, Tensor* output, diff --git a/onnxruntime/core/providers/webgpu/math/matmul_algorithm_scheduler.h b/onnxruntime/core/providers/webgpu/math/matmul_algorithm_scheduler.h index 1c134d7553d50..eb2f9dc809cff 100644 --- a/onnxruntime/core/providers/webgpu/math/matmul_algorithm_scheduler.h +++ b/onnxruntime/core/providers/webgpu/math/matmul_algorithm_scheduler.h @@ -20,6 +20,37 @@ struct MatMulAlgorithmSelectionParams { bool use_split_k = false; }; +struct MatMulAlgorithmPrerequisites { + bool can_use_subgroup_matrix = false; + bool has_intel_subgroup_capability = false; + bool split_k_configured = false; + bool deterministic_compute = false; + bool is_vec4 = false; + bool has_fused_activation = false; + bool split_k_bias_layout_supported = true; +}; + +inline bool MeetsMatMulAlgorithmPrerequisites( + MatMulAlgorithm algorithm, + const MatMulAlgorithmPrerequisites& prerequisites) { + switch (algorithm) { + case MatMulAlgorithm::SubgroupMatrix: + return prerequisites.can_use_subgroup_matrix; + case MatMulAlgorithm::IntelSubgroup: + return prerequisites.has_intel_subgroup_capability; + case MatMulAlgorithm::PackedSplitK: + return prerequisites.split_k_configured && + !prerequisites.deterministic_compute && + prerequisites.is_vec4 && + !prerequisites.has_fused_activation && + prerequisites.split_k_bias_layout_supported; + case MatMulAlgorithm::Naive: + case MatMulAlgorithm::Packed: + return true; + } + return false; +} + class MatMulAlgorithmScheduler { public: virtual ~MatMulAlgorithmScheduler() = default; diff --git a/onnxruntime/core/providers/webgpu/math/subgroup_matrix_matmul.cc b/onnxruntime/core/providers/webgpu/math/subgroup_matrix_matmul.cc index f7361167594e7..f46b887463c9e 100644 --- a/onnxruntime/core/providers/webgpu/math/subgroup_matrix_matmul.cc +++ b/onnxruntime/core/providers/webgpu/math/subgroup_matrix_matmul.cc @@ -86,15 +86,10 @@ class SubgroupMatrixMatMulImpl final : public MatMulOptImpl { : config_index_(config_index), tiling_selector_(std::move(tiling_selector)) {} - Status Compute(ComputeContext& context, - const std::vector& inputs, - Tensor* output, - const Activation& activation, - bool is_channels_last, - bool b_is_constant, - /*out*/ bool& handled) override { - handled = false; - + bool CanApply(const ComputeContext& context, + const std::vector& inputs, + bool is_channels_last, + bool b_is_constant) const override { const auto* a = inputs[0]; const auto* b = inputs[1]; const auto& a_shape = a->Shape(); @@ -104,55 +99,102 @@ class SubgroupMatrixMatMulImpl final : public MatMulOptImpl { const size_t b_rank = b_shape.NumDimensions(); if ((!is_channels_last && has_bias) || a_rank < 2 || b_rank < 2 || !a->IsDataType() || !b->IsDataType()) { - return Status::OK(); + return false; } const uint32_t K = narrow(a_shape[a_rank - 1]); if (K == 0) { - return Status::OK(); + return false; } uint32_t M = 0; uint32_t N = 0; uint32_t batch = 1; if (b_rank == 2) { - ORT_ENFORCE(narrow(b_shape[0]) == K, - "MatMul contraction dim mismatch: A K=", K, " vs B rows=", b_shape[0]); + if (narrow(b_shape[0]) != K) { + return false; + } M = narrow(a_shape.Size() / static_cast(K)); N = narrow(b_shape[1]); } else { - if (a_rank != b_rank) { - return Status::OK(); + if (a_rank != b_rank || narrow(b_shape[b_rank - 2]) != K) { + return false; } - ORT_ENFORCE(narrow(b_shape[b_rank - 2]) == K, - "MatMul contraction dim mismatch: A K=", K, - " vs B rows=", b_shape[b_rank - 2]); M = narrow(a_shape[a_rank - 2]); N = narrow(b_shape[b_rank - 1]); for (size_t i = 0; i + 2 < a_rank; ++i) { if (a_shape[i] != b_shape[i]) { - return Status::OK(); + return false; } } batch = narrow(a_shape.SizeToDimension(a_rank - 2)); } if (M == 0 || N == 0) { - return Status::OK(); + return false; } const std::optional tiling = tiling_selector_(context, M, N, K, batch); if (!tiling) { - return Status::OK(); + return false; } + const auto& config = supported_subgroup_matrix_configs[config_index_]; + const bool needs_padded_b = N % 2 != 0; + if (needs_padded_b && (!b_is_constant || N == std::numeric_limits::max())) { + return false; + } + const uint32_t n_b = needs_padded_b ? N + 1 : N; + return config.K != 0 && K % config.K == 0 && + M >= tiling->tile_m && n_b >= tiling->tile_n; + } + + Status Compute(ComputeContext& context, + const std::vector& inputs, + Tensor* output, + const Activation& activation, + bool is_channels_last, + bool b_is_constant) override { + ORT_RETURN_IF_NOT(CanApply(context, inputs, is_channels_last, b_is_constant), + "MatMul algorithm subgroup_matrix does not support these inputs or this device."); + + const auto* a = inputs[0]; + const auto* b = inputs[1]; + const auto& a_shape = a->Shape(); + const auto& b_shape = b->Shape(); + const bool has_bias = inputs.size() > 2; + const size_t a_rank = a_shape.NumDimensions(); + const size_t b_rank = b_shape.NumDimensions(); + const uint32_t K = narrow(a_shape[a_rank - 1]); + + uint32_t M = 0; + uint32_t N = 0; + uint32_t batch = 1; + if (b_rank == 2) { + ORT_ENFORCE(narrow(b_shape[0]) == K, + "MatMul contraction dim mismatch: A K=", K, " vs B rows=", b_shape[0]); + M = narrow(a_shape.Size() / static_cast(K)); + N = narrow(b_shape[1]); + } else { + ORT_ENFORCE(narrow(b_shape[b_rank - 2]) == K, + "MatMul contraction dim mismatch: A K=", K, + " vs B rows=", b_shape[b_rank - 2]); + M = narrow(a_shape[a_rank - 2]); + N = narrow(b_shape[b_rank - 1]); + for (size_t i = 0; i + 2 < a_rank; ++i) { + ORT_ENFORCE(a_shape[i] == b_shape[i]); + } + batch = narrow(a_shape.SizeToDimension(a_rank - 2)); + } + + const std::optional tiling = tiling_selector_(context, M, N, K, batch); + ORT_ENFORCE(tiling.has_value()); + const auto& config = supported_subgroup_matrix_configs[config_index_]; const bool needs_padded_b = N % 2 != 0; // Require whole subgroup-matrix K blocks. An odd-width B must be constant // because its padded copy is cached by this implementation. - if (config.K == 0 || K % config.K != 0 || - (needs_padded_b && !b_is_constant)) { - return Status::OK(); - } + ORT_ENFORCE(config.K != 0 && K % config.K == 0 && + (!needs_padded_b || b_is_constant)); // N_b is just N rounded up to even - compute it before doing any padding work so // the tile-fit check below can bail out without a wasted pad dispatch. @@ -165,9 +207,7 @@ class SubgroupMatrixMatMulImpl final : public MatMulOptImpl { // The kernel keeps its operand loads in bounds by shifting a trailing partial // tile back, which is only possible when the tile fits within M and N. - if (M < tiling->tile_m || N_b < tiling->tile_n) { - return Status::OK(); - } + ORT_ENFORCE(M >= tiling->tile_m && N_b >= tiling->tile_n); // The optimized path will run: now materialize the even-strided B for odd N. const Tensor* b_used = b; @@ -207,7 +247,6 @@ class SubgroupMatrixMatMulImpl final : public MatMulOptImpl { } ORT_RETURN_IF_ERROR(context.RunProgram(program)); - handled = true; return Status::OK(); } diff --git a/onnxruntime/core/providers/webgpu/vendor/intel/math/matmul.cc b/onnxruntime/core/providers/webgpu/vendor/intel/math/matmul.cc index 9bb66fc551242..70a8e760751d4 100644 --- a/onnxruntime/core/providers/webgpu/vendor/intel/math/matmul.cc +++ b/onnxruntime/core/providers/webgpu/vendor/intel/math/matmul.cc @@ -36,8 +36,9 @@ Status MatMulSubgroupProgram::GenerateShaderCode(ShaderHelper& shader) const { return Status::OK(); } -bool CanApplyMatMulIntel(const ComputeContext& context, int64_t M, int64_t N, int64_t K) { - return CanApplySubgroup(context, M, N, K); +bool HasMatMulIntelCapability(const ComputeContext& context) { + return context.AdapterInfo().vendor == std::string_view{"intel"} && + context.HasFeature(wgpu::FeatureName::Subgroups); } Status ApplyMatMulIntel(ComputeContext& context, diff --git a/onnxruntime/core/providers/webgpu/vendor/intel/math/matmul.h b/onnxruntime/core/providers/webgpu/vendor/intel/math/matmul.h index 0887e0539a8b0..691cce96e8ace 100644 --- a/onnxruntime/core/providers/webgpu/vendor/intel/math/matmul.h +++ b/onnxruntime/core/providers/webgpu/vendor/intel/math/matmul.h @@ -46,7 +46,7 @@ class MatMulSubgroupProgram final : public Program { const InlinedVector elements_per_thread_; }; -bool CanApplyMatMulIntel(const ComputeContext& context, int64_t M, int64_t N, int64_t K); +bool HasMatMulIntelCapability(const ComputeContext& context); Status ApplyMatMulIntel(ComputeContext& context, const Activation& activation, diff --git a/onnxruntime/test/providers/webgpu/matmul_algorithm_scheduler_test.cc b/onnxruntime/test/providers/webgpu/matmul_algorithm_scheduler_test.cc index 6f0df83ccc54b..afb110d08850f 100644 --- a/onnxruntime/test/providers/webgpu/matmul_algorithm_scheduler_test.cc +++ b/onnxruntime/test/providers/webgpu/matmul_algorithm_scheduler_test.cc @@ -98,6 +98,43 @@ TEST(MatMulAlgorithmSchedulerTest, SplitKPrecedesPackedFallback) { EXPECT_EQ(scheduler.Select(params), MatMulAlgorithm::Packed); } +TEST(MatMulAlgorithmPrerequisiteTest, SplitKRejectsEachHardConstraint) { + MatMulAlgorithmPrerequisites prerequisites{}; + prerequisites.split_k_configured = true; + prerequisites.is_vec4 = true; + prerequisites.split_k_bias_layout_supported = true; + EXPECT_TRUE(MeetsMatMulAlgorithmPrerequisites(MatMulAlgorithm::PackedSplitK, prerequisites)); + + prerequisites.deterministic_compute = true; + EXPECT_FALSE(MeetsMatMulAlgorithmPrerequisites(MatMulAlgorithm::PackedSplitK, prerequisites)); + prerequisites.deterministic_compute = false; + + prerequisites.is_vec4 = false; + EXPECT_FALSE(MeetsMatMulAlgorithmPrerequisites(MatMulAlgorithm::PackedSplitK, prerequisites)); + prerequisites.is_vec4 = true; + + prerequisites.has_fused_activation = true; + EXPECT_FALSE(MeetsMatMulAlgorithmPrerequisites(MatMulAlgorithm::PackedSplitK, prerequisites)); + prerequisites.has_fused_activation = false; + + prerequisites.split_k_bias_layout_supported = false; + EXPECT_FALSE(MeetsMatMulAlgorithmPrerequisites(MatMulAlgorithm::PackedSplitK, prerequisites)); +} + +TEST(MatMulAlgorithmPrerequisiteTest, IntelCapabilityDoesNotIncludeAutomaticThresholds) { + MatMulAlgorithmPrerequisites prerequisites{}; + prerequisites.has_intel_subgroup_capability = true; + EXPECT_TRUE(MeetsMatMulAlgorithmPrerequisites(MatMulAlgorithm::IntelSubgroup, prerequisites)); + + MatMulAlgorithmSelectionParams below_heuristic_threshold{}; + below_heuristic_threshold.m = 1; + below_heuristic_threshold.n = 1; + below_heuristic_threshold.k = 1; + below_heuristic_threshold.has_intel_subgroup_capability = true; + intel::IntelMatMulAlgorithmScheduler scheduler; + EXPECT_NE(scheduler.Select(below_heuristic_threshold), MatMulAlgorithm::IntelSubgroup); +} + } // namespace test } // namespace webgpu } // namespace onnxruntime From 430487c2af57af769eda22751cbbc0b961224d1a Mon Sep 17 00:00:00 2001 From: Jiajia Qin Date: Wed, 16 Sep 2026 13:24:22 +0800 Subject: [PATCH 07/25] build: load Vulkan from system directory --- cmake/external/onnxruntime_external_deps.cmake | 2 ++ 1 file changed, 2 insertions(+) diff --git a/cmake/external/onnxruntime_external_deps.cmake b/cmake/external/onnxruntime_external_deps.cmake index cdd98b34c358c..6bf0b9a5bb97a 100644 --- a/cmake/external/onnxruntime_external_deps.cmake +++ b/cmake/external/onnxruntime_external_deps.cmake @@ -761,10 +761,12 @@ if (onnxruntime_USE_WEBGPU) set(DAWN_ENABLE_VULKAN OFF CACHE BOOL "" FORCE) endif() if (onnxruntime_ENABLE_DAWN_BACKEND_D3D12) + set(DAWN_FORCE_SYSTEM_COMPONENT_LOAD OFF CACHE BOOL "" FORCE) set(DAWN_USE_BUILT_DXC ON CACHE BOOL "" FORCE) set(TINT_BUILD_HLSL_WRITER ON CACHE BOOL "" FORCE) set(DAWN_ENABLE_D3D12 ON CACHE BOOL "" FORCE) else() + set(DAWN_FORCE_SYSTEM_COMPONENT_LOAD ON CACHE BOOL "" FORCE) set(DAWN_USE_BUILT_DXC OFF CACHE BOOL "" FORCE) set(TINT_BUILD_HLSL_WRITER OFF CACHE BOOL "" FORCE) set(DAWN_ENABLE_D3D12 OFF CACHE BOOL "" FORCE) From e83158fa62044e8f4c494ae08b6bfc6afeb34e6b Mon Sep 17 00:00:00 2001 From: Jiajia Qin Date: Wed, 16 Sep 2026 13:24:27 +0800 Subject: [PATCH 08/25] test: force WebGPU MatMul algorithms --- .../providers/webgpu/matmul_large_test.cc | 59 ++++++++++++++++++- 1 file changed, 56 insertions(+), 3 deletions(-) diff --git a/onnxruntime/test/providers/webgpu/matmul_large_test.cc b/onnxruntime/test/providers/webgpu/matmul_large_test.cc index 0f10d1b125cc4..381dd5dbfbe68 100644 --- a/onnxruntime/test/providers/webgpu/matmul_large_test.cc +++ b/onnxruntime/test/providers/webgpu/matmul_large_test.cc @@ -4,6 +4,8 @@ #include "gtest/gtest.h" #include "core/providers/cpu/math/matmul_helper.h" +#include "core/providers/webgpu/math/matmul_algorithm.h" +#include "core/providers/webgpu/webgpu_provider_options.h" #include "test/providers/provider_test_utils.h" #include "test/common/tensor_op_test_utils.h" #include "default_providers.h" @@ -42,10 +44,25 @@ static void ComputeExpectedResult(const std::vector& a_vals, const std::v template void RunTestTyped(std::initializer_list a_dims, std::initializer_list b_dims, - bool b_is_constant = false) { + bool b_is_constant = false, + std::optional forced_algorithm = std::nullopt, + OpTester::ExpectResult expected_result = OpTester::ExpectResult::kExpectSuccess, + const char* expected_error = nullptr) { static_assert(std::is_same_v || std::is_same_v, "unexpected type for T"); - auto webgpu_ep = DefaultWebGpuExecutionProvider(); + std::unique_ptr webgpu_ep; + if (forced_algorithm.has_value()) { + ConfigOptions config_options{}; + const std::string algorithm_name{webgpu::MatMulAlgorithmName(*forced_algorithm)}; + ASSERT_STATUS_OK(config_options.AddConfigEntry(webgpu::options::kDawnBackendType, + webgpu::options::kDawnBackendType_Vulkan)); + ASSERT_STATUS_OK(config_options.AddConfigEntry( + webgpu::options::kForceMatMulAlgorithm, + algorithm_name.c_str())); + webgpu_ep = WebGpuExecutionProviderWithOptions(config_options); + } else { + webgpu_ep = DefaultWebGpuExecutionProvider(); + } if (!webgpu_ep) { GTEST_SKIP() << "WebGPU execution provider is not available."; } @@ -79,7 +96,9 @@ void RunTestTyped(std::initializer_list a_dims, std::initializer_list @@ -107,6 +126,40 @@ TEST(MatMulProgramTest, VectorFallbackExecution) { RunTestTyped({2, 2, 8}, {8}); } +#if defined(_WIN32) +TEST(WebGpuMatMulAlgorithmTest, RejectsUnknownForcedAlgorithm) { + ConfigOptions config_options{}; + ASSERT_STATUS_OK(config_options.AddConfigEntry(webgpu::options::kForceMatMulAlgorithm, "unknown")); + EXPECT_THROW(WebGpuExecutionProviderWithOptions(config_options), OnnxRuntimeException); +} + +TEST(WebGpuMatMulAlgorithmTest, ForcedNaive) { + RunTestTyped({8, 8}, {8, 8}, false, webgpu::MatMulAlgorithm::Naive); +} + +TEST(WebGpuMatMulAlgorithmTest, ForcedPacked) { + RunTestTyped({2, 2}, {2, 2}, false, webgpu::MatMulAlgorithm::Packed); +} + +TEST(WebGpuMatMulAlgorithmTest, ForcedIntelSubgroup) { + RunTestTyped({64, 32}, {32, 64}, false, webgpu::MatMulAlgorithm::IntelSubgroup); +} + +TEST(WebGpuMatMulAlgorithmTest, ForcedPackedSplitK) { + RunTestTyped({1, 1024}, {1024, 16}, false, webgpu::MatMulAlgorithm::PackedSplitK); +} + +TEST(WebGpuMatMulAlgorithmTest, ForcedSubgroupMatrix) { + RunTestTyped({32, 16}, {16, 32}, false, webgpu::MatMulAlgorithm::SubgroupMatrix); +} + +TEST(WebGpuMatMulAlgorithmTest, ForcedSubgroupMatrixRejectsFloatInputs) { + RunTestTyped({32, 16}, {16, 32}, false, webgpu::MatMulAlgorithm::SubgroupMatrix, + OpTester::ExpectResult::kExpectFailure, + "MatMul algorithm subgroup_matrix"); +} +#endif + // 2D aligned baseline shapes. TEST(MatMul_Large, DISABLED_Aligned) { RunBothTypes({128, 64}, {64, 1024}); From 145b259a66e14e4ad9ef6cadb178165afeb3f1a3 Mon Sep 17 00:00:00 2001 From: Jiajia Qin Date: Fri, 18 Sep 2026 09:54:05 +0800 Subject: [PATCH 09/25] test: guard Vulkan MatMul algorithm tests --- onnxruntime/test/providers/webgpu/matmul_large_test.cc | 2 ++ 1 file changed, 2 insertions(+) diff --git a/onnxruntime/test/providers/webgpu/matmul_large_test.cc b/onnxruntime/test/providers/webgpu/matmul_large_test.cc index 381dd5dbfbe68..0bf8a8c628d27 100644 --- a/onnxruntime/test/providers/webgpu/matmul_large_test.cc +++ b/onnxruntime/test/providers/webgpu/matmul_large_test.cc @@ -133,6 +133,7 @@ TEST(WebGpuMatMulAlgorithmTest, RejectsUnknownForcedAlgorithm) { EXPECT_THROW(WebGpuExecutionProviderWithOptions(config_options), OnnxRuntimeException); } +#if defined(DAWN_ENABLE_VULKAN) TEST(WebGpuMatMulAlgorithmTest, ForcedNaive) { RunTestTyped({8, 8}, {8, 8}, false, webgpu::MatMulAlgorithm::Naive); } @@ -158,6 +159,7 @@ TEST(WebGpuMatMulAlgorithmTest, ForcedSubgroupMatrixRejectsFloatInputs) { OpTester::ExpectResult::kExpectFailure, "MatMul algorithm subgroup_matrix"); } +#endif // defined(DAWN_ENABLE_VULKAN) #endif // 2D aligned baseline shapes. From b6f7cbc25672f483dcd765a5535ce546f1d75b88 Mon Sep 17 00:00:00 2001 From: Jiajia Qin Date: Fri, 18 Sep 2026 11:20:24 +0800 Subject: [PATCH 10/25] test: skip unsupported forced MatMul paths --- .../providers/webgpu/matmul_large_test.cc | 94 +++++++++++++++++++ 1 file changed, 94 insertions(+) diff --git a/onnxruntime/test/providers/webgpu/matmul_large_test.cc b/onnxruntime/test/providers/webgpu/matmul_large_test.cc index 0bf8a8c628d27..17fbded94b997 100644 --- a/onnxruntime/test/providers/webgpu/matmul_large_test.cc +++ b/onnxruntime/test/providers/webgpu/matmul_large_test.cc @@ -5,6 +5,10 @@ #include "core/providers/cpu/math/matmul_helper.h" #include "core/providers/webgpu/math/matmul_algorithm.h" +#if !defined(ORT_USE_EP_API_ADAPTERS) +#include "core/providers/webgpu/math/subgroup_matrix_config.h" +#include "core/providers/webgpu/webgpu_context.h" +#endif #include "core/providers/webgpu/webgpu_provider_options.h" #include "test/providers/provider_test_utils.h" #include "test/common/tensor_op_test_utils.h" @@ -42,6 +46,88 @@ static void ComputeExpectedResult(const std::vector& a_vals, const std::v } } +#if defined(_WIN32) && defined(DAWN_ENABLE_VULKAN) +static std::optional GetForcedAlgorithmUnsupportedReason( + const IExecutionProvider& ep, + webgpu::MatMulAlgorithm algorithm) { +#if defined(ORT_USE_EP_API_ADAPTERS) + ORT_UNUSED_PARAMETER(ep); + switch (algorithm) { + case webgpu::MatMulAlgorithm::IntelSubgroup: + case webgpu::MatMulAlgorithm::PackedSplitK: + case webgpu::MatMulAlgorithm::SubgroupMatrix: + return "hardware-specific forced MatMul tests require direct adapter capability inspection."; + case webgpu::MatMulAlgorithm::Naive: + case webgpu::MatMulAlgorithm::Packed: + return std::nullopt; + } + return std::nullopt; +#else + auto& context = webgpu::WebGpuContextFactory::GetContext(ep.GetDeviceId()); + + switch (algorithm) { + case webgpu::MatMulAlgorithm::IntelSubgroup: + if (context.AdapterInfo().vendor != std::string_view{"intel"}) { + return "intel_subgroup requires an Intel adapter."; + } + if (!context.DeviceHasFeature(wgpu::FeatureName::Subgroups)) { + return "intel_subgroup requires the WebGPU Subgroups feature."; + } + break; + case webgpu::MatMulAlgorithm::PackedSplitK: + if (context.GetSplitKConfig().GetSplitDimInner() == 0) { + return "packed_split_k is not configured for the selected adapter."; + } + break; + case webgpu::MatMulAlgorithm::SubgroupMatrix: { + if (!context.DeviceHasFeature(wgpu::FeatureName::ChromiumExperimentalSubgroupMatrix)) { + return "subgroup_matrix requires the WebGPU subgroup-matrix feature."; + } + + const auto& adapter_info = context.AdapterInfo(); + const auto& device_configs = context.SubgroupMatrixConfigs(); + bool has_required_config = false; + for (const auto& required_config : webgpu::supported_subgroup_matrix_configs) { + if (!required_config.Is(8, 16, 16) || + required_config.componentType != wgpu::SubgroupMatrixComponentType::F16 || + required_config.resultComponentType != wgpu::SubgroupMatrixComponentType::F16) { + continue; + } + for (size_t i = 0; i < device_configs.configCount; ++i) { + const auto& device_config = device_configs.configs[i]; + if (device_config.componentType == required_config.componentType && + device_config.resultComponentType == required_config.resultComponentType && + device_config.M == required_config.M && + device_config.N == required_config.N && + device_config.K == required_config.K && + adapter_info.subgroupMinSize == required_config.subgroupMinSize && + adapter_info.subgroupMaxSize == required_config.subgroupMaxSize) { + has_required_config = true; + break; + } + } + if (has_required_config) { + break; + } + } + if (!has_required_config) { + return "subgroup_matrix requires an 8x16x16 F16 configuration with subgroup range 16-32."; + } + if (!context.DeviceHasFeature(wgpu::FeatureName::SubgroupSizeControl)) { + return "subgroup_matrix requires the WebGPU SubgroupSizeControl feature."; + } + break; + } + case webgpu::MatMulAlgorithm::Naive: + case webgpu::MatMulAlgorithm::Packed: + break; + } + + return std::nullopt; +#endif +} +#endif + template void RunTestTyped(std::initializer_list a_dims, std::initializer_list b_dims, bool b_is_constant = false, @@ -66,6 +152,14 @@ void RunTestTyped(std::initializer_list a_dims, std::initializer_list Date: Sun, 20 Sep 2026 09:11:12 +0800 Subject: [PATCH 11/25] build: keep Vulkan configuration test-only --- cmake/external/onnxruntime_external_deps.cmake | 10 ++++------ 1 file changed, 4 insertions(+), 6 deletions(-) diff --git a/cmake/external/onnxruntime_external_deps.cmake b/cmake/external/onnxruntime_external_deps.cmake index 6bf0b9a5bb97a..7e56d910c3fb2 100644 --- a/cmake/external/onnxruntime_external_deps.cmake +++ b/cmake/external/onnxruntime_external_deps.cmake @@ -751,6 +751,10 @@ if (onnxruntime_USE_WEBGPU) set(DAWN_ENABLE_SPIRV_VALIDATION OFF CACHE BOOL "" FORCE) if (WIN32) + # building this requires the HLSL writer to be enabled in Tint. TBD if that we need either of these to be ON. + set(DAWN_USE_BUILT_DXC ON CACHE BOOL "" FORCE) + set(TINT_BUILD_HLSL_WRITER ON CACHE BOOL "" FORCE) + if ((NOT onnxruntime_ENABLE_DAWN_BACKEND_VULKAN) AND (NOT onnxruntime_ENABLE_DAWN_BACKEND_D3D12)) message(FATAL_ERROR "At least one of onnxruntime_ENABLE_DAWN_BACKEND_VULKAN or onnxruntime_ENABLE_DAWN_BACKEND_D3D12 must be enabled when using Dawn on Windows.") endif() @@ -761,14 +765,8 @@ if (onnxruntime_USE_WEBGPU) set(DAWN_ENABLE_VULKAN OFF CACHE BOOL "" FORCE) endif() if (onnxruntime_ENABLE_DAWN_BACKEND_D3D12) - set(DAWN_FORCE_SYSTEM_COMPONENT_LOAD OFF CACHE BOOL "" FORCE) - set(DAWN_USE_BUILT_DXC ON CACHE BOOL "" FORCE) - set(TINT_BUILD_HLSL_WRITER ON CACHE BOOL "" FORCE) set(DAWN_ENABLE_D3D12 ON CACHE BOOL "" FORCE) else() - set(DAWN_FORCE_SYSTEM_COMPONENT_LOAD ON CACHE BOOL "" FORCE) - set(DAWN_USE_BUILT_DXC OFF CACHE BOOL "" FORCE) - set(TINT_BUILD_HLSL_WRITER OFF CACHE BOOL "" FORCE) set(DAWN_ENABLE_D3D12 OFF CACHE BOOL "" FORCE) endif() # We are currently always using the D3D12 backend. From 0b03a49a52509833696cf327c093f05be7ca2afd Mon Sep 17 00:00:00 2001 From: Jiajia Qin Date: Sun, 20 Sep 2026 09:16:44 +0800 Subject: [PATCH 12/25] docs: move WebGPU MatMul scheduler design --- .../webgpu_matmul_algorithm_scheduler.md} | 0 ...09-16-webgpu-matmul-algorithm-scheduler.md | 223 ------------------ 2 files changed, 223 deletions(-) rename docs/{superpowers/specs/2026-09-16-webgpu-matmul-algorithm-scheduler-design.md => design/webgpu_matmul_algorithm_scheduler.md} (100%) delete mode 100644 docs/superpowers/plans/2026-09-16-webgpu-matmul-algorithm-scheduler.md diff --git a/docs/superpowers/specs/2026-09-16-webgpu-matmul-algorithm-scheduler-design.md b/docs/design/webgpu_matmul_algorithm_scheduler.md similarity index 100% rename from docs/superpowers/specs/2026-09-16-webgpu-matmul-algorithm-scheduler-design.md rename to docs/design/webgpu_matmul_algorithm_scheduler.md diff --git a/docs/superpowers/plans/2026-09-16-webgpu-matmul-algorithm-scheduler.md b/docs/superpowers/plans/2026-09-16-webgpu-matmul-algorithm-scheduler.md deleted file mode 100644 index 9e8be4e90db7a..0000000000000 --- a/docs/superpowers/plans/2026-09-16-webgpu-matmul-algorithm-scheduler.md +++ /dev/null @@ -1,223 +0,0 @@ -# WebGPU MatMul Algorithm Scheduler Implementation Plan - -> **For agentic workers:** REQUIRED SUB-SKILL: Use superpowers:subagent-driven-development (recommended) or superpowers:executing-plans to implement this plan task-by-task. Steps use checkbox (`- [ ]`) syntax for tracking. - -**Goal:** Replace conditional trial dispatch in `ComputeMatMul` with an explicit, vendor-extensible algorithm scheduler and a test-only forced-algorithm session option. - -**Architecture:** A lightweight algorithm enum is shared by WebGPU provider configuration and MatMul. A pure scheduler preserves current precedence and delegates vendor rules to an Intel subclass; `ComputeMatMul` dispatches the chosen enum through one switch and validates hard prerequisites when a test forces a path. - -**Tech Stack:** C++17, ONNX Runtime WebGPU EP, GoogleTest, CMake/Visual Studio on Windows - -**Spec:** `docs/superpowers/specs/2026-09-16-webgpu-matmul-algorithm-scheduler-design.md` - -## Global Constraints - -- Automatic selection must preserve current behavior and rule precedence. -- Forced selection bypasses heuristics but never bypasses correctness or hardware prerequisites. -- Unsupported forced algorithms fail explicitly and never silently fall back. -- `PackedSplitK` is independently selectable from `Packed`. -- The option remains internal/test-only and is not added to public documentation. -- Existing MatMul, pointwise Conv, and Attention callers continue through `ComputeMatMul`. -- The Windows validation build enables Dawn Vulkan, disables Dawn D3D12, and tests explicitly select `dawnBackendType=Vulkan`. - ---- - -### Task 1: Algorithm Identity and Configuration Plumbing - -**Files:** -- Create: `onnxruntime/core/providers/webgpu/math/matmul_algorithm.h` -- Modify: `onnxruntime/core/providers/webgpu/webgpu_provider_options.h` -- Modify: `onnxruntime/core/providers/webgpu/webgpu_execution_provider.h` -- Modify: `onnxruntime/core/providers/webgpu/webgpu_execution_provider.cc` -- Modify: `onnxruntime/core/providers/webgpu/webgpu_provider_factory.cc` -- Modify: `onnxruntime/core/providers/webgpu/compute_context.h` -- Test: `onnxruntime/test/providers/webgpu/matmul_algorithm_scheduler_test.cc` - -**Interfaces:** -- Produces: `enum class MatMulAlgorithm { SubgroupMatrix, Naive, IntelSubgroup, Packed, PackedSplitK }`. -- Produces: `std::optional ParseMatMulAlgorithm(std::string_view)` and `std::string_view MatMulAlgorithmName(MatMulAlgorithm)`. -- Produces: `ComputeContextBase::ForcedMatMulAlgorithm() const` returning `std::optional`. - -- [ ] **Step 1: Write the failing enum parsing tests** - -```cpp -EXPECT_EQ(ParseMatMulAlgorithm("packed_split_k"), MatMulAlgorithm::PackedSplitK); -EXPECT_EQ(MatMulAlgorithmName(MatMulAlgorithm::PackedSplitK), "packed_split_k"); -EXPECT_EQ(ParseMatMulAlgorithm("unknown"), std::nullopt); -``` - -- [ ] **Step 2: Build to verify the new tests fail** - -```powershell -python tools/ci_build/build.py --config RelWithDebInfo --build_dir .\build\WGPU-Vulkan --use_webgpu --update --build --parallel --target onnxruntime_provider_test --cmake_extra_defines onnxruntime_ENABLE_DAWN_BACKEND_VULKAN=ON onnxruntime_ENABLE_DAWN_BACKEND_D3D12=OFF -``` - -Expected: compilation fails because the algorithm interface is missing. - -- [ ] **Step 3: Implement enum conversion and option plumbing** - -Use exact case-sensitive names `subgroup_matrix`, `naive`, `intel_subgroup`, `packed`, and `packed_split_k`. Declare `kForceMatMulAlgorithm = "ep.webgpuexecutionprovider.forceMatmulAlgorithm"`, parse it in `ParseEpConfig`, reject invalid values with the accepted-value list, store an `std::optional` in `WebGpuExecutionProvider`, and expose it through `ComputeContextBase`. - -- [ ] **Step 4: Build and run parser tests** - -Build with Step 2, locate `onnxruntime_provider_test.exe`, and run `--gtest_filter="MatMulAlgorithmParsingTest.*"` from its directory. Confirm a non-zero test count and zero failures. - -- [ ] **Step 5: Commit** - -```powershell -git add onnxruntime/core/providers/webgpu onnxruntime/test/providers/webgpu/matmul_algorithm_scheduler_test.cc -git commit -m "webgpu: add forced MatMul algorithm option" -``` - -### Task 2: Common and Intel Algorithm Schedulers - -**Files:** -- Modify: `onnxruntime/core/providers/webgpu/math/matmul.h` -- Modify: `onnxruntime/core/providers/webgpu/math/matmul.cc` -- Modify: `onnxruntime/core/providers/webgpu/vendor/intel/math/matmul.h` -- Modify: `onnxruntime/core/providers/webgpu/vendor/intel/math/matmul.cc` -- Test: `onnxruntime/test/providers/webgpu/matmul_algorithm_scheduler_test.cc` - -**Interfaces:** -- Consumes: `MatMulAlgorithm` and the optional forced value from Task 1. -- Produces: `MatMulAlgorithmSelectionParams` with common rule facts. -- Produces: `MatMulAlgorithmScheduler::Select(const MatMulAlgorithmSelectionParams&, std::optional) const`. -- Produces: an Intel scheduler overriding the vendor rule hook. - -- [ ] **Step 1: Write failing scheduler tests** - -Cover forced precedence, subgroup-matrix precedence, exact `N < 8 && K < 8` boundaries, Intel selection, Split-K selection, and packed fallback. - -```cpp -EXPECT_EQ(base.Select(small, std::nullopt), MatMulAlgorithm::Naive); -EXPECT_EQ(intel.Select(intel_problem, std::nullopt), MatMulAlgorithm::IntelSubgroup); -EXPECT_EQ(base.Select(split_k_problem, std::nullopt), MatMulAlgorithm::PackedSplitK); -EXPECT_EQ(base.Select({}, MatMulAlgorithm::Packed), MatMulAlgorithm::Packed); -``` - -- [ ] **Step 2: Build and confirm compilation fails for missing scheduler types** - -Run the Task 1 build command and retain the compiler failure as the red TDD result. - -- [ ] **Step 3: Implement the common scheduler** - -Implement a non-virtual `Select` that checks forced selection, subgroup matrix, naive, a protected virtual vendor hook, Split-K, and packed in that order. The base vendor hook returns `std::nullopt`. - -- [ ] **Step 4: Implement the Intel scheduler** - -The Intel override returns `IntelSubgroup` only for the supplied current-rule fact. Add a scheduler factory selected by `context.AdapterInfo().vendor`. - -- [ ] **Step 5: Build, run `--gtest_filter="MatMulAlgorithmSchedulerTest.*"`, and commit** - -```powershell -git add onnxruntime/core/providers/webgpu/math/matmul.h onnxruntime/core/providers/webgpu/math/matmul.cc onnxruntime/core/providers/webgpu/vendor/intel/math/matmul.h onnxruntime/core/providers/webgpu/vendor/intel/math/matmul.cc onnxruntime/test/providers/webgpu/matmul_algorithm_scheduler_test.cc -git commit -m "refactor: add WebGPU MatMul algorithm scheduler" -``` - -### Task 3: Direct Enum Dispatch - -**Files:** -- Modify: `onnxruntime/core/providers/webgpu/math/matmul.h` -- Modify: `onnxruntime/core/providers/webgpu/math/matmul.cc` -- Modify: `onnxruntime/core/providers/webgpu/math/subgroup_matrix_matmul.cc` -- Modify: `onnxruntime/core/providers/webgpu/vendor/intel/math/gemm_subgroup.h` -- Modify: `onnxruntime/core/providers/webgpu/vendor/intel/math/gemm_subgroup.cc` -- Modify: `onnxruntime/core/providers/webgpu/vendor/intel/math/matmul.h` -- Modify: `onnxruntime/core/providers/webgpu/vendor/intel/math/matmul.cc` -- Test: `onnxruntime/test/providers/webgpu/matmul_algorithm_scheduler_test.cc` - -**Interfaces:** -- Produces: `MatMulOptImpl::CanApply(...) const` and `Compute(...)` without a `handled` output. -- Produces: separate Intel hard-capability and automatic-heuristic checks. -- Produces: one `switch (algorithm)` in `ComputeMatMul` and one helper per implementation. - -- [ ] **Step 1: Write failing prerequisite tests** - -Cover forced Split-K rejection for deterministic compute and incompatible packing/activation/layout. Cover Intel hard capability separately from its `M >= 64 && N >= 512 && K >= 32` automatic heuristic. - -- [ ] **Step 2: Run focused tests and preserve the red result** - -Run `onnxruntime_provider_test.exe --gtest_filter="MatMulAlgorithm*"`; expect missing validation interfaces or failing assertions. - -- [ ] **Step 3: Split subgroup-matrix applicability from execution** - -Move non-mutating early-decline checks into `CanApply`, make `Compute` error if called inapplicably, remove `handled`, and preserve odd-N cached padding. - -- [ ] **Step 4: Separate Intel capability from heuristic policy** - -Hard capability is `vendor == intel && Subgroups`. Automatic policy adds the existing `M >= 64 && N >= 512 && K >= 32` thresholds. Forced mode uses only hard capability. - -- [ ] **Step 5: Extract naive and packed helpers** - -Make the packed helper accept explicit `bool use_split_k`. False never runs the Split-K heuristic; true validates a configured split size, non-deterministic compute, vec4 packing, no fused activation, and compatible bias layout before atomic accumulation. - -- [ ] **Step 6: Replace conditional execution with one switch** - -```cpp -switch (algorithm) { - case MatMulAlgorithm::SubgroupMatrix: return subgroup_impl->Compute(...); - case MatMulAlgorithm::Naive: return ApplyMatMulNaive(...); - case MatMulAlgorithm::IntelSubgroup: return intel::ApplyMatMulIntel(...); - case MatMulAlgorithm::Packed: return ApplyMatMulPacked(..., false); - case MatMulAlgorithm::PackedSplitK: return ApplyMatMulPacked(..., true); -} -``` - -Every forced-prerequisite error includes `MatMulAlgorithmName(algorithm)`. - -- [ ] **Step 7: Build, run `--gtest_filter="MatMulAlgorithm*"`, and commit** - -```powershell -git add onnxruntime/core/providers/webgpu/math onnxruntime/core/providers/webgpu/vendor/intel/math onnxruntime/test/providers/webgpu/matmul_algorithm_scheduler_test.cc -git commit -m "refactor: dispatch WebGPU MatMul by algorithm" -``` - -### Task 4: Forced-Path Integration and Verification - -**Files:** -- Modify: `onnxruntime/test/providers/webgpu/matmul_large_test.cc` -- Modify: `onnxruntime/test/providers/webgpu/matmul_algorithm_scheduler_test.cc` - -**Interfaces:** -- Consumes: `kForceMatMulAlgorithm` and `WebGpuExecutionProviderWithOptions`. -- Produces: end-to-end proof that forcing changes dispatch and preserves output. - -- [ ] **Step 1: Write failing integration tests** - -Create each EP with `kDawnBackendType=Vulkan`. Use a non-small shape forced to `Naive`, a small shape forced to `Packed`, a compatible large shape forced to `PackedSplitK`, and a compatible f16 shape forced to `SubgroupMatrix`. Add an expected-failure case with deliberately unmet hard prerequisites. - -- [ ] **Step 2: Run `--gtest_filter="WebGpuMatMulAlgorithmTest.*"` and preserve the failing result** - -Confirm a non-zero selected test count and at least one expected pre-implementation failure. - -- [ ] **Step 3: Complete only wiring or diagnostics exposed by the tests** - -Ensure test EP construction passes the option through `ConfigOptions`, the same path used by WebGPU session creation. - -- [ ] **Step 4: Run focused and regression tests** - -```powershell -.\onnxruntime_provider_test.exe --gtest_filter="MatMulAlgorithm*:WebGpuMatMulAlgorithmTest.*" -.\onnxruntime_provider_test.exe --gtest_filter="MathOpTest.MatMulFloatType:MathOpTest.MatMul_Float16:WebGpuMatMulLargeTest.*" -``` - -Confirm non-zero counts and zero failures. The subgroup-matrix case must verify that the Vulkan adapter advertises the required subgroup-matrix configuration and subgroup-size control; otherwise it may skip only with the exact missing capability in its reason. - -- [ ] **Step 5: Run formatting and diff checks** - -Run the repository C++ formatter/linter for changed files, then: - -```powershell -git diff --check main...HEAD -git status --short -git diff --stat main...HEAD -``` - -- [ ] **Step 6: Commit and report limitations** - -```powershell -git add onnxruntime/test/providers/webgpu/matmul_large_test.cc onnxruntime/test/providers/webgpu/matmul_algorithm_scheduler_test.cc -git commit -m "test: force WebGPU MatMul algorithms" -``` - -Report exact commands, counts, Vulkan adapter identity, and observed subgroup-matrix capabilities. State that macOS-arm64 Metal CI remains additional cross-backend validation and that lavapipe cannot execute this family reliably. From 50e2cc992749e68aa15ebc7163317c15ba109893 Mon Sep 17 00:00:00 2001 From: Jiajia Qin Date: Sun, 20 Sep 2026 11:25:47 +0800 Subject: [PATCH 13/25] fix: validate forced WebGPU MatMul edge cases --- .../core/providers/webgpu/math/matmul.cc | 1 + .../webgpu/math/matmul_algorithm_scheduler.h | 9 +++-- .../webgpu/vendor/intel/math/gemm_subgroup.cc | 12 +++++++ .../webgpu/vendor/intel/math/gemm_subgroup.h | 4 +++ .../webgpu/vendor/intel/math/matmul.cc | 5 +-- .../webgpu/matmul_algorithm_scheduler_test.cc | 33 +++++++++++++++++++ .../providers/webgpu/matmul_large_test.cc | 8 ++++- 7 files changed, 66 insertions(+), 6 deletions(-) diff --git a/onnxruntime/core/providers/webgpu/math/matmul.cc b/onnxruntime/core/providers/webgpu/math/matmul.cc index d4f5a066d988f..995b9d03aa941 100644 --- a/onnxruntime/core/providers/webgpu/math/matmul.cc +++ b/onnxruntime/core/providers/webgpu/math/matmul.cc @@ -403,6 +403,7 @@ Status ComputeMatMul(ComputeContext* context, MatMulAlgorithmPrerequisites prerequisites{}; prerequisites.can_use_subgroup_matrix = can_use_subgroup_matrix; prerequisites.has_intel_subgroup_capability = has_intel_subgroup_capability; + prerequisites.has_nonzero_k = helper.K() > 0; prerequisites.split_k_configured = context->GetSplitKConfig().GetSplitDimInner() != 0; prerequisites.deterministic_compute = context->KernelContext().GetUseDeterministicCompute(); prerequisites.is_vec4 = helper.K() % 4 == 0 && helper.N() % 4 == 0; diff --git a/onnxruntime/core/providers/webgpu/math/matmul_algorithm_scheduler.h b/onnxruntime/core/providers/webgpu/math/matmul_algorithm_scheduler.h index eb2f9dc809cff..9537809110e9a 100644 --- a/onnxruntime/core/providers/webgpu/math/matmul_algorithm_scheduler.h +++ b/onnxruntime/core/providers/webgpu/math/matmul_algorithm_scheduler.h @@ -23,6 +23,7 @@ struct MatMulAlgorithmSelectionParams { struct MatMulAlgorithmPrerequisites { bool can_use_subgroup_matrix = false; bool has_intel_subgroup_capability = false; + bool has_nonzero_k = false; bool split_k_configured = false; bool deterministic_compute = false; bool is_vec4 = false; @@ -39,13 +40,15 @@ inline bool MeetsMatMulAlgorithmPrerequisites( case MatMulAlgorithm::IntelSubgroup: return prerequisites.has_intel_subgroup_capability; case MatMulAlgorithm::PackedSplitK: - return prerequisites.split_k_configured && + return prerequisites.has_nonzero_k && + prerequisites.split_k_configured && !prerequisites.deterministic_compute && prerequisites.is_vec4 && !prerequisites.has_fused_activation && prerequisites.split_k_bias_layout_supported; - case MatMulAlgorithm::Naive: case MatMulAlgorithm::Packed: + return prerequisites.has_nonzero_k; + case MatMulAlgorithm::Naive: return true; } return false; @@ -64,7 +67,7 @@ class MatMulAlgorithmScheduler { if (params.can_use_subgroup_matrix) { return MatMulAlgorithm::SubgroupMatrix; } - if (params.n < 8 && params.k < 8) { + if (params.k == 0 || (params.n < 8 && params.k < 8)) { return MatMulAlgorithm::Naive; } if (const auto vendor_algorithm = SelectVendorAlgorithm(params); vendor_algorithm.has_value()) { diff --git a/onnxruntime/core/providers/webgpu/vendor/intel/math/gemm_subgroup.cc b/onnxruntime/core/providers/webgpu/vendor/intel/math/gemm_subgroup.cc index 5b17f06e08070..b4e6c195881e5 100644 --- a/onnxruntime/core/providers/webgpu/vendor/intel/math/gemm_subgroup.cc +++ b/onnxruntime/core/providers/webgpu/vendor/intel/math/gemm_subgroup.cc @@ -161,6 +161,18 @@ int64_t ElementsPerThreadY(ComputeContext& context, uint32_t M) { return M <= 8 ? 1 : (M <= 16 ? 2 : (M <= 32 ? 4 : (is_xe_lpg_or_xe_3lpg ? 4 : 8))); } +bool CanUseAVec4CooperativeLoad(std::string_view architecture, + uint32_t dim_inner, + int64_t elements_per_thread_y) { + // A 32-wide subgroup has four 8-lane cooperative-load groups. Every generated + // subgroup-size branch must distribute rows evenly across its lane groups. + constexpr int64_t max_lane_groups = 4; + return architecture == gpu_arch::kXe3Lpg && + dim_inner % 4 == 0 && + elements_per_thread_y > 0 && + elements_per_thread_y % max_lane_groups == 0; +} + Status MakeMatMulSubgroupSource(ShaderHelper& shader, const InlinedVector& elements_per_thread, const ShaderIndicesHelper* batch_dims, diff --git a/onnxruntime/core/providers/webgpu/vendor/intel/math/gemm_subgroup.h b/onnxruntime/core/providers/webgpu/vendor/intel/math/gemm_subgroup.h index 72babd38cd06d..86853467e60ec 100644 --- a/onnxruntime/core/providers/webgpu/vendor/intel/math/gemm_subgroup.h +++ b/onnxruntime/core/providers/webgpu/vendor/intel/math/gemm_subgroup.h @@ -22,6 +22,10 @@ bool CanApplySubgroup(const ComputeContext& context, int64_t M, int64_t N, int64 int64_t ElementsPerThreadY(ComputeContext& context, uint32_t M); +bool CanUseAVec4CooperativeLoad(std::string_view architecture, + uint32_t dim_inner, + int64_t elements_per_thread_y); + Status MakeMatMulSubgroupSource(ShaderHelper& shader, const InlinedVector& elements_per_thread, const ShaderIndicesHelper* batch_dims, diff --git a/onnxruntime/core/providers/webgpu/vendor/intel/math/matmul.cc b/onnxruntime/core/providers/webgpu/vendor/intel/math/matmul.cc index 70a8e760751d4..318ba4b3f35e9 100644 --- a/onnxruntime/core/providers/webgpu/vendor/intel/math/matmul.cc +++ b/onnxruntime/core/providers/webgpu/vendor/intel/math/matmul.cc @@ -107,11 +107,12 @@ Status ApplyMatMulIntel(ComputeContext& context, const bool is_vec4 = dim_b_outer % 4 == 0; // vec4 A loads and double-buffering of the B tile are only enabled on Xe-3LPG. const bool is_xe_3lpg = arch == gpu_arch::kXe3Lpg; - // Load A from global memory as vec4 when K is a multiple of 4; otherwise fall back to scalar load. - const bool a_vec4 = is_xe_3lpg && dim_inner % 4 == 0; // Double-buffering of the B tile (held in workgroup memory) is only enabled for float16 B inputs. const bool b_is_fp16 = is_xe_3lpg && b->GetElementType() == ONNX_NAMESPACE::TensorProto_DataType_FLOAT16; InlinedVector elements_per_thread = InlinedVector({4, ElementsPerThreadY(context, dim_a_outer), 1}); + // Fall back to scalar A loads when rows cannot be distributed evenly among + // the cooperative vec4 lane groups (for example, forced low-M execution). + const bool a_vec4 = CanUseAVec4CooperativeLoad(arch, dim_inner, elements_per_thread[1]); const uint32_t dispatch_x = narrow((dim_b_outer + kSubgroupLogicalWorkGroupSizeX * elements_per_thread[0] - 1) / (kSubgroupLogicalWorkGroupSizeX * elements_per_thread[0])); diff --git a/onnxruntime/test/providers/webgpu/matmul_algorithm_scheduler_test.cc b/onnxruntime/test/providers/webgpu/matmul_algorithm_scheduler_test.cc index afb110d08850f..51c9fe92bfe4c 100644 --- a/onnxruntime/test/providers/webgpu/matmul_algorithm_scheduler_test.cc +++ b/onnxruntime/test/providers/webgpu/matmul_algorithm_scheduler_test.cc @@ -5,6 +5,7 @@ #include "core/providers/webgpu/math/matmul_algorithm.h" #include "core/providers/webgpu/math/matmul_algorithm_scheduler.h" +#include "core/providers/webgpu/vendor/intel/math/gemm_subgroup.h" #include "core/providers/webgpu/vendor/intel/math/matmul_algorithm_scheduler.h" namespace onnxruntime { @@ -72,6 +73,15 @@ TEST(MatMulAlgorithmSchedulerTest, NaiveUsesStrictSmallDimensionBoundaries) { EXPECT_EQ(scheduler.Select(k_boundary), MatMulAlgorithm::Packed); } +TEST(MatMulAlgorithmSchedulerTest, ZeroContractionDimensionUsesNaive) { + MatMulAlgorithmScheduler scheduler; + MatMulAlgorithmSelectionParams params{}; + params.n = 8; + params.k = 0; + + EXPECT_EQ(scheduler.Select(params), MatMulAlgorithm::Naive); +} + TEST(MatMulAlgorithmSchedulerTest, IntelSchedulerAppliesCurrentVendorRule) { intel::IntelMatMulAlgorithmScheduler scheduler; MatMulAlgorithmSelectionParams params{}; @@ -100,6 +110,7 @@ TEST(MatMulAlgorithmSchedulerTest, SplitKPrecedesPackedFallback) { TEST(MatMulAlgorithmPrerequisiteTest, SplitKRejectsEachHardConstraint) { MatMulAlgorithmPrerequisites prerequisites{}; + prerequisites.has_nonzero_k = true; prerequisites.split_k_configured = true; prerequisites.is_vec4 = true; prerequisites.split_k_bias_layout_supported = true; @@ -121,6 +132,28 @@ TEST(MatMulAlgorithmPrerequisiteTest, SplitKRejectsEachHardConstraint) { EXPECT_FALSE(MeetsMatMulAlgorithmPrerequisites(MatMulAlgorithm::PackedSplitK, prerequisites)); } +TEST(MatMulAlgorithmPrerequisiteTest, PackedAlgorithmsRejectZeroContractionDimension) { + MatMulAlgorithmPrerequisites prerequisites{}; + prerequisites.split_k_configured = true; + prerequisites.is_vec4 = true; + prerequisites.split_k_bias_layout_supported = true; + + EXPECT_FALSE(MeetsMatMulAlgorithmPrerequisites(MatMulAlgorithm::Packed, prerequisites)); + EXPECT_FALSE(MeetsMatMulAlgorithmPrerequisites(MatMulAlgorithm::PackedSplitK, prerequisites)); + + prerequisites.has_nonzero_k = true; + EXPECT_TRUE(MeetsMatMulAlgorithmPrerequisites(MatMulAlgorithm::Packed, prerequisites)); + EXPECT_TRUE(MeetsMatMulAlgorithmPrerequisites(MatMulAlgorithm::PackedSplitK, prerequisites)); +} + +TEST(MatMulAlgorithmPrerequisiteTest, IntelAVec4RequiresCompatibleRowsPerThread) { + EXPECT_FALSE(intel::CanUseAVec4CooperativeLoad(intel::gpu_arch::kXe3Lpg, 32, 1)); + EXPECT_FALSE(intel::CanUseAVec4CooperativeLoad(intel::gpu_arch::kXe3Lpg, 32, 2)); + EXPECT_TRUE(intel::CanUseAVec4CooperativeLoad(intel::gpu_arch::kXe3Lpg, 32, 4)); + EXPECT_FALSE(intel::CanUseAVec4CooperativeLoad(intel::gpu_arch::kXe3Lpg, 31, 4)); + EXPECT_FALSE(intel::CanUseAVec4CooperativeLoad(intel::gpu_arch::kXeLpg, 32, 4)); +} + TEST(MatMulAlgorithmPrerequisiteTest, IntelCapabilityDoesNotIncludeAutomaticThresholds) { MatMulAlgorithmPrerequisites prerequisites{}; prerequisites.has_intel_subgroup_capability = true; diff --git a/onnxruntime/test/providers/webgpu/matmul_large_test.cc b/onnxruntime/test/providers/webgpu/matmul_large_test.cc index 17fbded94b997..9dbc3fcccf63e 100644 --- a/onnxruntime/test/providers/webgpu/matmul_large_test.cc +++ b/onnxruntime/test/providers/webgpu/matmul_large_test.cc @@ -236,8 +236,14 @@ TEST(WebGpuMatMulAlgorithmTest, ForcedPacked) { RunTestTyped({2, 2}, {2, 2}, false, webgpu::MatMulAlgorithm::Packed); } +TEST(WebGpuMatMulAlgorithmTest, ForcedPackedRejectsZeroContractionDimension) { + RunTestTyped({1, 0}, {0, 1}, false, webgpu::MatMulAlgorithm::Packed, + OpTester::ExpectResult::kExpectFailure, + "MatMul algorithm packed"); +} + TEST(WebGpuMatMulAlgorithmTest, ForcedIntelSubgroup) { - RunTestTyped({64, 32}, {32, 64}, false, webgpu::MatMulAlgorithm::IntelSubgroup); + RunTestTyped({8, 32}, {32, 64}, false, webgpu::MatMulAlgorithm::IntelSubgroup); } TEST(WebGpuMatMulAlgorithmTest, ForcedPackedSplitK) { From 614bd012a20b8d86f603b02b94dfb03bc2e40e26 Mon Sep 17 00:00:00 2001 From: Jiajia Qin Date: Mon, 21 Sep 2026 12:06:42 +0800 Subject: [PATCH 14/25] fix: address WebGPU MatMul review and CI failures --- .../webgpu/math/matmul_algorithm_scheduler.h | 3 ++- .../webgpu/vendor/intel/math/gemm_subgroup.h | 10 +------- .../vendor/intel/math/gemm_subgroup_utils.h | 24 +++++++++++++++++++ .../webgpu/matmul_algorithm_scheduler_test.cc | 13 +++++++++- .../providers/webgpu/matmul_large_test.cc | 20 +++++++++++----- 5 files changed, 53 insertions(+), 17 deletions(-) create mode 100644 onnxruntime/core/providers/webgpu/vendor/intel/math/gemm_subgroup_utils.h diff --git a/onnxruntime/core/providers/webgpu/math/matmul_algorithm_scheduler.h b/onnxruntime/core/providers/webgpu/math/matmul_algorithm_scheduler.h index 9537809110e9a..269ee589354bf 100644 --- a/onnxruntime/core/providers/webgpu/math/matmul_algorithm_scheduler.h +++ b/onnxruntime/core/providers/webgpu/math/matmul_algorithm_scheduler.h @@ -38,7 +38,8 @@ inline bool MeetsMatMulAlgorithmPrerequisites( case MatMulAlgorithm::SubgroupMatrix: return prerequisites.can_use_subgroup_matrix; case MatMulAlgorithm::IntelSubgroup: - return prerequisites.has_intel_subgroup_capability; + return prerequisites.has_intel_subgroup_capability && + prerequisites.has_nonzero_k; case MatMulAlgorithm::PackedSplitK: return prerequisites.has_nonzero_k && prerequisites.split_k_configured && diff --git a/onnxruntime/core/providers/webgpu/vendor/intel/math/gemm_subgroup.h b/onnxruntime/core/providers/webgpu/vendor/intel/math/gemm_subgroup.h index 86853467e60ec..dbe6b645af03b 100644 --- a/onnxruntime/core/providers/webgpu/vendor/intel/math/gemm_subgroup.h +++ b/onnxruntime/core/providers/webgpu/vendor/intel/math/gemm_subgroup.h @@ -4,16 +4,12 @@ #pragma once #include "core/providers/webgpu/shader_helper.h" +#include "core/providers/webgpu/vendor/intel/math/gemm_subgroup_utils.h" namespace onnxruntime { namespace webgpu { namespace intel { -namespace gpu_arch { -inline constexpr std::string_view kXeLpg = "xe-lpg"; -inline constexpr std::string_view kXe3Lpg = "xe-3lpg"; -} // namespace gpu_arch - const uint32_t kSubgroupLogicalWorkGroupSizeX = 32; const uint32_t kSubgroupLogicalWorkGroupSizeY = 8; const uint32_t kSubgroupLogicalWorkGroupSizeZ = 1; @@ -22,10 +18,6 @@ bool CanApplySubgroup(const ComputeContext& context, int64_t M, int64_t N, int64 int64_t ElementsPerThreadY(ComputeContext& context, uint32_t M); -bool CanUseAVec4CooperativeLoad(std::string_view architecture, - uint32_t dim_inner, - int64_t elements_per_thread_y); - Status MakeMatMulSubgroupSource(ShaderHelper& shader, const InlinedVector& elements_per_thread, const ShaderIndicesHelper* batch_dims, diff --git a/onnxruntime/core/providers/webgpu/vendor/intel/math/gemm_subgroup_utils.h b/onnxruntime/core/providers/webgpu/vendor/intel/math/gemm_subgroup_utils.h new file mode 100644 index 0000000000000..69e5d2907d7c2 --- /dev/null +++ b/onnxruntime/core/providers/webgpu/vendor/intel/math/gemm_subgroup_utils.h @@ -0,0 +1,24 @@ +// Copyright (c) Microsoft Corporation. All rights reserved. +// Licensed under the MIT License. + +#pragma once + +#include +#include + +namespace onnxruntime { +namespace webgpu { +namespace intel { + +namespace gpu_arch { +inline constexpr std::string_view kXeLpg = "xe-lpg"; +inline constexpr std::string_view kXe3Lpg = "xe-3lpg"; +} // namespace gpu_arch + +bool CanUseAVec4CooperativeLoad(std::string_view architecture, + uint32_t dim_inner, + int64_t elements_per_thread_y); + +} // namespace intel +} // namespace webgpu +} // namespace onnxruntime diff --git a/onnxruntime/test/providers/webgpu/matmul_algorithm_scheduler_test.cc b/onnxruntime/test/providers/webgpu/matmul_algorithm_scheduler_test.cc index 51c9fe92bfe4c..b052d4c69efe8 100644 --- a/onnxruntime/test/providers/webgpu/matmul_algorithm_scheduler_test.cc +++ b/onnxruntime/test/providers/webgpu/matmul_algorithm_scheduler_test.cc @@ -5,7 +5,7 @@ #include "core/providers/webgpu/math/matmul_algorithm.h" #include "core/providers/webgpu/math/matmul_algorithm_scheduler.h" -#include "core/providers/webgpu/vendor/intel/math/gemm_subgroup.h" +#include "core/providers/webgpu/vendor/intel/math/gemm_subgroup_utils.h" #include "core/providers/webgpu/vendor/intel/math/matmul_algorithm_scheduler.h" namespace onnxruntime { @@ -157,6 +157,7 @@ TEST(MatMulAlgorithmPrerequisiteTest, IntelAVec4RequiresCompatibleRowsPerThread) TEST(MatMulAlgorithmPrerequisiteTest, IntelCapabilityDoesNotIncludeAutomaticThresholds) { MatMulAlgorithmPrerequisites prerequisites{}; prerequisites.has_intel_subgroup_capability = true; + prerequisites.has_nonzero_k = true; EXPECT_TRUE(MeetsMatMulAlgorithmPrerequisites(MatMulAlgorithm::IntelSubgroup, prerequisites)); MatMulAlgorithmSelectionParams below_heuristic_threshold{}; @@ -168,6 +169,16 @@ TEST(MatMulAlgorithmPrerequisiteTest, IntelCapabilityDoesNotIncludeAutomaticThre EXPECT_NE(scheduler.Select(below_heuristic_threshold), MatMulAlgorithm::IntelSubgroup); } +TEST(MatMulAlgorithmPrerequisiteTest, IntelSubgroupRejectsZeroContractionDimension) { + MatMulAlgorithmPrerequisites prerequisites{}; + prerequisites.has_intel_subgroup_capability = true; + + EXPECT_FALSE(MeetsMatMulAlgorithmPrerequisites(MatMulAlgorithm::IntelSubgroup, prerequisites)); + + prerequisites.has_nonzero_k = true; + EXPECT_TRUE(MeetsMatMulAlgorithmPrerequisites(MatMulAlgorithm::IntelSubgroup, prerequisites)); +} + } // namespace test } // namespace webgpu } // namespace onnxruntime diff --git a/onnxruntime/test/providers/webgpu/matmul_large_test.cc b/onnxruntime/test/providers/webgpu/matmul_large_test.cc index 9dbc3fcccf63e..7f18532d678e5 100644 --- a/onnxruntime/test/providers/webgpu/matmul_large_test.cc +++ b/onnxruntime/test/providers/webgpu/matmul_large_test.cc @@ -25,6 +25,9 @@ static void ComputeExpectedResult(const std::vector& a_vals, const std::v const auto M = helper.M(); const auto K = helper.K(); const auto N = helper.N(); + if (K == 0) { + return; + } const auto& left_offsets = helper.LeftOffsets(); const auto& right_offsets = helper.RightOffsets(); const auto& output_offsets = helper.OutputOffsets(); @@ -100,8 +103,10 @@ static std::optional GetForcedAlgorithmUnsupportedReason( device_config.M == required_config.M && device_config.N == required_config.N && device_config.K == required_config.K && - adapter_info.subgroupMinSize == required_config.subgroupMinSize && - adapter_info.subgroupMaxSize == required_config.subgroupMaxSize) { + webgpu::IsSubgroupSizeSupported( + adapter_info.subgroupMinSize, adapter_info.subgroupMaxSize, + required_config.subgroupSize, + context.DeviceHasFeature(wgpu::FeatureName::SubgroupSizeControl))) { has_required_config = true; break; } @@ -111,10 +116,7 @@ static std::optional GetForcedAlgorithmUnsupportedReason( } } if (!has_required_config) { - return "subgroup_matrix requires an 8x16x16 F16 configuration with subgroup range 16-32."; - } - if (!context.DeviceHasFeature(wgpu::FeatureName::SubgroupSizeControl)) { - return "subgroup_matrix requires the WebGPU SubgroupSizeControl feature."; + return "subgroup_matrix requires an 8x16x16 F16 configuration with subgroup size 32."; } break; } @@ -246,6 +248,12 @@ TEST(WebGpuMatMulAlgorithmTest, ForcedIntelSubgroup) { RunTestTyped({8, 32}, {32, 64}, false, webgpu::MatMulAlgorithm::IntelSubgroup); } +TEST(WebGpuMatMulAlgorithmTest, ForcedIntelSubgroupRejectsZeroContractionDimension) { + RunTestTyped({1, 0}, {0, 1}, false, webgpu::MatMulAlgorithm::IntelSubgroup, + OpTester::ExpectResult::kExpectFailure, + "MatMul algorithm intel_subgroup"); +} + TEST(WebGpuMatMulAlgorithmTest, ForcedPackedSplitK) { RunTestTyped({1, 1024}, {1024, 16}, false, webgpu::MatMulAlgorithm::PackedSplitK); } From 4f0bef6c775532b564783d0fe5743234f0d75545 Mon Sep 17 00:00:00 2001 From: Jiajia Qin Date: Mon, 21 Sep 2026 12:16:05 +0800 Subject: [PATCH 15/25] docs: align MatMul scheduler validation details --- docs/design/webgpu_matmul_algorithm_scheduler.md | 2 +- onnxruntime/test/providers/webgpu/matmul_large_test.cc | 3 +++ 2 files changed, 4 insertions(+), 1 deletion(-) diff --git a/docs/design/webgpu_matmul_algorithm_scheduler.md b/docs/design/webgpu_matmul_algorithm_scheduler.md index c70ce892c2970..db25fd0f3ded4 100644 --- a/docs/design/webgpu_matmul_algorithm_scheduler.md +++ b/docs/design/webgpu_matmul_algorithm_scheduler.md @@ -21,7 +21,7 @@ There is no `Auto` algorithm value. Automatic versus forced selection is represe Add a `MatMulAlgorithmScheduler` base class. Its common rule order mirrors the existing `ComputeMatMul` condition order: 1. Select `SubgroupMatrix` when its implementation reports that it can handle the problem. -2. Select `Naive` when `N < 8 && K < 8`. +2. Select `Naive` when `K == 0` or when `N < 8 && K < 8`. 3. Ask a virtual vendor-policy hook for a vendor algorithm. 4. Select `PackedSplitK` when the existing `SplitKConfig::UseSplitK` rule succeeds. 5. Fall back to `Packed`. diff --git a/onnxruntime/test/providers/webgpu/matmul_large_test.cc b/onnxruntime/test/providers/webgpu/matmul_large_test.cc index 7f18532d678e5..a0bad0794e6fd 100644 --- a/onnxruntime/test/providers/webgpu/matmul_large_test.cc +++ b/onnxruntime/test/providers/webgpu/matmul_large_test.cc @@ -118,6 +118,9 @@ static std::optional GetForcedAlgorithmUnsupportedReason( if (!has_required_config) { return "subgroup_matrix requires an 8x16x16 F16 configuration with subgroup size 32."; } + if (!context.DeviceHasFeature(wgpu::FeatureName::SubgroupSizeControl)) { + return "subgroup_matrix requires the WebGPU SubgroupSizeControl feature."; + } break; } case webgpu::MatMulAlgorithm::Naive: From 296503f98201cc83e38abac245c0f58995cd8cf0 Mon Sep 17 00:00:00 2001 From: Jiajia Qin Date: Mon, 21 Sep 2026 12:32:23 +0800 Subject: [PATCH 16/25] fix: handle zero-K WebGPU MatMul inputs --- .../core/providers/webgpu/math/matmul.cc | 75 +++++++++++-------- .../core/providers/webgpu/math/matmul.h | 6 +- .../providers/webgpu/matmul_large_test.cc | 4 + 3 files changed, 51 insertions(+), 34 deletions(-) diff --git a/onnxruntime/core/providers/webgpu/math/matmul.cc b/onnxruntime/core/providers/webgpu/math/matmul.cc index 995b9d03aa941..6e8e246d7a9fc 100644 --- a/onnxruntime/core/providers/webgpu/math/matmul.cc +++ b/onnxruntime/core/providers/webgpu/math/matmul.cc @@ -73,13 +73,17 @@ static std::string CalcResult(int64_t components, int64_t a_components, int64_t } Status MatMulNaiveProgram::GenerateShaderCode(ShaderHelper& shader) const { - const auto& a = shader.AddInput("a", ShaderUsage::UseUniform | ShaderUsage::UseIndicesTypeAlias | - ShaderUsage::UseValueTypeAlias | ShaderUsage::UseElementTypeAlias); - const auto& b = shader.AddInput("b", ShaderUsage::UseUniform | ShaderUsage::UseIndicesTypeAlias | - ShaderUsage::UseValueTypeAlias | ShaderUsage::UseElementTypeAlias); - - const int a_components = a.NumComponents(); - const int components = b.NumComponents(); // components of N + const ShaderVariableHelper* a = nullptr; + const ShaderVariableHelper* b = nullptr; + int a_components = 1; + if (!is_zero_k_) { + a = &shader.AddInput("a", ShaderUsage::UseUniform | ShaderUsage::UseIndicesTypeAlias | + ShaderUsage::UseValueTypeAlias | ShaderUsage::UseElementTypeAlias); + b = &shader.AddInput("b", ShaderUsage::UseUniform | ShaderUsage::UseIndicesTypeAlias | + ShaderUsage::UseValueTypeAlias | ShaderUsage::UseElementTypeAlias); + a_components = a->NumComponents(); + } + const int components = NumberOfComponents(Outputs()[0].var_type); // components of N std::string process_bias; if (has_bias_) { @@ -93,32 +97,34 @@ Status MatMulNaiveProgram::GenerateShaderCode(ShaderHelper& shader) const { const auto& output = shader.AddOutput("output", ShaderUsage::UseUniform | ShaderUsage::UseIndicesTypeAlias | ShaderUsage::UseValueTypeAlias | ShaderUsage::UseElementTypeAlias); shader.AdditionalImplementation() << GetActivationDeclaration(activation_, "output_value_t", "output_element_t"); - const auto& batch_dims = shader.AddIndices("batch_dims"); shader.MainFunctionBody() << shader.GuardAgainstOutOfBoundsWorkgroupSizes("uniforms.output_size") << "let col = (global_idx % (uniforms.N / " << components << ")) * " << components << ";\n" << "var index1 = global_idx / (uniforms.N / " << components << ");\n" << "let stride1 = uniforms.M / " << output_number_ << ";\n" << "let row = (index1 % stride1) * " << output_number_ << ";\n" - << "let batch = index1 / stride1;\n"; - if (output_rank_ != 2) { - shader.MainFunctionBody() << "let batch_indices = " << batch_dims.OffsetToIndices("batch") << ";\n"; + << "let batch = index1 / stride1;\n" + << "var values: array;\n"; + if (!is_zero_k_) { + const auto& batch_dims = shader.AddIndices("batch_dims"); + if (output_rank_ != 2) { + shader.MainFunctionBody() << "let batch_indices = " << batch_dims.OffsetToIndices("batch") << ";\n"; + } + shader.MainFunctionBody() << "var a_indices: a_indices_t;\n" + << ConvertOutputBatchIndicesToInputBatchIndices("a", *a, a->Rank() - 2, batch_dims.Rank(), "batch_indices") + << a->IndicesSet("a_indices", a->Rank() - 2, 0) << "\n" + << a->IndicesSet("a_indices", a->Rank() - 1, 0) << "\n" + << "let a_offset = " << a->IndicesToOffset("a_indices") << "*" << a_components << ";\n" + << "var b_indices: b_indices_t;\n" + << ConvertOutputBatchIndicesToInputBatchIndices("b", *b, b->Rank() - 2, batch_dims.Rank(), "batch_indices") + << b->IndicesSet("b_indices", b->Rank() - 2, 0) << "\n" + << b->IndicesSet("b_indices", b->Rank() - 1, 0) << "\n" + << "let b_offset = " << b->IndicesToOffset("b_indices") << " * " << components << ";\n" + << "for (var k: u32 = 0u; k < uniforms.K; k = k + " << a_components << ") {\n" + << CalcResult(components, a_components, output_number_) << "\n" + << "}\n"; } - shader.MainFunctionBody() << "var a_indices: a_indices_t;\n" - << ConvertOutputBatchIndicesToInputBatchIndices("a", a, a.Rank() - 2, batch_dims.Rank(), "batch_indices") - << a.IndicesSet("a_indices", a.Rank() - 2, 0) << "\n" - << a.IndicesSet("a_indices", a.Rank() - 1, 0) << "\n" - << "let a_offset = " << a.IndicesToOffset("a_indices") << "*" << a_components << ";\n" - << "var b_indices: b_indices_t;\n" - << ConvertOutputBatchIndicesToInputBatchIndices("b", b, b.Rank() - 2, batch_dims.Rank(), "batch_indices") - << b.IndicesSet("b_indices", b.Rank() - 2, 0) << "\n" - << b.IndicesSet("b_indices", b.Rank() - 1, 0) << "\n" - << "let b_offset = " << b.IndicesToOffset("b_indices") << " * " << components << ";\n" - << "var values: array;\n" - << "for (var k: u32 = 0u; k < uniforms.K; k = k + " << a_components << ") {\n" - << CalcResult(components, a_components, output_number_) << "\n" - << "}\n" - << "for (var i = 0u; i < " << output_number_ << "u; i++) {\n" + shader.MainFunctionBody() << "for (var i = 0u; i < " << output_number_ << "u; i++) {\n" << " var value = values[i];\n" << process_bias << "\n" << apply_activation << "\n" @@ -180,8 +186,9 @@ static Status ApplyMatMulNaive(ComputeContext& context, const uint32_t m = narrow(helper.M()); const uint32_t n = narrow(helper.N()); const uint32_t k = narrow(helper.K()); + const bool is_zero_k = k == 0; const int components = GetMaxComponents(n); - const int a_components = GetMaxComponents(k); + const int a_components = is_zero_k ? 1 : GetMaxComponents(k); const int64_t output_number = GetMaxComponents(m); const TensorShape& logical_output_shape = helper.OutputShape(); const size_t output_rank = logical_output_shape.NumDimensions(); @@ -193,13 +200,18 @@ static Status ApplyMatMulNaive(ComputeContext& context, const uint32_t output_size = narrow(logical_output_shape.Size() / components / output_number); - MatMulNaiveProgram program{activation, output_rank, output_number, has_bias, is_channels_last}; + MatMulNaiveProgram program{activation, output_rank, output_number, has_bias, + is_channels_last, is_zero_k}; program .CacheHint(activation.CacheKey(), std::to_string(components), std::to_string(a_components), std::to_string(output_number), - std::to_string(is_channels_last)) - .AddInputs({{a, ProgramTensorMetadataDependency::TypeAndRank, a_components}, - {b, ProgramTensorMetadataDependency::TypeAndRank, components}}); + std::to_string(is_channels_last), std::to_string(is_zero_k)); + if (!is_zero_k) { + program + .AddInputs({{a, ProgramTensorMetadataDependency::TypeAndRank, a_components}, + {b, ProgramTensorMetadataDependency::TypeAndRank, components}}) + .AddIndices(outer_dims); + } if (has_bias) { const int bias_components = is_channels_last ? components : 1; program.AddInput({inputs[2], ProgramTensorMetadataDependency::Rank, bias_components}); @@ -208,7 +220,6 @@ static Status ApplyMatMulNaive(ComputeContext& context, .AddOutputs({{output_tensor, ProgramTensorMetadataDependency::None, output_program_shape, components}}) .SetDispatchGroupSize(CeilDiv(output_size, 64u)) - .AddIndices(outer_dims) .AddUniformVariables({{output_size}, {m}, {n}, {k}}); AppendActivationUniformsData(activation, program); return context.RunProgram(program); diff --git a/onnxruntime/core/providers/webgpu/math/matmul.h b/onnxruntime/core/providers/webgpu/math/matmul.h index fb572f0b2aee1..bc3a0ed17f211 100644 --- a/onnxruntime/core/providers/webgpu/math/matmul.h +++ b/onnxruntime/core/providers/webgpu/math/matmul.h @@ -86,8 +86,9 @@ class MatMul final : public WebGpuKernel { class MatMulNaiveProgram final : public Program { public: - MatMulNaiveProgram(const Activation& activation, const size_t output_rank, int64_t output_number, bool has_bias, bool is_channels_last = false) - : Program{"MatMulNaive"}, activation_(activation), output_rank_(output_rank), output_number_(output_number), has_bias_{has_bias}, is_channels_last_(is_channels_last) { + MatMulNaiveProgram(const Activation& activation, const size_t output_rank, int64_t output_number, + bool has_bias, bool is_channels_last = false, bool is_zero_k = false) + : Program{"MatMulNaive"}, activation_(activation), output_rank_(output_rank), output_number_(output_number), has_bias_{has_bias}, is_channels_last_(is_channels_last), is_zero_k_(is_zero_k) { } Status GenerateShaderCode(ShaderHelper& sh) const override; @@ -104,6 +105,7 @@ class MatMulNaiveProgram final : public Program { const int64_t output_number_; const bool has_bias_; const bool is_channels_last_; + const bool is_zero_k_; }; } // namespace webgpu diff --git a/onnxruntime/test/providers/webgpu/matmul_large_test.cc b/onnxruntime/test/providers/webgpu/matmul_large_test.cc index a0bad0794e6fd..0e5eb0a96982a 100644 --- a/onnxruntime/test/providers/webgpu/matmul_large_test.cc +++ b/onnxruntime/test/providers/webgpu/matmul_large_test.cc @@ -217,6 +217,10 @@ TEST(MatMulNaiveProgramTest, VectorExecution) { RunTestTyped({3}, {3}); } +TEST(MatMulNaiveProgramTest, ZeroContractionDimensionExecution) { + RunTestTyped({1, 0}, {0, 8}); +} + TEST(MatMulProgramTest, VectorFallbackExecution) { RunTestTyped({8}, {8, 3}); RunTestTyped({2, 8}, {8}); From a3e7c2dd2a6b513c5be4202506d56ed3f8e49620 Mon Sep 17 00:00:00 2001 From: Jiajia Qin Date: Mon, 21 Sep 2026 15:54:46 +0800 Subject: [PATCH 17/25] Refine WebGPU MatMul vendor scheduling and tuning --- .../webgpu_matmul_algorithm_scheduler.md | 32 +-- .../core/providers/webgpu/math/matmul.cc | 109 +++++++--- .../webgpu/math/matmul_algorithm_scheduler.h | 158 +++++++++++++- .../providers/webgpu/math/matmul_packed.cc | 7 +- .../providers/webgpu/math/matmul_packed.h | 23 +- .../intel/math/matmul_algorithm_scheduler.h | 3 + .../webgpu/matmul_algorithm_scheduler_test.cc | 198 +++++++++++++++++- 7 files changed, 466 insertions(+), 64 deletions(-) diff --git a/docs/design/webgpu_matmul_algorithm_scheduler.md b/docs/design/webgpu_matmul_algorithm_scheduler.md index db25fd0f3ded4..09c7040fdd679 100644 --- a/docs/design/webgpu_matmul_algorithm_scheduler.md +++ b/docs/design/webgpu_matmul_algorithm_scheduler.md @@ -2,7 +2,7 @@ ## Goal -Make every `ComputeMatMul` implementation path explicit and independently testable without changing the default runtime behavior. Selection policy must be separated from execution, preserve the current rule order, and allow vendor-specific policy to extend the common rules. +Make every `ComputeMatMul` implementation path explicit and independently testable. Selection policy must be separated from execution, retain the existing common rules as a fallback, and allow vendor-specific policy to override any performance heuristic. ## Algorithms @@ -18,17 +18,25 @@ There is no `Auto` algorithm value. Automatic versus forced selection is represe ## Selection Architecture -Add a `MatMulAlgorithmScheduler` base class. Its common rule order mirrors the existing `ComputeMatMul` condition order: +Add a `MatMulAlgorithmScheduler` base class. Automatic selection uses this order: -1. Select `SubgroupMatrix` when its implementation reports that it can handle the problem. -2. Select `Naive` when `K == 0` or when `N < 8 && K < 8`. -3. Ask a virtual vendor-policy hook for a vendor algorithm. -4. Select `PackedSplitK` when the existing `SplitKConfig::UseSplitK` rule succeeds. -5. Fall back to `Packed`. +1. Select `Naive` when `K == 0`. This is a correctness rule and cannot be overridden by vendor policy. +2. Ask the virtual `SelectVendorAlgorithm` policy hook for any algorithm. A vendor may use its own thresholds for every algorithm, not only vendor-specific implementations. +3. If the vendor returns no selection, call the private, non-virtual `SelectCommonAlgorithm` fallback. It selects `SubgroupMatrix` when supported, `Naive` when `N < 8 && K < 8`, `PackedSplitK` when the existing `SplitKConfig::UseSplitK` rule succeeds, and otherwise `Packed`. -An Intel-derived scheduler implements the vendor hook and selects `IntelSubgroup` under the current Intel subgroup rule. Other vendors use the base scheduler unchanged. Future vendor policies can derive from the base scheduler without adding vendor conditionals to `ComputeMatMul`. +An Intel-derived scheduler implements the vendor hook. It preserves the original policy by selecting `SubgroupMatrix` first when applicable, then selecting `IntelSubgroup` under the current Intel subgroup rule. Other vendors use the base scheduler unchanged. Future vendor policies can derive from the base scheduler, override as many performance ranges as needed, and return no selection to delegate the remaining ranges to the common fallback without adding vendor conditionals to `ComputeMatMul`. -The scheduler accepts already-computed selection facts rather than owning tensor execution. This keeps it deterministic and unit-testable without a WebGPU device. Vendor policy may inspect dimensions and device-derived capability facts, but it must return only a `MatMulAlgorithm`. +The scheduler accepts immutable problem facts rather than mutable policy decisions: logical and packed dimensions, batch sizes, adapter architecture, input types, packing and layout facts, deterministic-compute state, activation and bias facts, device capabilities, and the configured Split-K size. The common `SplitKConfig::UseSplitK` result is passed separately as the common fallback recommendation. A vendor may ignore that recommendation and apply independent thresholds from the raw facts. This keeps the scheduler deterministic and unit-testable without a WebGPU device. + +## Execution Configuration + +Algorithm selection and execution tuning are separate decisions. After selecting one enum, the scheduler creates a `MatMulExecutionPlan` containing that enum and a typed algorithm configuration. It first asks the protected `SelectVendorConfiguration` hook for tuning, then uses private common defaults when the vendor declines. The tuning hook runs for both automatic and forced algorithms, so the test-only forcing option controls the implementation path without disabling real device tuning. + +The packed configuration initially contains workgroup size, elements per thread, inner tile size, and Split-K size. `ApplyMatMulPacked` consumes those values directly and includes shader-affecting values in its cache key. The common configuration preserves the existing `8x8x1` workgroup, `4x1x1` or `4x4x1` elements-per-thread rule, inner tile size 32, and adapter Split-K size. A vendor may replace any of these values without changing `ComputeMatMul` or the packed implementation. + +Packed tuning is validated at the execution boundary. The current shader requires both Z workgroup dimensions to remain one because Z identifies a batch or Split-K slice rather than a tiled output axis. Split-K sizes must be greater than one and aligned to the inner tile so adjacent workgroups cannot overlap their K ranges. Dispatch counts use widened, overflow-safe arithmetic and are range-checked before conversion to WebGPU's 32-bit dimensions. + +Configuration is represented by an algorithm-specific variant rather than a bag of unrelated optional fields. The dispatcher rejects an algorithm/configuration type mismatch before execution. Empty configuration types reserve the same typed boundary for algorithms whose tuning remains inside their existing implementation. In particular, subgroup-matrix MatMul already receives a vendor-specific `SubgroupMatrixTilingSelector`; that existing selector continues to choose tile M, tile N, and split K without coupling the scheduler to device or shader classes. ## Forced Test Selection @@ -42,7 +50,7 @@ Invalid option strings fail during WebGPU provider creation and list accepted va ## Dispatch and Implementation Boundaries -`ComputeMatMul` computes shared shape facts once, asks the scheduler for exactly one enum, and switches directly to the matching implementation. Algorithm bodies are extracted into focused helpers where necessary. The generic packed helper takes an explicit Split-K mode; it does not re-run the selection heuristic. +`ComputeMatMul` computes shared shape facts once, asks the scheduler for one execution plan, validates its hard prerequisites and configuration type, and switches directly to the matching implementation. Algorithm bodies are extracted into focused helpers where necessary. The generic packed helper takes an explicit Split-K mode and packed configuration; it does not re-run selection or tuning heuristics. The subgroup-matrix optional implementation gains a non-mutating applicability query and an execution method that no longer communicates selection through a `handled` output. This removes trial execution as a dispatch mechanism. @@ -50,13 +58,13 @@ The existing per-kernel cache continues to own device-dependent subgroup-matrix ## Compatibility -With no forcing option, the selected algorithm and precedence remain equivalent to the current code. Existing call sites keep using `ComputeMatMul`; the refactor does not change the operator API or model semantics. +With no forcing option or vendor override, the common scheduler remains equivalent to the previous selection order. Existing call sites keep using `ComputeMatMul`; the refactor does not change the operator API or model semantics. The option is intentionally internal and test-only: it is declared with WebGPU provider options for configuration plumbing but is not added to public user documentation. ## Testing -- Add device-independent scheduler unit tests covering every common branch, precedence, Intel override, default fallback, and forced override. +- Add device-independent scheduler unit tests covering every common branch, forced-over-vendor precedence, the zero-K correctness guard, vendor-over-common precedence, Intel override, default fallback, independent vendor thresholds, common packed defaults, and vendor tuning of a forced algorithm. - Add parser/configuration tests for every accepted value and invalid input. - Add WebGPU MatMul tests that choose shapes which normally select a different path, force a compatible algorithm, and verify numerical output. Hardware-specific forced algorithms are tested only when their hard capabilities are present; strict-failure tests cover unsupported forced choices. - Build Dawn and the WebGPU provider on Windows with the Vulkan backend enabled and D3D12 disabled. Tests explicitly request `dawnBackendType=Vulkan`. diff --git a/onnxruntime/core/providers/webgpu/math/matmul.cc b/onnxruntime/core/providers/webgpu/math/matmul.cc index 6e8e246d7a9fc..a32b5197df49d 100644 --- a/onnxruntime/core/providers/webgpu/math/matmul.cc +++ b/onnxruntime/core/providers/webgpu/math/matmul.cc @@ -265,6 +265,7 @@ static Status ApplyMatMulPacked(ComputeContext& context, Tensor* output_tensor, bool is_channels_last, const MatMulComputeHelper& helper, + const MatMulPackedConfiguration& configuration, bool use_split_k) { const auto* a = inputs[0]; const auto* b = inputs[1]; @@ -299,19 +300,26 @@ static Status ApplyMatMulPacked(ComputeContext& context, const uint32_t dim_b_outer = narrow(b_shape[b_shape.NumDimensions() - 1]); const bool is_vec4 = dim_inner % 4 == 0 && dim_b_outer % 4 == 0; - InlinedVector elements_per_thread = dim_a_outer <= 8 - ? InlinedVector({4, 1, 1}) - : InlinedVector({4, 4, 1}); - const uint32_t dispatch_x = narrow( - (dim_b_outer + MatMul::MATMUL_PACKED_WORKGROUP_SIZE_X * elements_per_thread[0] - 1) / - (MatMul::MATMUL_PACKED_WORKGROUP_SIZE_X * elements_per_thread[0])); - const uint32_t dispatch_y = narrow( - (dim_a_outer + MatMul::MATMUL_PACKED_WORKGROUP_SIZE_Y * elements_per_thread[1] - 1) / - (MatMul::MATMUL_PACKED_WORKGROUP_SIZE_Y * elements_per_thread[1])); - uint32_t dispatch_z = narrow( - (static_cast(batch_size) + - MatMul::MATMUL_PACKED_WORKGROUP_SIZE_Z * elements_per_thread[2] - 1) / - (MatMul::MATMUL_PACKED_WORKGROUP_SIZE_Z * elements_per_thread[2])); + ORT_RETURN_IF_NOT(IsMatMulPackedConfigurationValid(configuration, use_split_k), + "MatMul packed configuration is invalid for ", + use_split_k ? "Split-K." : "the packed algorithm."); + InlinedVector elements_per_thread{ + configuration.elements_per_thread[0], + configuration.elements_per_thread[1], + configuration.elements_per_thread[2]}; + const auto dispatch_x_value = TryGetMatMulPackedDispatchGroupCount( + dim_b_outer, configuration.workgroup_size[0], configuration.elements_per_thread[0]); + const auto dispatch_y_value = TryGetMatMulPackedDispatchGroupCount( + dim_a_outer, configuration.workgroup_size[1], configuration.elements_per_thread[1]); + const auto dispatch_z_value = TryGetMatMulPackedDispatchGroupCount( + batch_size, configuration.workgroup_size[2], configuration.elements_per_thread[2]); + ORT_RETURN_IF_NOT(dispatch_x_value.has_value() && + dispatch_y_value.has_value() && + dispatch_z_value.has_value(), + "MatMul packed dispatch dimensions exceed uint32 limits."); + const uint32_t dispatch_x = *dispatch_x_value; + const uint32_t dispatch_y = *dispatch_y_value; + uint32_t dispatch_z = *dispatch_z_value; const int components = is_vec4 ? 4 : 1; const TensorShape a_shape_temp = @@ -328,7 +336,7 @@ static Status ApplyMatMulPacked(ComputeContext& context, if (use_split_k) { ORT_RETURN_IF(context.KernelContext().GetUseDeterministicCompute(), "MatMul algorithm packed_split_k does not support deterministic compute."); - ORT_RETURN_IF_NOT(context.GetSplitKConfig().GetSplitDimInner() != 0, + ORT_RETURN_IF_NOT(configuration.split_dim_inner > 1, "MatMul algorithm packed_split_k is not configured for this adapter."); ORT_RETURN_IF_NOT(is_vec4, "MatMul algorithm packed_split_k requires vec4 packing."); @@ -342,8 +350,12 @@ static Status ApplyMatMulPacked(ComputeContext& context, /*output_components=*/4, output_shape_temp, narrow(batch_size)); ORT_RETURN_IF_ERROR(context.RunProgram(fill_bias_program)); use_bias_in_matmul = false; - split_dim_inner = context.GetSplitKConfig().GetSplitDimInner(); - splits_per_batch = (dim_inner + split_dim_inner - 1) / split_dim_inner; + split_dim_inner = configuration.split_dim_inner; + const auto splits_per_batch_value = TryGetMatMulPackedDispatchGroupCount( + dim_inner, split_dim_inner, 1); + ORT_RETURN_IF_NOT(splits_per_batch_value.has_value(), + "MatMul packed Split-K count exceeds uint32 limits."); + splits_per_batch = *splits_per_batch_value; const uint64_t dispatch_z_u64 = static_cast(batch_size) * static_cast(splits_per_batch); ORT_RETURN_IF_NOT(dispatch_z_u64 <= static_cast(std::numeric_limits::max()), @@ -353,18 +365,20 @@ static Status ApplyMatMulPacked(ComputeContext& context, } MatMulProgram program{activation, use_bias_in_matmul, is_vec4, elements_per_thread, - is_channels_last, split_dim_inner}; + is_channels_last, split_dim_inner, configuration.tile_inner}; program .CacheHint(activation.CacheKey(), absl::StrJoin(elements_per_thread, "-"), - std::to_string(is_vec4), components, is_channels_last, split_dim_inner) + absl::StrJoin(configuration.workgroup_size, "-"), + std::to_string(is_vec4), components, is_channels_last, + split_dim_inner, configuration.tile_inner) .AddInputs({{a, ProgramTensorMetadataDependency::TypeAndRank, a_shape_temp, components}, {b, ProgramTensorMetadataDependency::TypeAndRank, b_shape_temp, components}}) .AddUniformVariables({{dim_a_outer}, {dim_b_outer}, {dim_inner}, {dispatch_x}, {dispatch_y}, {dispatch_z}, {splits_per_batch}}) .AddIndices(outer_dims) .SetDispatchGroupSize(dispatch_x, dispatch_y, dispatch_z) - .SetWorkgroupSize(MatMul::MATMUL_PACKED_WORKGROUP_SIZE_X, - MatMul::MATMUL_PACKED_WORKGROUP_SIZE_Y, - MatMul::MATMUL_PACKED_WORKGROUP_SIZE_Z) + .SetWorkgroupSize(configuration.workgroup_size[0], + configuration.workgroup_size[1], + configuration.workgroup_size[2]) .AddOutput(std::move(output)); AppendActivationUniformsData(activation, program); @@ -400,25 +414,60 @@ Status ComputeMatMul(ComputeContext* context, const bool use_split_k = ShouldUsePackedSplitK(*context, activation, inputs, is_channels_last, helper); + const int64_t batch_a = + logical_a_shape.NumDimensions() > 2 + ? logical_a_shape.SizeToDimension(logical_a_shape.NumDimensions() - 2) + : 1; + const int64_t batch_b = + logical_b_shape.NumDimensions() > 2 + ? logical_b_shape.SizeToDimension(logical_b_shape.NumDimensions() - 2) + : 1; + const TensorShape& logical_output_shape = helper.OutputShape(); + const uint64_t batch_size = logical_output_shape.NumDimensions() > 2 + ? narrow(logical_output_shape.SizeToDimension( + logical_output_shape.NumDimensions() - 2)) + : 1; + const bool folds_batch_into_m = batch_a != 1 && batch_b == 1; + MatMulAlgorithmSelectionParams selection_params{}; selection_params.m = helper.M(); selection_params.n = helper.N(); selection_params.k = helper.K(); + selection_params.packed_m = folds_batch_into_m + ? logical_a_shape.SizeToDimension(logical_a_shape.NumDimensions() - 1) + : helper.M(); + selection_params.batch_size = batch_size; + selection_params.packed_batch_size = folds_batch_into_m ? 1 : batch_size; + selection_params.adapter_architecture = context->AdapterInfo().architecture; + selection_params.a_data_type = a->GetElementType(); + selection_params.b_data_type = b->GetElementType(); selection_params.can_use_subgroup_matrix = can_use_subgroup_matrix; selection_params.has_intel_subgroup_capability = has_intel_subgroup_capability; - selection_params.use_split_k = use_split_k; - - const MatMulAlgorithm algorithm = - cache.GetOrCreateScheduler(*context).Select(selection_params, context->ForcedMatMulAlgorithm()); + selection_params.is_vec4 = helper.K() % 4 == 0 && helper.N() % 4 == 0; + selection_params.deterministic_compute = context->KernelContext().GetUseDeterministicCompute(); + selection_params.has_fused_activation = activation.activation_kind_ != ActivationKind::None; + selection_params.has_bias = has_bias; + selection_params.is_channels_last = is_channels_last; + selection_params.common_use_split_k = use_split_k; + selection_params.split_dim_inner = context->GetSplitKConfig().GetSplitDimInner(); + + const MatMulExecutionPlan plan = cache.GetOrCreateScheduler(*context).CreateExecutionPlan( + selection_params, context->ForcedMatMulAlgorithm()); + const MatMulAlgorithm algorithm = plan.algorithm; + ORT_RETURN_IF_NOT(IsMatMulAlgorithmConfigurationCompatible(plan), + "MatMul algorithm ", MatMulAlgorithmName(algorithm), + " received an incompatible vendor configuration."); + const auto* packed_configuration = std::get_if(&plan.configuration); MatMulAlgorithmPrerequisites prerequisites{}; prerequisites.can_use_subgroup_matrix = can_use_subgroup_matrix; prerequisites.has_intel_subgroup_capability = has_intel_subgroup_capability; prerequisites.has_nonzero_k = helper.K() > 0; - prerequisites.split_k_configured = context->GetSplitKConfig().GetSplitDimInner() != 0; - prerequisites.deterministic_compute = context->KernelContext().GetUseDeterministicCompute(); - prerequisites.is_vec4 = helper.K() % 4 == 0 && helper.N() % 4 == 0; - prerequisites.has_fused_activation = activation.activation_kind_ != ActivationKind::None; + prerequisites.split_k_configured = + packed_configuration != nullptr && packed_configuration->split_dim_inner > 1; + prerequisites.deterministic_compute = selection_params.deterministic_compute; + prerequisites.is_vec4 = selection_params.is_vec4; + prerequisites.has_fused_activation = selection_params.has_fused_activation; prerequisites.split_k_bias_layout_supported = !has_bias || is_channels_last; ORT_RETURN_IF_NOT(MeetsMatMulAlgorithmPrerequisites(algorithm, prerequisites), "MatMul algorithm ", MatMulAlgorithmName(algorithm), @@ -439,10 +488,12 @@ Status ComputeMatMul(ComputeContext* context, case MatMulAlgorithm::Packed: return ApplyMatMulPacked( *context, activation, inputs, output_tensor, is_channels_last, helper, + *packed_configuration, /*use_split_k=*/false); case MatMulAlgorithm::PackedSplitK: return ApplyMatMulPacked( *context, activation, inputs, output_tensor, is_channels_last, helper, + *packed_configuration, /*use_split_k=*/true); } diff --git a/onnxruntime/core/providers/webgpu/math/matmul_algorithm_scheduler.h b/onnxruntime/core/providers/webgpu/math/matmul_algorithm_scheduler.h index 269ee589354bf..4c3dbecad9561 100644 --- a/onnxruntime/core/providers/webgpu/math/matmul_algorithm_scheduler.h +++ b/onnxruntime/core/providers/webgpu/math/matmul_algorithm_scheduler.h @@ -3,8 +3,13 @@ #pragma once +#include #include +#include #include +#include +#include +#include #include "core/providers/webgpu/math/matmul_algorithm.h" @@ -15,11 +20,98 @@ struct MatMulAlgorithmSelectionParams { int64_t m = 0; int64_t n = 0; int64_t k = 0; + int64_t packed_m = 0; + uint64_t batch_size = 1; + uint64_t packed_batch_size = 1; + std::string_view adapter_architecture; + int32_t a_data_type = 0; + int32_t b_data_type = 0; bool can_use_subgroup_matrix = false; bool has_intel_subgroup_capability = false; - bool use_split_k = false; + bool is_vec4 = false; + bool deterministic_compute = false; + bool has_fused_activation = false; + bool has_bias = false; + bool is_channels_last = true; + bool common_use_split_k = false; + uint32_t split_dim_inner = 0; +}; + +struct MatMulSubgroupMatrixConfiguration {}; +struct MatMulNaiveConfiguration {}; +struct MatMulIntelSubgroupConfiguration {}; + +struct MatMulPackedConfiguration { + std::array workgroup_size{8, 8, 1}; + std::array elements_per_thread{4, 4, 1}; + uint32_t tile_inner = 32; + uint32_t split_dim_inner = 1; +}; + +inline bool IsMatMulPackedConfigurationValid( + const MatMulPackedConfiguration& configuration, + bool use_split_k) { + if (configuration.workgroup_size[0] == 0 || + configuration.workgroup_size[1] == 0 || + configuration.workgroup_size[2] != 1 || + configuration.elements_per_thread[0] == 0 || + configuration.elements_per_thread[1] == 0 || + configuration.elements_per_thread[2] != 1 || + configuration.tile_inner == 0) { + return false; + } + + return !use_split_k || + (configuration.split_dim_inner > 1 && + configuration.split_dim_inner % configuration.tile_inner == 0); +} + +inline std::optional TryGetMatMulPackedDispatchGroupCount( + uint64_t dimension, + uint32_t workgroup_size, + uint32_t elements_per_thread) { + if (workgroup_size == 0 || elements_per_thread == 0) { + return std::nullopt; + } + + const uint64_t elements_per_group = + static_cast(workgroup_size) * elements_per_thread; + const uint64_t group_count = + dimension == 0 ? 0 : 1 + (dimension - 1) / elements_per_group; + if (group_count > std::numeric_limits::max()) { + return std::nullopt; + } + + return static_cast(group_count); +} + +using MatMulAlgorithmConfiguration = + std::variant; + +struct MatMulExecutionPlan { + MatMulAlgorithm algorithm; + MatMulAlgorithmConfiguration configuration; }; +inline bool IsMatMulAlgorithmConfigurationCompatible( + const MatMulExecutionPlan& plan) { + switch (plan.algorithm) { + case MatMulAlgorithm::SubgroupMatrix: + return std::holds_alternative(plan.configuration); + case MatMulAlgorithm::Naive: + return std::holds_alternative(plan.configuration); + case MatMulAlgorithm::IntelSubgroup: + return std::holds_alternative(plan.configuration); + case MatMulAlgorithm::Packed: + case MatMulAlgorithm::PackedSplitK: + return std::holds_alternative(plan.configuration); + } + return false; +} + struct MatMulAlgorithmPrerequisites { bool can_use_subgroup_matrix = false; bool has_intel_subgroup_capability = false; @@ -65,19 +157,25 @@ class MatMulAlgorithmScheduler { if (forced_algorithm.has_value()) { return *forced_algorithm; } - if (params.can_use_subgroup_matrix) { - return MatMulAlgorithm::SubgroupMatrix; - } - if (params.k == 0 || (params.n < 8 && params.k < 8)) { + if (params.k == 0) { return MatMulAlgorithm::Naive; } if (const auto vendor_algorithm = SelectVendorAlgorithm(params); vendor_algorithm.has_value()) { return *vendor_algorithm; } - if (params.use_split_k) { - return MatMulAlgorithm::PackedSplitK; + return SelectCommonAlgorithm(params); + } + + MatMulExecutionPlan CreateExecutionPlan( + const MatMulAlgorithmSelectionParams& params, + std::optional forced_algorithm = std::nullopt) const { + const MatMulAlgorithm algorithm = Select(params, forced_algorithm); + std::optional configuration = + SelectVendorConfiguration(algorithm, params); + if (!configuration.has_value()) { + configuration = SelectCommonConfiguration(algorithm, params); } - return MatMulAlgorithm::Packed; + return MatMulExecutionPlan{algorithm, std::move(*configuration)}; } protected: @@ -85,6 +183,50 @@ class MatMulAlgorithmScheduler { const MatMulAlgorithmSelectionParams& /*params*/) const { return std::nullopt; } + + virtual std::optional SelectVendorConfiguration( + MatMulAlgorithm /*algorithm*/, + const MatMulAlgorithmSelectionParams& /*params*/) const { + return std::nullopt; + } + + private: + MatMulAlgorithm SelectCommonAlgorithm(const MatMulAlgorithmSelectionParams& params) const { + if (params.can_use_subgroup_matrix) { + return MatMulAlgorithm::SubgroupMatrix; + } + if (params.n < 8 && params.k < 8) { + return MatMulAlgorithm::Naive; + } + if (params.common_use_split_k) { + return MatMulAlgorithm::PackedSplitK; + } + return MatMulAlgorithm::Packed; + } + + MatMulAlgorithmConfiguration SelectCommonConfiguration( + MatMulAlgorithm algorithm, + const MatMulAlgorithmSelectionParams& params) const { + switch (algorithm) { + case MatMulAlgorithm::SubgroupMatrix: + return MatMulSubgroupMatrixConfiguration{}; + case MatMulAlgorithm::Naive: + return MatMulNaiveConfiguration{}; + case MatMulAlgorithm::IntelSubgroup: + return MatMulIntelSubgroupConfiguration{}; + case MatMulAlgorithm::Packed: + case MatMulAlgorithm::PackedSplitK: { + MatMulPackedConfiguration configuration{}; + configuration.elements_per_thread = + params.packed_m <= 8 ? std::array{4, 1, 1} + : std::array{4, 4, 1}; + configuration.split_dim_inner = + algorithm == MatMulAlgorithm::PackedSplitK ? params.split_dim_inner : 1; + return configuration; + } + } + return MatMulNaiveConfiguration{}; + } }; } // namespace webgpu diff --git a/onnxruntime/core/providers/webgpu/math/matmul_packed.cc b/onnxruntime/core/providers/webgpu/math/matmul_packed.cc index 6cd2f692a5cd3..59a40c6c3f12a 100644 --- a/onnxruntime/core/providers/webgpu/math/matmul_packed.cc +++ b/onnxruntime/core/providers/webgpu/math/matmul_packed.cc @@ -47,9 +47,12 @@ Status MatMulProgram::GenerateShaderCode(ShaderHelper& shader) const { ORT_RETURN_IF_ERROR(MakeMatMulPackedVec4Source( shader, elements_per_thread_, WorkgroupSizeX(), WorkgroupSizeY(), data_type, &batch_dims, /*transA = */ false, /*transB = */ false, /*alpha = */ 1.f, /*need_handle_matmul = */ true, - /*output_components = */ 4, /*tile_inner = */ 32, need_split_k, split_dim_inner_)); + /*output_components = */ 4, tile_inner_, need_split_k, split_dim_inner_)); } else { - ORT_RETURN_IF_ERROR(MakeMatMulPackedSource(shader, elements_per_thread_, WorkgroupSizeX(), WorkgroupSizeY(), data_type, &batch_dims)); + ORT_RETURN_IF_ERROR(MakeMatMulPackedSource( + shader, elements_per_thread_, WorkgroupSizeX(), WorkgroupSizeY(), data_type, &batch_dims, + /*transpose_a = */ false, /*transpose_b = */ false, /*alpha = */ 1.f, + /*need_handle_matmul = */ true, tile_inner_)); } return Status::OK(); } diff --git a/onnxruntime/core/providers/webgpu/math/matmul_packed.h b/onnxruntime/core/providers/webgpu/math/matmul_packed.h index 833305e9cddc8..e1be676298fca 100644 --- a/onnxruntime/core/providers/webgpu/math/matmul_packed.h +++ b/onnxruntime/core/providers/webgpu/math/matmul_packed.h @@ -13,13 +13,21 @@ namespace onnxruntime { namespace webgpu { class MatMulProgram final : public Program { public: - MatMulProgram(const Activation& activation, bool bias, bool is_vec4, const gsl::span& elements_per_thread, bool is_channels_last = false, uint32_t split_dim_inner = 1) : Program{"MatMul"}, - activation_(activation), - has_bias_{bias}, - is_vec4_{is_vec4}, - elements_per_thread_(elements_per_thread.begin(), elements_per_thread.end()), - is_channels_last_(is_channels_last), - split_dim_inner_(split_dim_inner) {} + MatMulProgram(const Activation& activation, + bool bias, + bool is_vec4, + const gsl::span& elements_per_thread, + bool is_channels_last = false, + uint32_t split_dim_inner = 1, + uint32_t tile_inner = 32) + : Program{"MatMul"}, + activation_(activation), + has_bias_{bias}, + is_vec4_{is_vec4}, + elements_per_thread_(elements_per_thread.begin(), elements_per_thread.end()), + is_channels_last_(is_channels_last), + split_dim_inner_(split_dim_inner), + tile_inner_(tile_inner) {} Status GenerateShaderCode(ShaderHelper& sh) const override; WEBGPU_PROGRAM_DEFINE_UNIFORM_VARIABLES({"dim_a_outer", ProgramUniformVariableDataType::Uint32}, @@ -40,6 +48,7 @@ class MatMulProgram final : public Program { const InlinedVector elements_per_thread_; bool is_channels_last_ = false; uint32_t split_dim_inner_ = 1; + uint32_t tile_inner_ = 32; }; // The program to initialize the output with 0 or bias before doing MatMul with Split-K. In Split-K, diff --git a/onnxruntime/core/providers/webgpu/vendor/intel/math/matmul_algorithm_scheduler.h b/onnxruntime/core/providers/webgpu/vendor/intel/math/matmul_algorithm_scheduler.h index 37bd97332fcb2..123c05c75033a 100644 --- a/onnxruntime/core/providers/webgpu/vendor/intel/math/matmul_algorithm_scheduler.h +++ b/onnxruntime/core/providers/webgpu/vendor/intel/math/matmul_algorithm_scheduler.h @@ -13,6 +13,9 @@ class IntelMatMulAlgorithmScheduler final : public MatMulAlgorithmScheduler { protected: std::optional SelectVendorAlgorithm( const MatMulAlgorithmSelectionParams& params) const override { + if (params.can_use_subgroup_matrix) { + return MatMulAlgorithm::SubgroupMatrix; + } if (params.has_intel_subgroup_capability && params.m >= 64 && params.n >= 512 && params.k >= 32) { return MatMulAlgorithm::IntelSubgroup; diff --git a/onnxruntime/test/providers/webgpu/matmul_algorithm_scheduler_test.cc b/onnxruntime/test/providers/webgpu/matmul_algorithm_scheduler_test.cc index b052d4c69efe8..e2b408d60ce74 100644 --- a/onnxruntime/test/providers/webgpu/matmul_algorithm_scheduler_test.cc +++ b/onnxruntime/test/providers/webgpu/matmul_algorithm_scheduler_test.cc @@ -1,6 +1,8 @@ // Copyright (c) Microsoft Corporation. All rights reserved. // Licensed under the MIT License. +#include + #include "gtest/gtest.h" #include "core/providers/webgpu/math/matmul_algorithm.h" @@ -12,6 +14,59 @@ namespace onnxruntime { namespace webgpu { namespace test { +namespace { + +class AlwaysPackedVendorScheduler final : public MatMulAlgorithmScheduler { + protected: + std::optional SelectVendorAlgorithm( + const MatMulAlgorithmSelectionParams& /*params*/) const override { + return MatMulAlgorithm::Packed; + } +}; + +class TunedPackedVendorScheduler final : public MatMulAlgorithmScheduler { + protected: + std::optional SelectVendorAlgorithm( + const MatMulAlgorithmSelectionParams& /*params*/) const override { + return MatMulAlgorithm::Naive; + } + + std::optional SelectVendorConfiguration( + MatMulAlgorithm algorithm, + const MatMulAlgorithmSelectionParams& params) const override { + const bool is_packed_algorithm = algorithm == MatMulAlgorithm::Packed || + algorithm == MatMulAlgorithm::PackedSplitK; + if (!is_packed_algorithm || + params.adapter_architecture != "test-architecture" || + params.a_data_type != 10 || params.b_data_type != 10) { + return std::nullopt; + } + + MatMulPackedConfiguration configuration{}; + configuration.workgroup_size = {16, 4, 1}; + configuration.elements_per_thread = {4, 2, 1}; + configuration.tile_inner = 16; + configuration.split_dim_inner = 128; + return configuration; + } +}; + +class VendorSplitKThresholdScheduler final : public MatMulAlgorithmScheduler { + protected: + std::optional SelectVendorAlgorithm( + const MatMulAlgorithmSelectionParams& params) const override { + if (params.batch_size <= 16 && params.is_vec4 && + !params.deterministic_compute && !params.has_fused_activation && + (!params.has_bias || params.is_channels_last) && + params.split_dim_inner != 0) { + return MatMulAlgorithm::PackedSplitK; + } + return std::nullopt; + } +}; + +} // namespace + TEST(MatMulAlgorithmParsingTest, RoundTripsEveryAlgorithmName) { struct TestCase { std::string_view name; @@ -38,24 +93,143 @@ TEST(MatMulAlgorithmParsingTest, RejectsUnknownAlgorithmName) { } TEST(MatMulAlgorithmSchedulerTest, ForcedAlgorithmTakesPrecedence) { - MatMulAlgorithmScheduler scheduler; + AlwaysPackedVendorScheduler scheduler; MatMulAlgorithmSelectionParams params{}; params.can_use_subgroup_matrix = true; - EXPECT_EQ(scheduler.Select(params, MatMulAlgorithm::Packed), MatMulAlgorithm::Packed); + EXPECT_EQ(scheduler.Select(params, MatMulAlgorithm::Naive), MatMulAlgorithm::Naive); +} + +TEST(MatMulAlgorithmSchedulerTest, VendorCanTuneForcedAlgorithmConfiguration) { + TunedPackedVendorScheduler scheduler; + MatMulAlgorithmSelectionParams params{}; + params.adapter_architecture = "test-architecture"; + params.a_data_type = 10; + params.b_data_type = 10; + + const MatMulExecutionPlan plan = + scheduler.CreateExecutionPlan(params, MatMulAlgorithm::Packed); + + EXPECT_EQ(plan.algorithm, MatMulAlgorithm::Packed); + const auto& configuration = + std::get(plan.configuration); + EXPECT_EQ(configuration.workgroup_size, + (std::array{16, 4, 1})); + EXPECT_EQ(configuration.elements_per_thread, + (std::array{4, 2, 1})); + EXPECT_EQ(configuration.tile_inner, 16u); + EXPECT_EQ(configuration.split_dim_inner, 128u); +} + +TEST(MatMulAlgorithmSchedulerTest, VendorCanSetIndependentSplitKThresholds) { + VendorSplitKThresholdScheduler scheduler; + MatMulAlgorithmSelectionParams params{}; + params.m = 32; + params.n = 64; + params.k = 1024; + params.batch_size = 16; + params.is_vec4 = true; + params.is_channels_last = true; + params.split_dim_inner = 128; + params.common_use_split_k = false; + + EXPECT_EQ(scheduler.Select(params), MatMulAlgorithm::PackedSplitK); + + params.batch_size = 17; + EXPECT_EQ(scheduler.Select(params), MatMulAlgorithm::Packed); } -TEST(MatMulAlgorithmSchedulerTest, SubgroupMatrixTakesAutomaticPrecedence) { +TEST(MatMulAlgorithmSchedulerTest, CommonPackedConfigurationPreservesCurrentTuning) { + MatMulAlgorithmScheduler scheduler; + MatMulAlgorithmSelectionParams params{}; + params.m = 8; + params.packed_m = 8; + params.n = 64; + params.k = 64; + + const MatMulExecutionPlan small_m_plan = scheduler.CreateExecutionPlan(params); + const auto& small_m_configuration = + std::get(small_m_plan.configuration); + EXPECT_EQ(small_m_configuration.workgroup_size, + (std::array{8, 8, 1})); + EXPECT_EQ(small_m_configuration.elements_per_thread, + (std::array{4, 1, 1})); + EXPECT_EQ(small_m_configuration.tile_inner, 32u); + + params.m = 9; + params.packed_m = 9; + const MatMulExecutionPlan large_m_plan = scheduler.CreateExecutionPlan(params); + const auto& large_m_configuration = + std::get(large_m_plan.configuration); + EXPECT_EQ(large_m_configuration.elements_per_thread, + (std::array{4, 4, 1})); +} + +TEST(MatMulAlgorithmConfigurationTest, PackedConfigurationKeepsBatchAxesUntiled) { + MatMulPackedConfiguration configuration{}; + EXPECT_TRUE(IsMatMulPackedConfigurationValid(configuration, /*use_split_k=*/false)); + + configuration.workgroup_size[2] = 2; + EXPECT_FALSE(IsMatMulPackedConfigurationValid(configuration, /*use_split_k=*/false)); + + configuration.workgroup_size[2] = 1; + configuration.elements_per_thread[2] = 2; + EXPECT_FALSE(IsMatMulPackedConfigurationValid(configuration, /*use_split_k=*/false)); +} + +TEST(MatMulAlgorithmConfigurationTest, SplitKConfigurationRequiresTileAlignedSplits) { + MatMulPackedConfiguration configuration{}; + configuration.tile_inner = 32; + configuration.split_dim_inner = 64; + EXPECT_TRUE(IsMatMulPackedConfigurationValid(configuration, /*use_split_k=*/true)); + + configuration.split_dim_inner = 48; + EXPECT_FALSE(IsMatMulPackedConfigurationValid(configuration, /*use_split_k=*/true)); +} + +TEST(MatMulAlgorithmConfigurationTest, PackedDispatchArithmeticIsOverflowSafe) { + const auto maximum_tuning_dispatch = TryGetMatMulPackedDispatchGroupCount( + std::numeric_limits::max(), + std::numeric_limits::max(), + std::numeric_limits::max()); + ASSERT_TRUE(maximum_tuning_dispatch.has_value()); + EXPECT_EQ(*maximum_tuning_dispatch, 1u); + + EXPECT_EQ(TryGetMatMulPackedDispatchGroupCount( + std::numeric_limits::max(), 1, 1), + std::nullopt); +} + +TEST(MatMulAlgorithmSchedulerTest, CommonFallbackPrefersSubgroupMatrix) { MatMulAlgorithmScheduler scheduler; MatMulAlgorithmSelectionParams params{}; params.n = 4; params.k = 4; params.can_use_subgroup_matrix = true; - params.use_split_k = true; + params.common_use_split_k = true; EXPECT_EQ(scheduler.Select(params), MatMulAlgorithm::SubgroupMatrix); } +TEST(MatMulAlgorithmSchedulerTest, VendorPolicyPrecedesCommonHeuristics) { + AlwaysPackedVendorScheduler scheduler; + MatMulAlgorithmSelectionParams params{}; + params.n = 4; + params.k = 4; + params.can_use_subgroup_matrix = true; + + EXPECT_EQ(scheduler.Select(params), MatMulAlgorithm::Packed); +} + +TEST(MatMulAlgorithmSchedulerTest, ZeroContractionDimensionPrecedesVendorPolicy) { + AlwaysPackedVendorScheduler scheduler; + MatMulAlgorithmSelectionParams params{}; + params.n = 8; + params.k = 0; + + EXPECT_EQ(scheduler.Select(params), MatMulAlgorithm::Naive); +} + TEST(MatMulAlgorithmSchedulerTest, NaiveUsesStrictSmallDimensionBoundaries) { MatMulAlgorithmScheduler scheduler; @@ -96,15 +270,27 @@ TEST(MatMulAlgorithmSchedulerTest, IntelSchedulerAppliesCurrentVendorRule) { EXPECT_EQ(scheduler.Select(params), MatMulAlgorithm::Packed); } +TEST(MatMulAlgorithmSchedulerTest, IntelSchedulerPreservesSubgroupMatrixPrecedence) { + intel::IntelMatMulAlgorithmScheduler scheduler; + MatMulAlgorithmSelectionParams params{}; + params.m = 64; + params.n = 512; + params.k = 32; + params.can_use_subgroup_matrix = true; + params.has_intel_subgroup_capability = true; + + EXPECT_EQ(scheduler.Select(params), MatMulAlgorithm::SubgroupMatrix); +} + TEST(MatMulAlgorithmSchedulerTest, SplitKPrecedesPackedFallback) { MatMulAlgorithmScheduler scheduler; MatMulAlgorithmSelectionParams params{}; params.n = 64; params.k = 1024; - params.use_split_k = true; + params.common_use_split_k = true; EXPECT_EQ(scheduler.Select(params), MatMulAlgorithm::PackedSplitK); - params.use_split_k = false; + params.common_use_split_k = false; EXPECT_EQ(scheduler.Select(params), MatMulAlgorithm::Packed); } From 2d529c6fb61eb494a85a1ad4f456282fb93e5ec0 Mon Sep 17 00:00:00 2001 From: Jiajia Qin Date: Mon, 21 Sep 2026 17:22:25 +0800 Subject: [PATCH 18/25] Address WebGPU MatMul review feedback --- .../core/providers/webgpu/math/matmul_algorithm_scheduler.h | 3 +++ .../providers/webgpu/matmul_algorithm_scheduler_test.cc | 6 ++++++ onnxruntime/test/providers/webgpu/matmul_large_test.cc | 5 +++++ 3 files changed, 14 insertions(+) diff --git a/onnxruntime/core/providers/webgpu/math/matmul_algorithm_scheduler.h b/onnxruntime/core/providers/webgpu/math/matmul_algorithm_scheduler.h index 4c3dbecad9561..de55e92d1267d 100644 --- a/onnxruntime/core/providers/webgpu/math/matmul_algorithm_scheduler.h +++ b/onnxruntime/core/providers/webgpu/math/matmul_algorithm_scheduler.h @@ -11,6 +11,7 @@ #include #include +#include "core/common/common.h" #include "core/providers/webgpu/math/matmul_algorithm.h" namespace onnxruntime { @@ -149,7 +150,9 @@ inline bool MeetsMatMulAlgorithmPrerequisites( class MatMulAlgorithmScheduler { public: + MatMulAlgorithmScheduler() = default; virtual ~MatMulAlgorithmScheduler() = default; + ORT_DISALLOW_COPY_ASSIGNMENT_AND_MOVE(MatMulAlgorithmScheduler); MatMulAlgorithm Select( const MatMulAlgorithmSelectionParams& params, diff --git a/onnxruntime/test/providers/webgpu/matmul_algorithm_scheduler_test.cc b/onnxruntime/test/providers/webgpu/matmul_algorithm_scheduler_test.cc index e2b408d60ce74..f5a3277086364 100644 --- a/onnxruntime/test/providers/webgpu/matmul_algorithm_scheduler_test.cc +++ b/onnxruntime/test/providers/webgpu/matmul_algorithm_scheduler_test.cc @@ -2,6 +2,7 @@ // Licensed under the MIT License. #include +#include #include "gtest/gtest.h" @@ -16,6 +17,11 @@ namespace test { namespace { +static_assert(!std::is_copy_constructible_v); +static_assert(!std::is_copy_assignable_v); +static_assert(!std::is_move_constructible_v); +static_assert(!std::is_move_assignable_v); + class AlwaysPackedVendorScheduler final : public MatMulAlgorithmScheduler { protected: std::optional SelectVendorAlgorithm( diff --git a/onnxruntime/test/providers/webgpu/matmul_large_test.cc b/onnxruntime/test/providers/webgpu/matmul_large_test.cc index 0e5eb0a96982a..3e6eea4a9b1e6 100644 --- a/onnxruntime/test/providers/webgpu/matmul_large_test.cc +++ b/onnxruntime/test/providers/webgpu/matmul_large_test.cc @@ -231,6 +231,11 @@ TEST(MatMulProgramTest, VectorFallbackExecution) { #if defined(_WIN32) TEST(WebGpuMatMulAlgorithmTest, RejectsUnknownForcedAlgorithm) { + ConfigOptions valid_config_options{}; + if (!WebGpuExecutionProviderWithOptions(valid_config_options)) { + GTEST_SKIP() << "WebGPU execution provider is unavailable in this build."; + } + ConfigOptions config_options{}; ASSERT_STATUS_OK(config_options.AddConfigEntry(webgpu::options::kForceMatMulAlgorithm, "unknown")); EXPECT_THROW(WebGpuExecutionProviderWithOptions(config_options), OnnxRuntimeException); From 65d4725e7fa029236a7c79e6ec1a70bdb8b46700 Mon Sep 17 00:00:00 2001 From: Jiajia Qin Date: Mon, 21 Sep 2026 23:00:10 +0800 Subject: [PATCH 19/25] Refactor WebGPU Split-K vendor policy --- .../webgpu_matmul_algorithm_scheduler.md | 6 +- .../core/providers/webgpu/math/matmul.cc | 44 +-------- .../webgpu/math/matmul_algorithm_scheduler.h | 30 +++++-- .../intel/math/matmul_algorithm_scheduler.h | 7 ++ .../vendor/intel/math/split_k_config.cc | 44 +++++++++ .../webgpu/vendor/intel/math/split_k_config.h | 18 ++++ .../core/providers/webgpu/webgpu_context.cc | 2 +- .../core/providers/webgpu/webgpu_utils.cc | 82 +++++++---------- .../core/providers/webgpu/webgpu_utils.h | 18 +++- .../webgpu/matmul_algorithm_scheduler_test.cc | 89 ++++++++++++++++--- 10 files changed, 225 insertions(+), 115 deletions(-) create mode 100644 onnxruntime/core/providers/webgpu/vendor/intel/math/split_k_config.cc create mode 100644 onnxruntime/core/providers/webgpu/vendor/intel/math/split_k_config.h diff --git a/docs/design/webgpu_matmul_algorithm_scheduler.md b/docs/design/webgpu_matmul_algorithm_scheduler.md index 09c7040fdd679..6852b719b7d64 100644 --- a/docs/design/webgpu_matmul_algorithm_scheduler.md +++ b/docs/design/webgpu_matmul_algorithm_scheduler.md @@ -26,7 +26,9 @@ Add a `MatMulAlgorithmScheduler` base class. Automatic selection uses this order An Intel-derived scheduler implements the vendor hook. It preserves the original policy by selecting `SubgroupMatrix` first when applicable, then selecting `IntelSubgroup` under the current Intel subgroup rule. Other vendors use the base scheduler unchanged. Future vendor policies can derive from the base scheduler, override as many performance ranges as needed, and return no selection to delegate the remaining ranges to the common fallback without adding vendor conditionals to `ComputeMatMul`. -The scheduler accepts immutable problem facts rather than mutable policy decisions: logical and packed dimensions, batch sizes, adapter architecture, input types, packing and layout facts, deterministic-compute state, activation and bias facts, device capabilities, and the configured Split-K size. The common `SplitKConfig::UseSplitK` result is passed separately as the common fallback recommendation. A vendor may ignore that recommendation and apply independent thresholds from the raw facts. This keeps the scheduler deterministic and unit-testable without a WebGPU device. +The scheduler accepts immutable problem facts rather than mutable policy decisions: logical and packed dimensions, batch sizes, adapter architecture, input types, packing and layout facts, deterministic-compute state, activation and bias facts, and device capabilities. It also owns a copy of the immutable `SplitKConfig` selected for the adapter. The private common fallback evaluates `SplitKConfig::UseSplitK` directly from the problem facts instead of receiving a precomputed decision from `ComputeMatMul`. A vendor may ignore that common recommendation and apply independent thresholds from the raw facts. This keeps the scheduler deterministic and unit-testable without a WebGPU device. + +`SplitKConfig` contains only generic Split-K eligibility evaluation and data. Adapter routing is handled by a small generic factory, while Intel architecture profiles and their measured threshold tables live under `vendor/intel`. `WebGpuContext` owns the selected configuration so GEMM and MatMul use the same profile; the MatMul scheduler receives that configuration when it is created. A future vendor can add its own profile builder and factory route without adding conditions to `ComputeMatMul` or changing the generic evaluator. ## Execution Configuration @@ -64,7 +66,7 @@ The option is intentionally internal and test-only: it is declared with WebGPU p ## Testing -- Add device-independent scheduler unit tests covering every common branch, forced-over-vendor precedence, the zero-K correctness guard, vendor-over-common precedence, Intel override, default fallback, independent vendor thresholds, common packed defaults, and vendor tuning of a forced algorithm. +- Add device-independent scheduler unit tests covering every common branch, forced-over-vendor precedence, the zero-K correctness guard, vendor-over-common precedence, Intel override, default fallback, independent vendor thresholds, common packed defaults, vendor tuning of a forced algorithm, Split-K profile routing, and current Intel architecture boundaries. - Add parser/configuration tests for every accepted value and invalid input. - Add WebGPU MatMul tests that choose shapes which normally select a different path, force a compatible algorithm, and verify numerical output. Hardware-specific forced algorithms are tested only when their hard capabilities are present; strict-failure tests cover unsupported forced choices. - Build Dawn and the WebGPU provider on Windows with the Vulkan backend enabled and D3D12 disabled. Tests explicitly request `dawnBackendType=Vulkan`. diff --git a/onnxruntime/core/providers/webgpu/math/matmul.cc b/onnxruntime/core/providers/webgpu/math/matmul.cc index a32b5197df49d..538eca1a24d8a 100644 --- a/onnxruntime/core/providers/webgpu/math/matmul.cc +++ b/onnxruntime/core/providers/webgpu/math/matmul.cc @@ -30,9 +30,10 @@ MatMulOptImpl* MatMulOptImplCache::GetOrCreate(const ComputeContextBase& context const MatMulAlgorithmScheduler& MatMulOptImplCache::GetOrCreateScheduler(const ComputeContextBase& context) { std::call_once(scheduler_init_flag_, [&]() { if (context.AdapterInfo().vendor == std::string_view{"intel"}) { - scheduler_ = std::make_unique(); + scheduler_ = std::make_unique( + context.GetSplitKConfig()); } else { - scheduler_ = std::make_unique(); + scheduler_ = std::make_unique(context.GetSplitKConfig()); } }); return *scheduler_; @@ -225,40 +226,6 @@ static Status ApplyMatMulNaive(ComputeContext& context, return context.RunProgram(program); } -static bool ShouldUsePackedSplitK(ComputeContext& context, - const Activation& activation, - const std::vector& inputs, - bool is_channels_last, - const MatMulComputeHelper& helper) { - if (context.KernelContext().GetUseDeterministicCompute()) { - return false; - } - - TensorShape a_shape = inputs[0]->Shape(); - TensorShape b_shape = inputs[1]->Shape(); - TensorShape output_shape = helper.OutputShape(); - const int64_t batch_a = - a_shape.NumDimensions() > 2 ? a_shape.SizeToDimension(a_shape.NumDimensions() - 2) : 1; - const int64_t batch_b = - b_shape.NumDimensions() > 2 ? b_shape.SizeToDimension(b_shape.NumDimensions() - 2) : 1; - if (batch_a != 1 && batch_b == 1) { - const int64_t batch_and_m = a_shape.SizeToDimension(a_shape.NumDimensions() - 1); - a_shape = TensorShape({batch_and_m, helper.K()}); - b_shape = TensorShape({helper.K(), helper.N()}); - output_shape = TensorShape({batch_and_m, helper.N()}); - } - - const int64_t batch_size = output_shape.NumDimensions() > 2 - ? output_shape.SizeToDimension(output_shape.NumDimensions() - 2) - : 1; - const uint32_t m = narrow(a_shape[a_shape.NumDimensions() - 2]); - const uint32_t k = narrow(a_shape[a_shape.NumDimensions() - 1]); - const uint32_t n = narrow(b_shape[b_shape.NumDimensions() - 1]); - const bool is_vec4 = k % 4 == 0 && n % 4 == 0; - return context.GetSplitKConfig().UseSplitK( - is_vec4, activation.activation_kind_, batch_size, m, n, k, is_channels_last); -} - static Status ApplyMatMulPacked(ComputeContext& context, const Activation& activation, const std::vector& inputs, @@ -411,9 +378,6 @@ Status ComputeMatMul(ComputeContext* context, subgroup_impl != nullptr && subgroup_impl->CanApply(*context, inputs, is_channels_last, b_is_constant); const bool has_intel_subgroup_capability = intel::HasMatMulIntelCapability(*context); - const bool use_split_k = - ShouldUsePackedSplitK(*context, activation, inputs, is_channels_last, helper); - const int64_t batch_a = logical_a_shape.NumDimensions() > 2 ? logical_a_shape.SizeToDimension(logical_a_shape.NumDimensions() - 2) @@ -448,8 +412,6 @@ Status ComputeMatMul(ComputeContext* context, selection_params.has_fused_activation = activation.activation_kind_ != ActivationKind::None; selection_params.has_bias = has_bias; selection_params.is_channels_last = is_channels_last; - selection_params.common_use_split_k = use_split_k; - selection_params.split_dim_inner = context->GetSplitKConfig().GetSplitDimInner(); const MatMulExecutionPlan plan = cache.GetOrCreateScheduler(*context).CreateExecutionPlan( selection_params, context->ForcedMatMulAlgorithm()); diff --git a/onnxruntime/core/providers/webgpu/math/matmul_algorithm_scheduler.h b/onnxruntime/core/providers/webgpu/math/matmul_algorithm_scheduler.h index de55e92d1267d..303d097418386 100644 --- a/onnxruntime/core/providers/webgpu/math/matmul_algorithm_scheduler.h +++ b/onnxruntime/core/providers/webgpu/math/matmul_algorithm_scheduler.h @@ -13,6 +13,7 @@ #include "core/common/common.h" #include "core/providers/webgpu/math/matmul_algorithm.h" +#include "core/providers/webgpu/webgpu_utils.h" namespace onnxruntime { namespace webgpu { @@ -34,8 +35,6 @@ struct MatMulAlgorithmSelectionParams { bool has_fused_activation = false; bool has_bias = false; bool is_channels_last = true; - bool common_use_split_k = false; - uint32_t split_dim_inner = 0; }; struct MatMulSubgroupMatrixConfiguration {}; @@ -150,7 +149,8 @@ inline bool MeetsMatMulAlgorithmPrerequisites( class MatMulAlgorithmScheduler { public: - MatMulAlgorithmScheduler() = default; + explicit MatMulAlgorithmScheduler(SplitKConfig split_k_config = {}) + : split_k_config_{std::move(split_k_config)} {} virtual ~MatMulAlgorithmScheduler() = default; ORT_DISALLOW_COPY_ASSIGNMENT_AND_MOVE(MatMulAlgorithmScheduler); @@ -201,7 +201,7 @@ class MatMulAlgorithmScheduler { if (params.n < 8 && params.k < 8) { return MatMulAlgorithm::Naive; } - if (params.common_use_split_k) { + if (ShouldUseSplitK(params)) { return MatMulAlgorithm::PackedSplitK; } return MatMulAlgorithm::Packed; @@ -224,12 +224,32 @@ class MatMulAlgorithmScheduler { params.packed_m <= 8 ? std::array{4, 1, 1} : std::array{4, 4, 1}; configuration.split_dim_inner = - algorithm == MatMulAlgorithm::PackedSplitK ? params.split_dim_inner : 1; + algorithm == MatMulAlgorithm::PackedSplitK + ? split_k_config_.GetSplitDimInner() + : 1; return configuration; } } return MatMulNaiveConfiguration{}; } + + bool ShouldUseSplitK(const MatMulAlgorithmSelectionParams& params) const { + if (params.deterministic_compute || params.has_fused_activation || + params.packed_m < 0 || params.n < 0 || params.k < 0) { + return false; + } + + return split_k_config_.UseSplitK( + params.is_vec4, + ActivationKind::None, + params.packed_batch_size, + static_cast(params.packed_m), + static_cast(params.n), + static_cast(params.k), + params.is_channels_last); + } + + SplitKConfig split_k_config_; }; } // namespace webgpu diff --git a/onnxruntime/core/providers/webgpu/vendor/intel/math/matmul_algorithm_scheduler.h b/onnxruntime/core/providers/webgpu/vendor/intel/math/matmul_algorithm_scheduler.h index 123c05c75033a..745e620c901a8 100644 --- a/onnxruntime/core/providers/webgpu/vendor/intel/math/matmul_algorithm_scheduler.h +++ b/onnxruntime/core/providers/webgpu/vendor/intel/math/matmul_algorithm_scheduler.h @@ -3,6 +3,8 @@ #pragma once +#include + #include "core/providers/webgpu/math/matmul_algorithm_scheduler.h" namespace onnxruntime { @@ -10,6 +12,11 @@ namespace webgpu { namespace intel { class IntelMatMulAlgorithmScheduler final : public MatMulAlgorithmScheduler { + public: + IntelMatMulAlgorithmScheduler() = default; + explicit IntelMatMulAlgorithmScheduler(SplitKConfig split_k_config) + : MatMulAlgorithmScheduler{std::move(split_k_config)} {} + protected: std::optional SelectVendorAlgorithm( const MatMulAlgorithmSelectionParams& params) const override { diff --git a/onnxruntime/core/providers/webgpu/vendor/intel/math/split_k_config.cc b/onnxruntime/core/providers/webgpu/vendor/intel/math/split_k_config.cc new file mode 100644 index 0000000000000..653a0fee09f68 --- /dev/null +++ b/onnxruntime/core/providers/webgpu/vendor/intel/math/split_k_config.cc @@ -0,0 +1,44 @@ +// Copyright (c) Microsoft Corporation. All rights reserved. +// Licensed under the MIT License. + +#include "core/providers/webgpu/vendor/intel/math/split_k_config.h" + +namespace onnxruntime { +namespace webgpu { +namespace intel { + +SplitKConfig CreateSplitKConfig(std::string_view architecture) { + // Disable Split-K on old Intel GPUs. + if (architecture == "gen-7" || architecture == "gen-8" || + architecture == "gen-9" || architecture == "gen-11") { + return {}; + } + + constexpr uint32_t max_batch_size = 8; + constexpr uint32_t split_dim_inner = 256; + constexpr uint32_t min_dim_inner_with_split_k = split_dim_inner * 2; + + if (architecture == "xe-2lpg" || architecture == "xe-2hpg" || + architecture == "gen-12hp") { + // These thresholds are verified on Intel discrete GPUs and Lunar Lake iGPUs. + return SplitKConfig{ + max_batch_size, split_dim_inner, min_dim_inner_with_split_k, + {{768, 52.0}, {2304, 35.0}, {3072, 21.5}, {4096, 16.0}}}; + } + + if (architecture == "xe-3lpg") { + // These thresholds are verified on Intel Panther Lake iGPUs (12Xe). + return SplitKConfig{ + max_batch_size, split_dim_inner, min_dim_inner_with_split_k, + {{768, 40.0}, {1792, 22.0}, {3072, 18.0}, {4096, 10.0}}}; + } + + // Default thresholds for newer Intel GPUs, chosen on a gen-12lp GPU with 32 EUs. + return SplitKConfig{ + max_batch_size, split_dim_inner, min_dim_inner_with_split_k, + {{768, 20.0}, {1792, 13.0}, {3072, 8.0}, {4096, 6.0}}}; +} + +} // namespace intel +} // namespace webgpu +} // namespace onnxruntime diff --git a/onnxruntime/core/providers/webgpu/vendor/intel/math/split_k_config.h b/onnxruntime/core/providers/webgpu/vendor/intel/math/split_k_config.h new file mode 100644 index 0000000000000..c2b826126222c --- /dev/null +++ b/onnxruntime/core/providers/webgpu/vendor/intel/math/split_k_config.h @@ -0,0 +1,18 @@ +// Copyright (c) Microsoft Corporation. All rights reserved. +// Licensed under the MIT License. + +#pragma once + +#include + +#include "core/providers/webgpu/webgpu_utils.h" + +namespace onnxruntime { +namespace webgpu { +namespace intel { + +SplitKConfig CreateSplitKConfig(std::string_view architecture); + +} // namespace intel +} // namespace webgpu +} // namespace onnxruntime diff --git a/onnxruntime/core/providers/webgpu/webgpu_context.cc b/onnxruntime/core/providers/webgpu/webgpu_context.cc index 38f0844136591..9ca0e070d22fc 100644 --- a/onnxruntime/core/providers/webgpu/webgpu_context.cc +++ b/onnxruntime/core/providers/webgpu/webgpu_context.cc @@ -256,7 +256,7 @@ void WebGpuContext::Initialize(const WebGpuContextConfig& config) { program_mgr_ = std::make_unique(*this); // create split-k config - split_k_config_ = std::make_unique(adapter_info_); + split_k_config_ = std::make_unique(CreateSplitKConfig(adapter_info_)); // set query type #if !defined(__wasm__) diff --git a/onnxruntime/core/providers/webgpu/webgpu_utils.cc b/onnxruntime/core/providers/webgpu/webgpu_utils.cc index 62d7b9934e164..a22b93a0b3cab 100644 --- a/onnxruntime/core/providers/webgpu/webgpu_utils.cc +++ b/onnxruntime/core/providers/webgpu/webgpu_utils.cc @@ -3,7 +3,10 @@ #include "core/providers/webgpu/webgpu_utils.h" #include +#include + #include "core/providers/webgpu/shader_variable.h" +#include "core/providers/webgpu/vendor/intel/math/split_k_config.h" namespace onnxruntime { namespace webgpu { @@ -25,55 +28,32 @@ TensorShape ReduceShapeByComponents(const TensorShape& shape, int64_t components return TensorShape(shape_vector); } -SplitKConfig::SplitKConfig(const wgpu::AdapterInfo& adapter_info) { - if (adapter_info.vendor == std::string_view{"intel"}) { - // Disable Split-K on old Intel GPUs. - if (adapter_info.architecture == std::string_view{"gen-7"} || - adapter_info.architecture == std::string_view{"gen-8"} || - adapter_info.architecture == std::string_view{"gen-9"} || - adapter_info.architecture == std::string_view{"gen-11"}) { - enable_split_k_ = false; - } else if (adapter_info.architecture == std::string_view{"xe-2lpg"} || - adapter_info.architecture == std::string_view{"xe-2hpg"} || - adapter_info.architecture == std::string_view{"gen-12hp"}) { - // Below thresholds are only verified on Intel discrete GPUs and Lunar Lake iGPUs. - enable_split_k_ = true; - - max_batch_size_ = 8; - split_dim_inner_ = 256; - min_dim_inner_with_split_k_ = split_dim_inner_ * 2; - - configs_per_dim_inner_range_.emplace_back(768, 52.0); - configs_per_dim_inner_range_.emplace_back(2304, 35.0); - configs_per_dim_inner_range_.emplace_back(3072, 21.5); - configs_per_dim_inner_range_.emplace_back(4096, 16.0); - } else if (adapter_info.architecture == std::string_view{"xe-3lpg"}) { - // Below thresholds are only verified on Intel Panther Lake iGPUs (12Xe). - enable_split_k_ = true; - - max_batch_size_ = 8; - split_dim_inner_ = 256; - min_dim_inner_with_split_k_ = split_dim_inner_ * 2; - - configs_per_dim_inner_range_.emplace_back(768, 40.0); - configs_per_dim_inner_range_.emplace_back(1792, 22.0); - configs_per_dim_inner_range_.emplace_back(3072, 18.0); - configs_per_dim_inner_range_.emplace_back(4096, 10.0); - } else { - // Below are the default thresholds on newer Intel GPUs. These values are chosen on - // Intel "gen-12lp" GPU with 32EUs. - enable_split_k_ = true; - - max_batch_size_ = 8; - split_dim_inner_ = 256; - min_dim_inner_with_split_k_ = split_dim_inner_ * 2; - - configs_per_dim_inner_range_.emplace_back(768, 20.0); - configs_per_dim_inner_range_.emplace_back(1792, 13.0); - configs_per_dim_inner_range_.emplace_back(3072, 8.0); - configs_per_dim_inner_range_.emplace_back(4096, 6.0); - } +SplitKConfig::SplitKConfig( + uint32_t max_batch_size, + uint32_t split_dim_inner, + uint32_t min_dim_inner_with_split_k, + std::initializer_list> configs_per_dim_inner_range) + : enable_split_k_{true}, + split_dim_inner_{split_dim_inner}, + min_dim_inner_with_split_k_{min_dim_inner_with_split_k}, + max_batch_size_{max_batch_size} { + configs_per_dim_inner_range_.reserve(configs_per_dim_inner_range.size()); + for (const auto& [max_dim_inner, rate] : configs_per_dim_inner_range) { + configs_per_dim_inner_range_.emplace_back(max_dim_inner, rate); + } +} + +SplitKConfig CreateSplitKConfig(const wgpu::AdapterInfo& adapter_info) { + return CreateSplitKConfig( + std::string_view{adapter_info.vendor}, + std::string_view{adapter_info.architecture}); +} + +SplitKConfig CreateSplitKConfig(std::string_view vendor, std::string_view architecture) { + if (vendor == "intel") { + return intel::CreateSplitKConfig(architecture); } + return {}; } SplitKConfig::ConfigAtRange::ConfigAtRange(uint32_t max_dim_inner, double rate) @@ -88,9 +68,9 @@ bool SplitKConfig::UseSplitK( bool is_vec4, ActivationKind activation_kind, uint64_t batch_size, - uint32_t dim_a_outer, - uint32_t dim_b_outer, - uint32_t dim_inner, + uint64_t dim_a_outer, + uint64_t dim_b_outer, + uint64_t dim_inner, bool is_channels_last) const { if (!enable_split_k_) { return false; diff --git a/onnxruntime/core/providers/webgpu/webgpu_utils.h b/onnxruntime/core/providers/webgpu/webgpu_utils.h index d4bb245e3e9e8..587bdd9f45d83 100644 --- a/onnxruntime/core/providers/webgpu/webgpu_utils.h +++ b/onnxruntime/core/providers/webgpu/webgpu_utils.h @@ -4,6 +4,11 @@ #pragma once #include +#include +#include +#include +#include + #include "core/common/common.h" #include "core/framework/tensor.h" #include "core/framework/tensor_shape.h" @@ -103,11 +108,17 @@ inline Tensor CreateTensorView(const Tensor& tensor, MLDataType new_data_type, c */ class SplitKConfig { public: - explicit SplitKConfig(const wgpu::AdapterInfo& adapter_info); + SplitKConfig() = default; + SplitKConfig( + uint32_t max_batch_size, + uint32_t split_dim_inner, + uint32_t min_dim_inner_with_split_k, + std::initializer_list> configs_per_dim_inner_range); bool UseSplitK( bool is_vec4, ActivationKind activation_kind, uint64_t batch_size, - uint32_t dim_a_outer, uint32_t dim_b_outer, uint32_t dim_inner, bool is_channels_last = true) const; + uint64_t dim_a_outer, uint64_t dim_b_outer, uint64_t dim_inner, + bool is_channels_last = true) const; uint32_t GetSplitDimInner() const; @@ -127,6 +138,9 @@ class SplitKConfig { std::vector configs_per_dim_inner_range_; }; +SplitKConfig CreateSplitKConfig(const wgpu::AdapterInfo& adapter_info); +SplitKConfig CreateSplitKConfig(std::string_view vendor, std::string_view architecture); + /** * Generates WGSL (WebGPU Shading Language) code for performing an atomic add operation * on a non-integer value (e.g., floating-point) in a shader. diff --git a/onnxruntime/test/providers/webgpu/matmul_algorithm_scheduler_test.cc b/onnxruntime/test/providers/webgpu/matmul_algorithm_scheduler_test.cc index f5a3277086364..04e0ced45f869 100644 --- a/onnxruntime/test/providers/webgpu/matmul_algorithm_scheduler_test.cc +++ b/onnxruntime/test/providers/webgpu/matmul_algorithm_scheduler_test.cc @@ -10,6 +10,7 @@ #include "core/providers/webgpu/math/matmul_algorithm_scheduler.h" #include "core/providers/webgpu/vendor/intel/math/gemm_subgroup_utils.h" #include "core/providers/webgpu/vendor/intel/math/matmul_algorithm_scheduler.h" +#include "core/providers/webgpu/vendor/intel/math/split_k_config.h" namespace onnxruntime { namespace webgpu { @@ -58,13 +59,20 @@ class TunedPackedVendorScheduler final : public MatMulAlgorithmScheduler { }; class VendorSplitKThresholdScheduler final : public MatMulAlgorithmScheduler { + public: + VendorSplitKThresholdScheduler() + : MatMulAlgorithmScheduler{SplitKConfig{ + /*max_batch_size=*/8, + /*split_dim_inner=*/128, + /*min_dim_inner_with_split_k=*/256, + {{4096, 16.0}}}} {} + protected: std::optional SelectVendorAlgorithm( const MatMulAlgorithmSelectionParams& params) const override { if (params.batch_size <= 16 && params.is_vec4 && !params.deterministic_compute && !params.has_fused_activation && - (!params.has_bias || params.is_channels_last) && - params.split_dim_inner != 0) { + (!params.has_bias || params.is_channels_last)) { return MatMulAlgorithm::PackedSplitK; } return std::nullopt; @@ -98,6 +106,38 @@ TEST(MatMulAlgorithmParsingTest, RejectsUnknownAlgorithmName) { EXPECT_EQ(ParseMatMulAlgorithm("unknown"), std::nullopt); } +TEST(SplitKConfigTest, IntelArchitectureProfilesPreserveCurrentBoundaries) { + const SplitKConfig discrete_config = intel::CreateSplitKConfig("xe-2lpg"); + EXPECT_EQ(discrete_config.GetSplitDimInner(), 256u); + EXPECT_TRUE(discrete_config.UseSplitK( + /*is_vec4=*/true, ActivationKind::None, /*batch_size=*/1, + /*dim_a_outer=*/192, /*dim_b_outer=*/160, /*dim_inner=*/1024)); + + const SplitKConfig xe3_config = intel::CreateSplitKConfig("xe-3lpg"); + EXPECT_FALSE(xe3_config.UseSplitK( + /*is_vec4=*/true, ActivationKind::None, /*batch_size=*/1, + /*dim_a_outer=*/192, /*dim_b_outer=*/160, /*dim_inner=*/1024)); + EXPECT_TRUE(xe3_config.UseSplitK( + /*is_vec4=*/true, ActivationKind::None, /*batch_size=*/1, + /*dim_a_outer=*/128, /*dim_b_outer=*/128, /*dim_inner=*/1024)); + + const SplitKConfig default_config = intel::CreateSplitKConfig("gen-12lp"); + EXPECT_FALSE(default_config.UseSplitK( + /*is_vec4=*/true, ActivationKind::None, /*batch_size=*/1, + /*dim_a_outer=*/128, /*dim_b_outer=*/128, /*dim_inner=*/1024)); + + const SplitKConfig legacy_config = intel::CreateSplitKConfig("gen-9"); + EXPECT_EQ(legacy_config.GetSplitDimInner(), 0u); + EXPECT_FALSE(legacy_config.UseSplitK( + /*is_vec4=*/true, ActivationKind::None, /*batch_size=*/1, + /*dim_a_outer=*/1, /*dim_b_outer=*/1, /*dim_inner=*/1024)); +} + +TEST(SplitKConfigTest, FactoryRoutesOnlySupportedVendorProfiles) { + EXPECT_EQ(CreateSplitKConfig("intel", "xe-2lpg").GetSplitDimInner(), 256u); + EXPECT_EQ(CreateSplitKConfig("nvidia", "pascal").GetSplitDimInner(), 0u); +} + TEST(MatMulAlgorithmSchedulerTest, ForcedAlgorithmTakesPrecedence) { AlwaysPackedVendorScheduler scheduler; MatMulAlgorithmSelectionParams params{}; @@ -134,17 +174,36 @@ TEST(MatMulAlgorithmSchedulerTest, VendorCanSetIndependentSplitKThresholds) { params.n = 64; params.k = 1024; params.batch_size = 16; + params.packed_batch_size = 16; params.is_vec4 = true; params.is_channels_last = true; - params.split_dim_inner = 128; - params.common_use_split_k = false; - EXPECT_EQ(scheduler.Select(params), MatMulAlgorithm::PackedSplitK); + const MatMulExecutionPlan plan = scheduler.CreateExecutionPlan(params); + EXPECT_EQ(plan.algorithm, MatMulAlgorithm::PackedSplitK); + EXPECT_EQ(std::get(plan.configuration).split_dim_inner, 128u); params.batch_size = 17; + params.packed_batch_size = 17; EXPECT_EQ(scheduler.Select(params), MatMulAlgorithm::Packed); } +TEST(MatMulAlgorithmSchedulerTest, InjectedSplitKPolicyCreatesPackedSplitKPlan) { + intel::IntelMatMulAlgorithmScheduler scheduler{intel::CreateSplitKConfig("xe-2lpg")}; + MatMulAlgorithmSelectionParams params{}; + params.m = 192; + params.packed_m = 192; + params.n = 160; + params.k = 1024; + params.is_vec4 = true; + + const MatMulExecutionPlan plan = scheduler.CreateExecutionPlan(params); + + EXPECT_EQ(plan.algorithm, MatMulAlgorithm::PackedSplitK); + const auto& configuration = + std::get(plan.configuration); + EXPECT_EQ(configuration.split_dim_inner, 256u); +} + TEST(MatMulAlgorithmSchedulerTest, CommonPackedConfigurationPreservesCurrentTuning) { MatMulAlgorithmScheduler scheduler; MatMulAlgorithmSelectionParams params{}; @@ -207,12 +266,14 @@ TEST(MatMulAlgorithmConfigurationTest, PackedDispatchArithmeticIsOverflowSafe) { } TEST(MatMulAlgorithmSchedulerTest, CommonFallbackPrefersSubgroupMatrix) { - MatMulAlgorithmScheduler scheduler; + MatMulAlgorithmScheduler scheduler{intel::CreateSplitKConfig("xe-2lpg")}; MatMulAlgorithmSelectionParams params{}; - params.n = 4; - params.k = 4; + params.m = 64; + params.packed_m = 64; + params.n = 64; + params.k = 1024; params.can_use_subgroup_matrix = true; - params.common_use_split_k = true; + params.is_vec4 = true; EXPECT_EQ(scheduler.Select(params), MatMulAlgorithm::SubgroupMatrix); } @@ -289,15 +350,17 @@ TEST(MatMulAlgorithmSchedulerTest, IntelSchedulerPreservesSubgroupMatrixPreceden } TEST(MatMulAlgorithmSchedulerTest, SplitKPrecedesPackedFallback) { - MatMulAlgorithmScheduler scheduler; + MatMulAlgorithmScheduler scheduler{intel::CreateSplitKConfig("xe-2lpg")}; MatMulAlgorithmSelectionParams params{}; + params.m = 64; + params.packed_m = 64; params.n = 64; params.k = 1024; - params.common_use_split_k = true; + params.is_vec4 = true; EXPECT_EQ(scheduler.Select(params), MatMulAlgorithm::PackedSplitK); - params.common_use_split_k = false; - EXPECT_EQ(scheduler.Select(params), MatMulAlgorithm::Packed); + MatMulAlgorithmScheduler disabled_scheduler; + EXPECT_EQ(disabled_scheduler.Select(params), MatMulAlgorithm::Packed); } TEST(MatMulAlgorithmPrerequisiteTest, SplitKRejectsEachHardConstraint) { From c9b6a60fc8fa418c28192a512d52dfdc15807b7f Mon Sep 17 00:00:00 2001 From: Jiajia Qin Date: Tue, 22 Sep 2026 10:42:32 +0800 Subject: [PATCH 20/25] Format Intel WebGPU Split-K profiles --- .../providers/webgpu/vendor/intel/math/split_k_config.cc | 9 +++------ 1 file changed, 3 insertions(+), 6 deletions(-) diff --git a/onnxruntime/core/providers/webgpu/vendor/intel/math/split_k_config.cc b/onnxruntime/core/providers/webgpu/vendor/intel/math/split_k_config.cc index 653a0fee09f68..4c5f8f02e45a5 100644 --- a/onnxruntime/core/providers/webgpu/vendor/intel/math/split_k_config.cc +++ b/onnxruntime/core/providers/webgpu/vendor/intel/math/split_k_config.cc @@ -22,21 +22,18 @@ SplitKConfig CreateSplitKConfig(std::string_view architecture) { architecture == "gen-12hp") { // These thresholds are verified on Intel discrete GPUs and Lunar Lake iGPUs. return SplitKConfig{ - max_batch_size, split_dim_inner, min_dim_inner_with_split_k, - {{768, 52.0}, {2304, 35.0}, {3072, 21.5}, {4096, 16.0}}}; + max_batch_size, split_dim_inner, min_dim_inner_with_split_k, {{768, 52.0}, {2304, 35.0}, {3072, 21.5}, {4096, 16.0}}}; } if (architecture == "xe-3lpg") { // These thresholds are verified on Intel Panther Lake iGPUs (12Xe). return SplitKConfig{ - max_batch_size, split_dim_inner, min_dim_inner_with_split_k, - {{768, 40.0}, {1792, 22.0}, {3072, 18.0}, {4096, 10.0}}}; + max_batch_size, split_dim_inner, min_dim_inner_with_split_k, {{768, 40.0}, {1792, 22.0}, {3072, 18.0}, {4096, 10.0}}}; } // Default thresholds for newer Intel GPUs, chosen on a gen-12lp GPU with 32 EUs. return SplitKConfig{ - max_batch_size, split_dim_inner, min_dim_inner_with_split_k, - {{768, 20.0}, {1792, 13.0}, {3072, 8.0}, {4096, 6.0}}}; + max_batch_size, split_dim_inner, min_dim_inner_with_split_k, {{768, 20.0}, {1792, 13.0}, {3072, 8.0}, {4096, 6.0}}}; } } // namespace intel From 4ccbbbc7b228394fbdc610349cfb2500cd5b7bf3 Mon Sep 17 00:00:00 2001 From: Jiajia Qin Date: Tue, 22 Sep 2026 15:58:02 +0800 Subject: [PATCH 21/25] docs: clarify WebGPU MatMul dispatch ownership --- .../design/webgpu_matmul_algorithm_scheduler.md | 17 +++++++++++++---- 1 file changed, 13 insertions(+), 4 deletions(-) diff --git a/docs/design/webgpu_matmul_algorithm_scheduler.md b/docs/design/webgpu_matmul_algorithm_scheduler.md index 6852b719b7d64..9df605a253e7f 100644 --- a/docs/design/webgpu_matmul_algorithm_scheduler.md +++ b/docs/design/webgpu_matmul_algorithm_scheduler.md @@ -52,15 +52,24 @@ Invalid option strings fail during WebGPU provider creation and list accepted va ## Dispatch and Implementation Boundaries -`ComputeMatMul` computes shared shape facts once, asks the scheduler for one execution plan, validates its hard prerequisites and configuration type, and switches directly to the matching implementation. Algorithm bodies are extracted into focused helpers where necessary. The generic packed helper takes an explicit Split-K mode and packed configuration; it does not re-run selection or tuning heuristics. +Introduce `MatMulComputeDispatcher` as the single compute entry point below the MatMul, pointwise Conv, and contrib Attention kernels. Each kernel owns one dispatcher for its lifetime. The dispatcher lazily creates adapter-dependent state from the first compute context and then owns: -The subgroup-matrix optional implementation gains a non-mutating applicability query and an execution method that no longer communicates selection through a `handled` output. This removes trial execution as a dispatch mechanism. +- one `MatMulAlgorithmScheduler`, which contains selection policy and immutable adapter tuning data; and +- an optional `SubgroupMatrixMatMulImpl`, which contains only the subgroup-matrix implementation's persistent device state, including cached padded constant weights. -The existing per-kernel cache continues to own device-dependent subgroup-matrix state and the scheduler, so MatMul, pointwise Conv, and Attention callers retain their current caching and behavior. +The dispatcher does not store a current or previously selected algorithm. Selection occurs for every invocation because shapes, input properties, activation, and bias can differ between calls. The session's forced-test configuration, when present, participates in each selection without becoming mutable dispatcher state. The resulting `MatMulExecutionPlan` is local to that invocation. + +For each call, `MatMulComputeDispatcher::Compute` computes shared shape and capability facts once, asks the scheduler for one execution plan, validates the selected algorithm's hard prerequisites and configuration type, and switches directly to the matching implementation. MatMul, pointwise Conv, and contrib Attention delegate through this same entry point rather than coordinating the scheduler and implementations themselves. + +The scheduler remains pure policy: it creates plans but does not create device programs, cache tensor data, or execute kernels. `SubgroupMatrixMatMulImpl` is named for the implementation it owns and is reached only when the plan selects `SubgroupMatrix`; its applicability query is non-mutating, and its execution method does not communicate selection through a `handled` output. This removes trial execution as a dispatch mechanism. + +The other implementations remain focused stateless functions or program builders unless they acquire persistent state in the future. The naive, Intel subgroup, packed, and packed Split-K paths therefore do not receive empty polymorphic wrapper classes. The generic packed helper takes an explicit Split-K mode and packed configuration; it does not re-run selection or tuning heuristics. WebGPU's existing program cache continues to own reusable compiled programs. + +This replaces `MatMulOptImplCache` and the generic `MatMulOptImpl` interface. If another algorithm later needs persistent implementation state, the dispatcher can own a separately named backend for that algorithm without changing scheduler policy or pretending that one object represents whichever algorithm happened to be selected most recently. ## Compatibility -With no forcing option or vendor override, the common scheduler remains equivalent to the previous selection order. Existing call sites keep using `ComputeMatMul`; the refactor does not change the operator API or model semantics. +With no forcing option or vendor override, the common scheduler remains equivalent to the previous selection order. The MatMul, pointwise Conv, and contrib Attention operator APIs and model semantics do not change; only their internal MatMul compute ownership moves behind the dispatcher. The option is intentionally internal and test-only: it is declared with WebGPU provider options for configuration plumbing but is not added to public user documentation. From c87348540de88b364b7aa4d593892cfbac3d8b9b Mon Sep 17 00:00:00 2001 From: Jiajia Qin Date: Tue, 22 Sep 2026 17:42:16 +0800 Subject: [PATCH 22/25] Refactor WebGPU MatMul compute dispatch --- .../contrib_ops/webgpu/bert/attention.cc | 11 +- .../contrib_ops/webgpu/bert/attention.h | 4 +- .../core/providers/webgpu/math/matmul.cc | 54 ++++----- .../core/providers/webgpu/math/matmul.h | 43 +------ .../webgpu/math/matmul_compute_dispatcher.h | 62 ++++++++++ .../webgpu/math/subgroup_matrix_matmul.cc | 8 +- onnxruntime/core/providers/webgpu/nn/conv.cc | 5 +- onnxruntime/core/providers/webgpu/nn/conv.h | 4 +- .../webgpu/matmul_algorithm_scheduler_test.cc | 22 ++++ .../providers/webgpu/matmul_large_test.cc | 107 +++++++++++++++++- 10 files changed, 232 insertions(+), 88 deletions(-) create mode 100644 onnxruntime/core/providers/webgpu/math/matmul_compute_dispatcher.h diff --git a/onnxruntime/contrib_ops/webgpu/bert/attention.cc b/onnxruntime/contrib_ops/webgpu/bert/attention.cc index decd943884548..63e2d0c4cdfa2 100644 --- a/onnxruntime/contrib_ops/webgpu/bert/attention.cc +++ b/onnxruntime/contrib_ops/webgpu/bert/attention.cc @@ -9,7 +9,6 @@ #include "contrib_ops/webgpu/webgpu_contrib_kernels.h" #include "core/providers/webgpu/webgpu_supported_types.h" #include "core/providers/webgpu/webgpu_utils.h" -#include "core/providers/webgpu/math/matmul.h" using namespace onnxruntime::webgpu; using namespace ::onnxruntime::common; using namespace ONNX_NAMESPACE; @@ -649,7 +648,7 @@ Attention::Attention(const OpKernelInfo& info) Status PrepareQKV(onnxruntime::webgpu::ComputeContext& context, const WebgpuAttentionParameters& parameters, const Tensor* input, const Tensor* weights, const Tensor* bias, Tensor* q, Tensor* k, Tensor* v, - MatMulOptImplCache& matmul_compute_cache, bool weights_are_constant) { + MatMulComputeDispatcher& matmul_compute_dispatcher, bool weights_are_constant) { // Use MatMul to compute packed QKV output: input * weights + bias // Then use SplitPackedQKV to split into Q, K, V in BSD format // Returns Q, K, V in BSD format @@ -663,9 +662,9 @@ Status PrepareQKV(onnxruntime::webgpu::ComputeContext& context, const WebgpuAtte std::vector matmul_inputs = {input, weights, bias}; // Call MatMul: packed_qkv = input * weights + bias - ORT_RETURN_IF_ERROR(onnxruntime::webgpu::ComputeMatMul( - &context, Activation(), matmul_inputs, &packed_qkv, /*is_channels_last=*/true, - matmul_compute_cache, weights_are_constant)); + ORT_RETURN_IF_ERROR(matmul_compute_dispatcher.Compute( + context, Activation(), matmul_inputs, &packed_qkv, /*is_channels_last=*/true, + weights_are_constant)); // Output Q, K, V in BSD format return SplitPackedQKV(context, parameters, &packed_qkv, q, k, v, parameters.hidden_size_); @@ -733,7 +732,7 @@ Status Attention::ComputeInternal(onnxruntime::webgpu::ComputeContext& context) // Compute Q, K, V from input, weights, and bias (returns BSD format) ORT_RETURN_IF_ERROR(PrepareQKV(context, parameters, input, weights, bias, &Q_bsd, &K_bsd, &V_bsd, - matmul_compute_cache_, weights_are_constant_)); + matmul_compute_dispatcher_, weights_are_constant_)); parameters.qkv_format_ = Q_K_V_BSNH; // Check if we can use flash attention diff --git a/onnxruntime/contrib_ops/webgpu/bert/attention.h b/onnxruntime/contrib_ops/webgpu/bert/attention.h index 7b5cf8da4a85e..a8c33e8d33bea 100644 --- a/onnxruntime/contrib_ops/webgpu/bert/attention.h +++ b/onnxruntime/contrib_ops/webgpu/bert/attention.h @@ -4,7 +4,7 @@ #pragma once #include "core/providers/webgpu/compute_context.h" -#include "core/providers/webgpu/math/matmul.h" +#include "core/providers/webgpu/math/matmul_compute_dispatcher.h" #include "core/providers/webgpu/program.h" #include "core/providers/webgpu/shader_helper.h" #include "core/providers/webgpu/webgpu_kernel.h" @@ -148,7 +148,7 @@ class Attention final : public WebGpuKernel, public onnxruntime::contrib::Attent Status ComputeInternal(onnxruntime::webgpu::ComputeContext& context) const override; private: - mutable MatMulOptImplCache matmul_compute_cache_; + mutable MatMulComputeDispatcher matmul_compute_dispatcher_; bool weights_are_constant_ = false; }; diff --git a/onnxruntime/core/providers/webgpu/math/matmul.cc b/onnxruntime/core/providers/webgpu/math/matmul.cc index 538eca1a24d8a..d8d7602f3c999 100644 --- a/onnxruntime/core/providers/webgpu/math/matmul.cc +++ b/onnxruntime/core/providers/webgpu/math/matmul.cc @@ -18,17 +18,11 @@ namespace onnxruntime { namespace webgpu { -std::unique_ptr CreateSubgroupMatrixMatMulImpl(const ComputeContextBase& context); +std::unique_ptr CreateSubgroupMatrixMatMulImpl(const ComputeContextBase& context); -MatMulOptImpl* MatMulOptImplCache::GetOrCreate(const ComputeContextBase& context) { - std::call_once(subgroup_impl_init_flag_, [&]() { - subgroup_impl_ = CreateSubgroupMatrixMatMulImpl(context); - }); - return subgroup_impl_.get(); -} - -const MatMulAlgorithmScheduler& MatMulOptImplCache::GetOrCreateScheduler(const ComputeContextBase& context) { - std::call_once(scheduler_init_flag_, [&]() { +void MatMulComputeDispatcher::Initialize(const ComputeContextBase& context) { + std::call_once(init_flag_, [&]() { + subgroup_matrix_impl_ = CreateSubgroupMatrixMatMulImpl(context); if (context.AdapterInfo().vendor == std::string_view{"intel"}) { scheduler_ = std::make_unique( context.GetSplitKConfig()); @@ -36,7 +30,6 @@ const MatMulAlgorithmScheduler& MatMulOptImplCache::GetOrCreateScheduler(const C scheduler_ = std::make_unique(context.GetSplitKConfig()); } }); - return *scheduler_; } ONNX_OPERATOR_VERSIONED_KERNEL_EX( @@ -171,8 +164,8 @@ Status MatMul::ComputeInternal(ComputeContext& context) const { inputs[1] = &promoted_b; } - return ComputeMatMul(&context, Activation(), inputs, output_tensor, - /*is_channels_last=*/true, compute_cache_, b_is_constant_); + return compute_dispatcher_.Compute(context, Activation(), inputs, output_tensor, + /*is_channels_last=*/true, b_is_constant_); } static Status ApplyMatMulNaive(ComputeContext& context, @@ -358,10 +351,12 @@ static Status ApplyMatMulPacked(ComputeContext& context, return context.RunProgram(program); } -Status ComputeMatMul(ComputeContext* context, - const Activation& activation, std::vector& inputs, Tensor* output_tensor, - bool is_channels_last, MatMulOptImplCache& cache, - bool b_is_constant) { +Status MatMulComputeDispatcher::Compute(ComputeContext& context, + const Activation& activation, + const std::vector& inputs, + Tensor* output_tensor, + bool is_channels_last, + bool b_is_constant) { const auto* a = inputs[0]; const auto* b = inputs[1]; const bool has_bias = inputs.size() > 2; @@ -373,11 +368,12 @@ Status ComputeMatMul(ComputeContext* context, MatMulComputeHelper helper; ORT_RETURN_IF_ERROR(helper.Compute(logical_a_shape, logical_b_shape)); - MatMulOptImpl* subgroup_impl = cache.GetOrCreate(*context); + Initialize(context); + SubgroupMatrixMatMulImpl* subgroup_impl = subgroup_matrix_impl_.get(); const bool can_use_subgroup_matrix = subgroup_impl != nullptr && - subgroup_impl->CanApply(*context, inputs, is_channels_last, b_is_constant); - const bool has_intel_subgroup_capability = intel::HasMatMulIntelCapability(*context); + subgroup_impl->CanApply(context, inputs, is_channels_last, b_is_constant); + const bool has_intel_subgroup_capability = intel::HasMatMulIntelCapability(context); const int64_t batch_a = logical_a_shape.NumDimensions() > 2 ? logical_a_shape.SizeToDimension(logical_a_shape.NumDimensions() - 2) @@ -402,19 +398,19 @@ Status ComputeMatMul(ComputeContext* context, : helper.M(); selection_params.batch_size = batch_size; selection_params.packed_batch_size = folds_batch_into_m ? 1 : batch_size; - selection_params.adapter_architecture = context->AdapterInfo().architecture; + selection_params.adapter_architecture = context.AdapterInfo().architecture; selection_params.a_data_type = a->GetElementType(); selection_params.b_data_type = b->GetElementType(); selection_params.can_use_subgroup_matrix = can_use_subgroup_matrix; selection_params.has_intel_subgroup_capability = has_intel_subgroup_capability; selection_params.is_vec4 = helper.K() % 4 == 0 && helper.N() % 4 == 0; - selection_params.deterministic_compute = context->KernelContext().GetUseDeterministicCompute(); + selection_params.deterministic_compute = context.KernelContext().GetUseDeterministicCompute(); selection_params.has_fused_activation = activation.activation_kind_ != ActivationKind::None; selection_params.has_bias = has_bias; selection_params.is_channels_last = is_channels_last; - const MatMulExecutionPlan plan = cache.GetOrCreateScheduler(*context).CreateExecutionPlan( - selection_params, context->ForcedMatMulAlgorithm()); + const MatMulExecutionPlan plan = scheduler_->CreateExecutionPlan( + selection_params, context.ForcedMatMulAlgorithm()); const MatMulAlgorithm algorithm = plan.algorithm; ORT_RETURN_IF_NOT(IsMatMulAlgorithmConfigurationCompatible(plan), "MatMul algorithm ", MatMulAlgorithmName(algorithm), @@ -440,21 +436,21 @@ Status ComputeMatMul(ComputeContext* context, ORT_RETURN_IF_NOT(subgroup_impl != nullptr, "MatMul algorithm subgroup_matrix is unavailable."); return subgroup_impl->Compute( - *context, inputs, output_tensor, activation, is_channels_last, b_is_constant); + context, inputs, output_tensor, activation, is_channels_last, b_is_constant); case MatMulAlgorithm::Naive: return ApplyMatMulNaive( - *context, activation, inputs, output_tensor, is_channels_last, helper); + context, activation, inputs, output_tensor, is_channels_last, helper); case MatMulAlgorithm::IntelSubgroup: return intel::ApplyMatMulIntel( - *context, activation, inputs, output_tensor, is_channels_last); + context, activation, inputs, output_tensor, is_channels_last); case MatMulAlgorithm::Packed: return ApplyMatMulPacked( - *context, activation, inputs, output_tensor, is_channels_last, helper, + context, activation, inputs, output_tensor, is_channels_last, helper, *packed_configuration, /*use_split_k=*/false); case MatMulAlgorithm::PackedSplitK: return ApplyMatMulPacked( - *context, activation, inputs, output_tensor, is_channels_last, helper, + context, activation, inputs, output_tensor, is_channels_last, helper, *packed_configuration, /*use_split_k=*/true); } diff --git a/onnxruntime/core/providers/webgpu/math/matmul.h b/onnxruntime/core/providers/webgpu/math/matmul.h index bc3a0ed17f211..86e8c3df0a915 100644 --- a/onnxruntime/core/providers/webgpu/math/matmul.h +++ b/onnxruntime/core/providers/webgpu/math/matmul.h @@ -3,13 +3,10 @@ #pragma once -#include -#include - #include "core/providers/webgpu/webgpu_kernel.h" #include "core/providers/webgpu/program.h" #include "core/providers/cpu/math/matmul_helper.h" -#include "core/providers/webgpu/math/matmul_algorithm_scheduler.h" +#include "core/providers/webgpu/math/matmul_compute_dispatcher.h" #include "core/providers/webgpu/math/matmul_utils.h" #include "core/providers/webgpu/math/matmul_packed.h" #include "core/providers/webgpu/webgpu_utils.h" @@ -18,42 +15,6 @@ namespace onnxruntime { namespace webgpu { -class MatMulOptImpl { - public: - virtual ~MatMulOptImpl() = default; - - virtual bool CanApply(const ComputeContext& context, - const std::vector& inputs, - bool is_channels_last, - bool b_is_constant) const = 0; - - virtual Status Compute(ComputeContext& context, - const std::vector& inputs, - Tensor* output, - const Activation& activation, - bool is_channels_last, - bool b_is_constant) = 0; -}; - -class MatMulOptImplCache { - public: - MatMulOptImplCache() = default; - ORT_DISALLOW_COPY_ASSIGNMENT_AND_MOVE(MatMulOptImplCache); - - MatMulOptImpl* GetOrCreate(const ComputeContextBase& context); - const MatMulAlgorithmScheduler& GetOrCreateScheduler(const ComputeContextBase& context); - - private: - std::once_flag subgroup_impl_init_flag_; - std::unique_ptr subgroup_impl_; - std::once_flag scheduler_init_flag_; - std::unique_ptr scheduler_; -}; - -Status ComputeMatMul(ComputeContext* context, const Activation& activation, std::vector& inputs, Tensor* output, - bool is_channels_last, MatMulOptImplCache& cache, - bool b_is_constant = false); - MatMulFillBiasOrZeroBeforeSplitKProgram CreateMatMulFillBiasOrZeroBeforeSplitKProgram( const Tensor* bias, Tensor* output, @@ -80,7 +41,7 @@ class MatMul final : public WebGpuKernel { constexpr static uint32_t MATMUL_PACKED_WORKGROUP_SIZE_Z = 1; private: - mutable MatMulOptImplCache compute_cache_; + mutable MatMulComputeDispatcher compute_dispatcher_; bool b_is_constant_ = false; }; diff --git a/onnxruntime/core/providers/webgpu/math/matmul_compute_dispatcher.h b/onnxruntime/core/providers/webgpu/math/matmul_compute_dispatcher.h new file mode 100644 index 0000000000000..45110df742633 --- /dev/null +++ b/onnxruntime/core/providers/webgpu/math/matmul_compute_dispatcher.h @@ -0,0 +1,62 @@ +// Copyright (c) Microsoft Corporation. All rights reserved. +// Licensed under the MIT License. + +#pragma once + +#include +#include +#include + +#include "core/common/common.h" +#include "core/common/status.h" +#include "core/providers/webgpu/math/matmul_algorithm_scheduler.h" + +namespace onnxruntime { + +class Tensor; + +namespace webgpu { + +struct Activation; +class ComputeContext; +class ComputeContextBase; + +class SubgroupMatrixMatMulImpl { + public: + virtual ~SubgroupMatrixMatMulImpl() = default; + + virtual bool CanApply(const ComputeContext& context, + const std::vector& inputs, + bool is_channels_last, + bool b_is_constant) const = 0; + + virtual Status Compute(ComputeContext& context, + const std::vector& inputs, + Tensor* output, + const Activation& activation, + bool is_channels_last, + bool b_is_constant) = 0; +}; + +class MatMulComputeDispatcher { + public: + MatMulComputeDispatcher() = default; + ORT_DISALLOW_COPY_ASSIGNMENT_AND_MOVE(MatMulComputeDispatcher); + + Status Compute(ComputeContext& context, + const Activation& activation, + const std::vector& inputs, + Tensor* output, + bool is_channels_last, + bool b_is_constant = false); + + private: + void Initialize(const ComputeContextBase& context); + + std::once_flag init_flag_; + std::unique_ptr subgroup_matrix_impl_; + std::unique_ptr scheduler_; +}; + +} // namespace webgpu +} // namespace onnxruntime diff --git a/onnxruntime/core/providers/webgpu/math/subgroup_matrix_matmul.cc b/onnxruntime/core/providers/webgpu/math/subgroup_matrix_matmul.cc index 607fa8c1defbb..818d91356b05c 100644 --- a/onnxruntime/core/providers/webgpu/math/subgroup_matrix_matmul.cc +++ b/onnxruntime/core/providers/webgpu/math/subgroup_matrix_matmul.cc @@ -76,9 +76,9 @@ class SubgroupMatrixMatMulPadBProgram final : public Program CreateSubgroupMatrixMatMulImpl(const ComputeContextBase& context) { +std::unique_ptr CreateSubgroupMatrixMatMulImpl(const ComputeContextBase& context) { // Only run on devices that report the fixed 8x16x16 F16 subgroup-matrix config // this kernel is implemented for. That config's adapters expose a 16-32 subgroup // size range, so the kernel's fixed 32 lanes per subgroup must be pinned with @@ -374,7 +374,7 @@ std::unique_ptr CreateSubgroupMatrixMatMulImpl(const ComputeConte if (!tiling_selector) { return nullptr; } - return std::make_unique(config_index, std::move(tiling_selector)); + return std::make_unique(config_index, std::move(tiling_selector)); } } // namespace webgpu diff --git a/onnxruntime/core/providers/webgpu/nn/conv.cc b/onnxruntime/core/providers/webgpu/nn/conv.cc index 4b7937e02f240..c522d5bdec6d6 100644 --- a/onnxruntime/core/providers/webgpu/nn/conv.cc +++ b/onnxruntime/core/providers/webgpu/nn/conv.cc @@ -9,7 +9,6 @@ #include "core/providers/webgpu/tensor/transpose.h" #include "core/providers/webgpu/nn/grouped_conv.h" #include "core/providers/webgpu/webgpu_utils.h" -#include "core/providers/webgpu/math/matmul.h" namespace onnxruntime { namespace webgpu { @@ -281,8 +280,8 @@ Status Conv::ComputeInternal(ComputeContext& context if (has_bias) { matmul_inputs.push_back(bias); } - return ComputeMatMul(&context, activation_, matmul_inputs, output, is_channels_last, - matmul_compute_cache_, matmul_b_is_constant); + return matmul_compute_dispatcher_.Compute(context, activation_, matmul_inputs, output, is_channels_last, + matmul_b_is_constant); } // Transpose weights when necessary Tensor transposed_kernel; diff --git a/onnxruntime/core/providers/webgpu/nn/conv.h b/onnxruntime/core/providers/webgpu/nn/conv.h index a64a206f1ef34..6684b7ca5259c 100644 --- a/onnxruntime/core/providers/webgpu/nn/conv.h +++ b/onnxruntime/core/providers/webgpu/nn/conv.h @@ -8,7 +8,7 @@ #include "core/providers/cpu/nn/conv_attributes.h" #include "core/providers/webgpu/program.h" #include "core/providers/webgpu/shader_helper.h" -#include "core/providers/webgpu/math/matmul.h" +#include "core/providers/webgpu/math/matmul_compute_dispatcher.h" #include "core/providers/webgpu/nn/fuse_utils.h" namespace onnxruntime { @@ -49,7 +49,7 @@ class Conv : public WebGpuKernel { // Layout of the tensor ComputeInternal ends up consuming -- `prepacked_kernel_` when it // is set, otherwise input 1. Stays `OIHW` while `prepacked_kernel_` is null. KernelLayout kernel_layout_{KernelLayout::OIHW}; - mutable MatMulOptImplCache matmul_compute_cache_; + mutable MatMulComputeDispatcher matmul_compute_dispatcher_; }; Status TransposeKernel(ComputeContext& context, const Tensor* kernel, const TensorShape& kernel_shape, Tensor* transposed_kernel, const InlinedVector& perm); diff --git a/onnxruntime/test/providers/webgpu/matmul_algorithm_scheduler_test.cc b/onnxruntime/test/providers/webgpu/matmul_algorithm_scheduler_test.cc index 04e0ced45f869..0ef6965860ad0 100644 --- a/onnxruntime/test/providers/webgpu/matmul_algorithm_scheduler_test.cc +++ b/onnxruntime/test/providers/webgpu/matmul_algorithm_scheduler_test.cc @@ -8,6 +8,7 @@ #include "core/providers/webgpu/math/matmul_algorithm.h" #include "core/providers/webgpu/math/matmul_algorithm_scheduler.h" +#include "core/providers/webgpu/math/matmul_compute_dispatcher.h" #include "core/providers/webgpu/vendor/intel/math/gemm_subgroup_utils.h" #include "core/providers/webgpu/vendor/intel/math/matmul_algorithm_scheduler.h" #include "core/providers/webgpu/vendor/intel/math/split_k_config.h" @@ -22,6 +23,10 @@ static_assert(!std::is_copy_constructible_v); static_assert(!std::is_copy_assignable_v); static_assert(!std::is_move_constructible_v); static_assert(!std::is_move_assignable_v); +static_assert(!std::is_copy_constructible_v); +static_assert(!std::is_copy_assignable_v); +static_assert(!std::is_move_constructible_v); +static_assert(!std::is_move_assignable_v); class AlwaysPackedVendorScheduler final : public MatMulAlgorithmScheduler { protected: @@ -146,6 +151,23 @@ TEST(MatMulAlgorithmSchedulerTest, ForcedAlgorithmTakesPrecedence) { EXPECT_EQ(scheduler.Select(params, MatMulAlgorithm::Naive), MatMulAlgorithm::Naive); } +TEST(MatMulAlgorithmSchedulerTest, ReevaluatesSelectionForEachRuntimeShape) { + MatMulAlgorithmScheduler scheduler; + MatMulAlgorithmSelectionParams params{}; + params.m = 4; + params.packed_m = 4; + params.n = 7; + params.k = 7; + + EXPECT_EQ(scheduler.CreateExecutionPlan(params).algorithm, MatMulAlgorithm::Naive); + + params.m = 64; + params.packed_m = 64; + params.n = 64; + params.k = 64; + EXPECT_EQ(scheduler.CreateExecutionPlan(params).algorithm, MatMulAlgorithm::Packed); +} + TEST(MatMulAlgorithmSchedulerTest, VendorCanTuneForcedAlgorithmConfiguration) { TunedPackedVendorScheduler scheduler; MatMulAlgorithmSelectionParams params{}; diff --git a/onnxruntime/test/providers/webgpu/matmul_large_test.cc b/onnxruntime/test/providers/webgpu/matmul_large_test.cc index 3e6eea4a9b1e6..d02be72b85a96 100644 --- a/onnxruntime/test/providers/webgpu/matmul_large_test.cc +++ b/onnxruntime/test/providers/webgpu/matmul_large_test.cc @@ -3,6 +3,7 @@ #include "gtest/gtest.h" +#include "core/graph/onnx_protobuf.h" #include "core/providers/cpu/math/matmul_helper.h" #include "core/providers/webgpu/math/matmul_algorithm.h" #if !defined(ORT_USE_EP_API_ADAPTERS) @@ -10,8 +11,12 @@ #include "core/providers/webgpu/webgpu_context.h" #endif #include "core/providers/webgpu/webgpu_provider_options.h" -#include "test/providers/provider_test_utils.h" #include "test/common/tensor_op_test_utils.h" +#include "test/providers/provider_test_utils.h" +#include "test/test_environment.h" +#include "test/unittest_util/framework_test_utils.h" +#include "test/util/include/asserts.h" +#include "test/util/include/inference_session_wrapper.h" #include "default_providers.h" namespace onnxruntime { @@ -242,6 +247,54 @@ TEST(WebGpuMatMulAlgorithmTest, RejectsUnknownForcedAlgorithm) { } #if defined(DAWN_ENABLE_VULKAN) +static std::string BuildDynamicMatMulModelBytes() { + ONNX_NAMESPACE::ModelProto model; + model.set_ir_version(ONNX_NAMESPACE::IR_VERSION); + auto* opset = model.add_opset_import(); + opset->set_domain(""); + opset->set_version(13); + + auto* graph = model.mutable_graph(); + graph->set_name("dynamic_matmul"); + + auto set_float_shape = [](ONNX_NAMESPACE::ValueInfoProto* value_info, + const char* first_dimension, + const char* second_dimension) { + auto* tensor_type = value_info->mutable_type()->mutable_tensor_type(); + tensor_type->set_elem_type(ONNX_NAMESPACE::TensorProto_DataType_FLOAT); + auto* shape = tensor_type->mutable_shape(); + shape->add_dim()->set_dim_param(first_dimension); + if (second_dimension != nullptr) { + shape->add_dim()->set_dim_param(second_dimension); + } else { + shape->add_dim()->set_dim_value(8); + } + }; + + auto* a = graph->add_input(); + a->set_name("A"); + set_float_shape(a, "M", "K"); + + auto* b = graph->add_input(); + b->set_name("B"); + set_float_shape(b, "K", nullptr); + + auto* y = graph->add_output(); + y->set_name("Y"); + set_float_shape(y, "M", nullptr); + + auto* node = graph->add_node(); + node->set_name("MatMul"); + node->set_op_type("MatMul"); + node->add_input("A"); + node->add_input("B"); + node->add_output("Y"); + + std::string bytes; + model.SerializeToString(&bytes); + return bytes; +} + TEST(WebGpuMatMulAlgorithmTest, ForcedNaive) { RunTestTyped({8, 8}, {8, 8}, false, webgpu::MatMulAlgorithm::Naive); } @@ -279,6 +332,58 @@ TEST(WebGpuMatMulAlgorithmTest, ForcedSubgroupMatrixRejectsFloatInputs) { OpTester::ExpectResult::kExpectFailure, "MatMul algorithm subgroup_matrix"); } + +TEST(WebGpuMatMulAlgorithmTest, ReselectsAlgorithmForEachDynamicShape) { + ConfigOptions config_options{}; + ASSERT_STATUS_OK(config_options.AddConfigEntry(webgpu::options::kDawnBackendType, + webgpu::options::kDawnBackendType_Vulkan)); + auto webgpu_ep = WebGpuExecutionProviderWithOptions(config_options); + if (!webgpu_ep) { + GTEST_SKIP() << "WebGPU execution provider is not available."; + } + + SessionOptions session_options; + session_options.session_logid = "WebGpuDynamicMatMulAlgorithmSelection"; + InferenceSessionWrapper session(session_options, GetEnvironment()); + ASSERT_STATUS_OK(session.RegisterExecutionProvider(std::move(webgpu_ep))); + + const std::string model_bytes = BuildDynamicMatMulModelBytes(); + ASSERT_STATUS_OK(session.Load(model_bytes.data(), static_cast(model_bytes.size()))); + ASSERT_STATUS_OK(session.Initialize()); + + const std::vector output_names{"Y"}; + auto run = [&](const TensorShape& a_shape, std::vector a_data, + const TensorShape& b_shape, std::vector b_data, + const TensorShape& expected_shape, float expected_value) { + OrtValue a_value; + OrtValue b_value; + CreateMLValue(a_shape.GetDims(), a_data.data(), OrtMemoryInfo(), &a_value); + CreateMLValue(b_shape.GetDims(), b_data.data(), OrtMemoryInfo(), &b_value); + + NameMLValMap feeds; + feeds.emplace("A", std::move(a_value)); + feeds.emplace("B", std::move(b_value)); + std::vector fetches; + ASSERT_STATUS_OK(session.Run(feeds, output_names, &fetches)); + ASSERT_EQ(fetches.size(), 1u); + const Tensor& output = fetches[0].Get(); + ASSERT_EQ(output.Shape(), expected_shape); + for (float value : output.DataAsSpan()) { + EXPECT_EQ(value, expected_value); + } + }; + + // The first invocation uses Packed. The second has K=0 and must reselect Naive on the same + // session and kernel; reusing the first execution plan would fail Packed's nonzero-K prerequisite. + run(TensorShape({8, 8}), std::vector(64, 1.0f), + TensorShape({8, 8}), std::vector(64, 1.0f), + TensorShape({8, 8}), 8.0f); + + const float unused_storage = 0.0f; + run(TensorShape({1, 0}), std::vector{unused_storage}, + TensorShape({0, 8}), std::vector{unused_storage}, + TensorShape({1, 8}), 0.0f); +} #endif // defined(DAWN_ENABLE_VULKAN) #endif From 7c796e2248bd0f6ab7bc030f14834cf6cf4433b4 Mon Sep 17 00:00:00 2001 From: Jiajia Qin Date: Mon, 28 Sep 2026 14:09:50 +0800 Subject: [PATCH 23/25] test: run WebGPU MatMul algorithms across backends --- docs/design/webgpu_matmul_algorithm_scheduler.md | 4 ++-- .../test/providers/webgpu/matmul_large_test.cc | 15 +-------------- 2 files changed, 3 insertions(+), 16 deletions(-) diff --git a/docs/design/webgpu_matmul_algorithm_scheduler.md b/docs/design/webgpu_matmul_algorithm_scheduler.md index 9df605a253e7f..ab6010c90f219 100644 --- a/docs/design/webgpu_matmul_algorithm_scheduler.md +++ b/docs/design/webgpu_matmul_algorithm_scheduler.md @@ -78,6 +78,6 @@ The option is intentionally internal and test-only: it is declared with WebGPU p - Add device-independent scheduler unit tests covering every common branch, forced-over-vendor precedence, the zero-K correctness guard, vendor-over-common precedence, Intel override, default fallback, independent vendor thresholds, common packed defaults, vendor tuning of a forced algorithm, Split-K profile routing, and current Intel architecture boundaries. - Add parser/configuration tests for every accepted value and invalid input. - Add WebGPU MatMul tests that choose shapes which normally select a different path, force a compatible algorithm, and verify numerical output. Hardware-specific forced algorithms are tested only when their hard capabilities are present; strict-failure tests cover unsupported forced choices. -- Build Dawn and the WebGPU provider on Windows with the Vulkan backend enabled and D3D12 disabled. Tests explicitly request `dawnBackendType=Vulkan`. -- Verify the selected Vulkan adapter exposes subgroup size control, f16, and the cooperative/subgroup-matrix configuration required by the 8x16x16 kernel before claiming subgroup-matrix execution coverage. +- Run hardware-backed tests on the build's default WebGPU backend. Hardware-specific tests inspect the selected adapter's capabilities and skip unsupported algorithms. +- Verify the selected adapter exposes subgroup size control, f16, and the cooperative/subgroup-matrix configuration required by the 8x16x16 kernel before claiming subgroup-matrix execution coverage. - Run scheduler/parser tests and hardware-backed MatMul tests on the local Intel Arc Vulkan adapter. macOS-arm64 Metal CI remains additional cross-backend coverage; lavapipe is not used because it cannot execute MatMul reliably. diff --git a/onnxruntime/test/providers/webgpu/matmul_large_test.cc b/onnxruntime/test/providers/webgpu/matmul_large_test.cc index 27254f95b4b43..58cc30722f8b6 100644 --- a/onnxruntime/test/providers/webgpu/matmul_large_test.cc +++ b/onnxruntime/test/providers/webgpu/matmul_large_test.cc @@ -89,7 +89,6 @@ static void ComputeExpectedResult(std::initializer_list a_dims, if (!b_is_vector) output_dims.push_back(N); } -#if defined(_WIN32) && defined(DAWN_ENABLE_VULKAN) static std::optional GetForcedAlgorithmUnsupportedReason( const IExecutionProvider& ep, webgpu::MatMulAlgorithm algorithm) { @@ -147,7 +146,6 @@ static std::optional GetForcedAlgorithmUnsupportedReason( return std::nullopt; #endif } -#endif template void RunTestTyped(std::initializer_list a_dims, std::initializer_list b_dims, @@ -161,8 +159,6 @@ void RunTestTyped(std::initializer_list a_dims, std::initializer_list a_dims, std::initializer_list a_vals(random.Gaussian(AsSpan(a_dims), 0.0f, 0.25f)); @@ -237,7 +231,6 @@ TEST(MatMulProgramTest, VectorFallbackExecution) { RunTestTyped({2, 2, 8}, {8}); } -#if defined(_WIN32) TEST(WebGpuMatMulAlgorithmTest, RejectsUnknownForcedAlgorithm) { ConfigOptions valid_config_options{}; if (!WebGpuExecutionProviderWithOptions(valid_config_options)) { @@ -249,7 +242,6 @@ TEST(WebGpuMatMulAlgorithmTest, RejectsUnknownForcedAlgorithm) { EXPECT_THROW(WebGpuExecutionProviderWithOptions(config_options), OnnxRuntimeException); } -#if defined(DAWN_ENABLE_VULKAN) static std::string BuildDynamicMatMulModelBytes() { ONNX_NAMESPACE::ModelProto model; model.set_ir_version(ONNX_NAMESPACE::IR_VERSION); @@ -337,10 +329,7 @@ TEST(WebGpuMatMulAlgorithmTest, ForcedSubgroupMatrixRejectsFloatInputs) { } TEST(WebGpuMatMulAlgorithmTest, ReselectsAlgorithmForEachDynamicShape) { - ConfigOptions config_options{}; - ASSERT_STATUS_OK(config_options.AddConfigEntry(webgpu::options::kDawnBackendType, - webgpu::options::kDawnBackendType_Vulkan)); - auto webgpu_ep = WebGpuExecutionProviderWithOptions(config_options); + auto webgpu_ep = DefaultWebGpuExecutionProvider(); if (!webgpu_ep) { GTEST_SKIP() << "WebGPU execution provider is not available."; } @@ -387,8 +376,6 @@ TEST(WebGpuMatMulAlgorithmTest, ReselectsAlgorithmForEachDynamicShape) { TensorShape({0, 8}), std::vector{unused_storage}, TensorShape({1, 8}), 0.0f); } -#endif // defined(DAWN_ENABLE_VULKAN) -#endif // 2D aligned baseline shapes. TEST(MatMul_Large, DISABLED_Aligned) { From 756e3913c5df4b663f8be61b266ce164fb8c1533 Mon Sep 17 00:00:00 2001 From: Jiajia Qin Date: Tue, 29 Sep 2026 14:43:09 +0800 Subject: [PATCH 24/25] Address WebGPU MatMul review feedback --- .../webgpu_matmul_algorithm_scheduler.md | 22 +- .../core/providers/webgpu/math/matmul.cc | 16 +- .../providers/webgpu/math/matmul_algorithm.cc | 45 +++++ .../providers/webgpu/math/matmul_algorithm.h | 37 +--- .../webgpu/math/matmul_algorithm_scheduler.cc | 189 ++++++++++++++++++ .../webgpu/math/matmul_algorithm_scheduler.h | 187 +++-------------- .../webgpu/math/matmul_compute_dispatcher.h | 2 + .../vendor/intel/math/gemm_subgroup_utils.h | 5 + .../webgpu/vendor/intel/math/matmul.cc | 55 ++++- .../webgpu/vendor/intel/math/matmul.h | 18 +- .../intel/math/matmul_algorithm_scheduler.cc | 29 +++ .../intel/math/matmul_algorithm_scheduler.h | 16 +- .../webgpu/webgpu_provider_factory.cc | 2 +- .../webgpu/matmul_algorithm_scheduler_test.cc | 46 +++-- .../providers/webgpu/matmul_large_test.cc | 33 +-- 15 files changed, 443 insertions(+), 259 deletions(-) create mode 100644 onnxruntime/core/providers/webgpu/math/matmul_algorithm.cc create mode 100644 onnxruntime/core/providers/webgpu/math/matmul_algorithm_scheduler.cc create mode 100644 onnxruntime/core/providers/webgpu/vendor/intel/math/matmul_algorithm_scheduler.cc diff --git a/docs/design/webgpu_matmul_algorithm_scheduler.md b/docs/design/webgpu_matmul_algorithm_scheduler.md index ab6010c90f219..15269933e4a86 100644 --- a/docs/design/webgpu_matmul_algorithm_scheduler.md +++ b/docs/design/webgpu_matmul_algorithm_scheduler.md @@ -2,7 +2,7 @@ ## Goal -Make every `ComputeMatMul` implementation path explicit and independently testable. Selection policy must be separated from execution, retain the existing common rules as a fallback, and allow vendor-specific policy to override any performance heuristic. +Make every `MatMulComputeDispatcher` implementation path explicit and independently testable. Selection policy must be separated from execution, retain the existing common rules as a fallback, and allow vendor-specific policy to override any performance heuristic. ## Algorithms @@ -10,7 +10,7 @@ Introduce `MatMulAlgorithm` with five concrete values: - `SubgroupMatrix`: the common subgroup-matrix implementation in `subgroup_matrix_matmul.cc`. - `Naive`: `MatMulNaiveProgram`. -- `IntelSubgroup`: the Intel subgroup implementation in `vendor/intel/math/matmul.cc`. +- `Subgroup`: the common subgroup implementation currently located in `vendor/intel/math/matmul.cc`. - `Packed`: the generic packed `MatMulProgram` without Split-K. - `PackedSplitK`: the same packed program with Split-K initialization and atomic accumulation. @@ -24,17 +24,21 @@ Add a `MatMulAlgorithmScheduler` base class. Automatic selection uses this order 2. Ask the virtual `SelectVendorAlgorithm` policy hook for any algorithm. A vendor may use its own thresholds for every algorithm, not only vendor-specific implementations. 3. If the vendor returns no selection, call the private, non-virtual `SelectCommonAlgorithm` fallback. It selects `SubgroupMatrix` when supported, `Naive` when `N < 8 && K < 8`, `PackedSplitK` when the existing `SplitKConfig::UseSplitK` rule succeeds, and otherwise `Packed`. -An Intel-derived scheduler implements the vendor hook. It preserves the original policy by selecting `SubgroupMatrix` first when applicable, then selecting `IntelSubgroup` under the current Intel subgroup rule. Other vendors use the base scheduler unchanged. Future vendor policies can derive from the base scheduler, override as many performance ranges as needed, and return no selection to delegate the remaining ranges to the common fallback without adding vendor conditionals to `ComputeMatMul`. +An Intel-derived scheduler implements the vendor hook. It preserves the original policy by selecting `SubgroupMatrix` first when applicable, then selecting the common `Subgroup` implementation under Intel-specific shape thresholds. Other vendors use the base scheduler unchanged. Future vendor policies can derive from the base scheduler, override as many performance ranges as needed, and return no selection to delegate the remaining ranges to the common fallback without adding vendor conditionals to `MatMulComputeDispatcher`. -The scheduler accepts immutable problem facts rather than mutable policy decisions: logical and packed dimensions, batch sizes, adapter architecture, input types, packing and layout facts, deterministic-compute state, activation and bias facts, and device capabilities. It also owns a copy of the immutable `SplitKConfig` selected for the adapter. The private common fallback evaluates `SplitKConfig::UseSplitK` directly from the problem facts instead of receiving a precomputed decision from `ComputeMatMul`. A vendor may ignore that common recommendation and apply independent thresholds from the raw facts. This keeps the scheduler deterministic and unit-testable without a WebGPU device. +The subgroup implementation is not inherently Intel-specific: forcing it requires the WebGPU `Subgroups` feature and a supported subgroup size. Its current source location under `vendor/intel` is historical and should move to the common math directory in a follow-up, while Intel-specific automatic-selection thresholds and tuning remain under `vendor/intel`. -`SplitKConfig` contains only generic Split-K eligibility evaluation and data. Adapter routing is handled by a small generic factory, while Intel architecture profiles and their measured threshold tables live under `vendor/intel`. `WebGpuContext` owns the selected configuration so GEMM and MatMul use the same profile; the MatMul scheduler receives that configuration when it is created. A future vendor can add its own profile builder and factory route without adding conditions to `ComputeMatMul` or changing the generic evaluator. +The subgroup shader implements subgroup-size branches for 8, 16, and 32 lanes. A fixed-size adapter may use any of those sizes directly. An adapter reporting a size range must expose subgroup-size control so the execution plan can select a supported size explicitly; otherwise the subgroup algorithm is unavailable. + +The scheduler accepts immutable problem facts rather than mutable policy decisions: logical and packed dimensions, batch sizes, adapter architecture, input types, packing and layout facts, deterministic-compute state, activation and bias facts, and device capabilities. It also owns a copy of the immutable `SplitKConfig` selected for the adapter. The private common fallback evaluates `SplitKConfig::UseSplitK` directly from the problem facts instead of receiving a precomputed decision from `MatMulComputeDispatcher`. A vendor may ignore that common recommendation and apply independent thresholds from the raw facts. This keeps the scheduler deterministic and unit-testable without a WebGPU device. + +`SplitKConfig` contains only generic Split-K eligibility evaluation and data. Adapter routing is handled by a small generic factory, while Intel architecture profiles and their measured threshold tables live under `vendor/intel`. `WebGpuContext` owns the selected configuration so GEMM and MatMul use the same profile; the MatMul scheduler receives that configuration when it is created. A future vendor can add its own profile builder and factory route without adding conditions to `MatMulComputeDispatcher` or changing the generic evaluator. ## Execution Configuration Algorithm selection and execution tuning are separate decisions. After selecting one enum, the scheduler creates a `MatMulExecutionPlan` containing that enum and a typed algorithm configuration. It first asks the protected `SelectVendorConfiguration` hook for tuning, then uses private common defaults when the vendor declines. The tuning hook runs for both automatic and forced algorithms, so the test-only forcing option controls the implementation path without disabling real device tuning. -The packed configuration initially contains workgroup size, elements per thread, inner tile size, and Split-K size. `ApplyMatMulPacked` consumes those values directly and includes shader-affecting values in its cache key. The common configuration preserves the existing `8x8x1` workgroup, `4x1x1` or `4x4x1` elements-per-thread rule, inner tile size 32, and adapter Split-K size. A vendor may replace any of these values without changing `ComputeMatMul` or the packed implementation. +The packed configuration initially contains workgroup size, elements per thread, inner tile size, and Split-K size. `ApplyMatMulPacked` consumes those values directly and includes shader-affecting values in its cache key. The common configuration preserves the existing `8x8x1` workgroup, `4x1x1` or `4x4x1` elements-per-thread rule, inner tile size 32, and adapter Split-K size. A vendor may replace any of these values without changing `MatMulComputeDispatcher` or the packed implementation. Packed tuning is validated at the execution boundary. The current shader requires both Z workgroup dimensions to remain one because Z identifies a batch or Split-K slice rather than a tiled output axis. Split-K sizes must be greater than one and aligned to the inner tile so adjacent workgroups cannot overlap their K ranges. Dispatch counts use widened, overflow-safe arithmetic and are range-checked before conversion to WebGPU's 32-bit dimensions. @@ -42,11 +46,11 @@ Configuration is represented by an algorithm-specific variant rather than a bag ## Forced Test Selection -Add the internal WebGPU session configuration key `ep.webgpuexecutionprovider.forceMatmulAlgorithm`. Accepted values are `subgroup_matrix`, `naive`, `intel_subgroup`, `packed`, and `packed_split_k`. The option is parsed when the WebGPU EP is created, stored as `std::optional`, and exposed read-only through `ComputeContextBase`. +Add the internal WebGPU session configuration key `ep.webgpuexecutionprovider.forceMatmulAlgorithm`. Accepted values are `subgroup_matrix`, `naive`, `subgroup`, `packed`, and `packed_split_k`. The option is parsed when the WebGPU EP is created, stored as `std::optional`, and exposed read-only through `ComputeContextBase`. When set, the scheduler returns the requested enum before applying heuristic rules. The dispatcher then validates the algorithm's hard prerequisites. Unsupported device features, data types, layouts, deterministic-compute settings, or other correctness constraints produce a descriptive failure naming the forced algorithm; forced mode never silently falls back. -Heuristic thresholds are not hard prerequisites. For example, forcing Intel subgroup bypasses its current `M/N/K` performance thresholds while still requiring an Intel adapter with subgroup support. Forcing Split-K bypasses performance thresholds while still requiring a usable Split-K configuration, non-deterministic compute, compatible packing/activation, and supported bias layout. +Heuristic thresholds are not hard prerequisites. For example, forcing subgroup bypasses Intel's current `M/N/K` performance thresholds while still requiring subgroup support. Forcing Split-K bypasses performance thresholds while still requiring a usable Split-K configuration, non-deterministic compute, compatible packing/activation, and supported bias layout. Invalid option strings fail during WebGPU provider creation and list accepted values. @@ -63,7 +67,7 @@ For each call, `MatMulComputeDispatcher::Compute` computes shared shape and capa The scheduler remains pure policy: it creates plans but does not create device programs, cache tensor data, or execute kernels. `SubgroupMatrixMatMulImpl` is named for the implementation it owns and is reached only when the plan selects `SubgroupMatrix`; its applicability query is non-mutating, and its execution method does not communicate selection through a `handled` output. This removes trial execution as a dispatch mechanism. -The other implementations remain focused stateless functions or program builders unless they acquire persistent state in the future. The naive, Intel subgroup, packed, and packed Split-K paths therefore do not receive empty polymorphic wrapper classes. The generic packed helper takes an explicit Split-K mode and packed configuration; it does not re-run selection or tuning heuristics. WebGPU's existing program cache continues to own reusable compiled programs. +The other implementations remain focused stateless functions or program builders unless they acquire persistent state in the future. The naive, subgroup, packed, and packed Split-K paths therefore do not receive empty polymorphic wrapper classes. The generic packed helper takes an explicit Split-K mode and packed configuration; it does not re-run selection or tuning heuristics. WebGPU's existing program cache continues to own reusable compiled programs. This replaces `MatMulOptImplCache` and the generic `MatMulOptImpl` interface. If another algorithm later needs persistent implementation state, the dispatcher can own a separately named backend for that algorithm without changing scheduler policy or pretending that one object represents whichever algorithm happened to be selected most recently. diff --git a/onnxruntime/core/providers/webgpu/math/matmul.cc b/onnxruntime/core/providers/webgpu/math/matmul.cc index d8d7602f3c999..803078b1cf125 100644 --- a/onnxruntime/core/providers/webgpu/math/matmul.cc +++ b/onnxruntime/core/providers/webgpu/math/matmul.cc @@ -373,7 +373,8 @@ Status MatMulComputeDispatcher::Compute(ComputeContext& context, const bool can_use_subgroup_matrix = subgroup_impl != nullptr && subgroup_impl->CanApply(context, inputs, is_channels_last, b_is_constant); - const bool has_intel_subgroup_capability = intel::HasMatMulIntelCapability(context); + const std::optional subgroup_size = intel::SelectMatMulSubgroupSize(context); + const bool has_subgroup_capability = subgroup_size.has_value(); const int64_t batch_a = logical_a_shape.NumDimensions() > 2 ? logical_a_shape.SizeToDimension(logical_a_shape.NumDimensions() - 2) @@ -402,7 +403,8 @@ Status MatMulComputeDispatcher::Compute(ComputeContext& context, selection_params.a_data_type = a->GetElementType(); selection_params.b_data_type = b->GetElementType(); selection_params.can_use_subgroup_matrix = can_use_subgroup_matrix; - selection_params.has_intel_subgroup_capability = has_intel_subgroup_capability; + selection_params.has_subgroup_capability = has_subgroup_capability; + selection_params.subgroup_size = subgroup_size.value_or(0); selection_params.is_vec4 = helper.K() % 4 == 0 && helper.N() % 4 == 0; selection_params.deterministic_compute = context.KernelContext().GetUseDeterministicCompute(); selection_params.has_fused_activation = activation.activation_kind_ != ActivationKind::None; @@ -416,10 +418,11 @@ Status MatMulComputeDispatcher::Compute(ComputeContext& context, "MatMul algorithm ", MatMulAlgorithmName(algorithm), " received an incompatible vendor configuration."); const auto* packed_configuration = std::get_if(&plan.configuration); + const auto* subgroup_configuration = std::get_if(&plan.configuration); MatMulAlgorithmPrerequisites prerequisites{}; prerequisites.can_use_subgroup_matrix = can_use_subgroup_matrix; - prerequisites.has_intel_subgroup_capability = has_intel_subgroup_capability; + prerequisites.has_subgroup_capability = has_subgroup_capability; prerequisites.has_nonzero_k = helper.K() > 0; prerequisites.split_k_configured = packed_configuration != nullptr && packed_configuration->split_dim_inner > 1; @@ -440,9 +443,10 @@ Status MatMulComputeDispatcher::Compute(ComputeContext& context, case MatMulAlgorithm::Naive: return ApplyMatMulNaive( context, activation, inputs, output_tensor, is_channels_last, helper); - case MatMulAlgorithm::IntelSubgroup: - return intel::ApplyMatMulIntel( - context, activation, inputs, output_tensor, is_channels_last); + case MatMulAlgorithm::Subgroup: + return intel::ApplyMatMulSubgroup( + context, activation, inputs, output_tensor, is_channels_last, + subgroup_configuration->subgroup_size); case MatMulAlgorithm::Packed: return ApplyMatMulPacked( context, activation, inputs, output_tensor, is_channels_last, helper, diff --git a/onnxruntime/core/providers/webgpu/math/matmul_algorithm.cc b/onnxruntime/core/providers/webgpu/math/matmul_algorithm.cc new file mode 100644 index 0000000000000..459f2ec57b771 --- /dev/null +++ b/onnxruntime/core/providers/webgpu/math/matmul_algorithm.cc @@ -0,0 +1,45 @@ +// Copyright (c) Microsoft Corporation. All rights reserved. +// Licensed under the MIT License. + +#include "core/providers/webgpu/math/matmul_algorithm.h" + +namespace onnxruntime { +namespace webgpu { + +std::optional ParseMatMulAlgorithm(std::string_view name) { + if (name == "subgroup_matrix") { + return MatMulAlgorithm::SubgroupMatrix; + } + if (name == "naive") { + return MatMulAlgorithm::Naive; + } + if (name == "subgroup") { + return MatMulAlgorithm::Subgroup; + } + if (name == "packed") { + return MatMulAlgorithm::Packed; + } + if (name == "packed_split_k") { + return MatMulAlgorithm::PackedSplitK; + } + return std::nullopt; +} + +std::string_view MatMulAlgorithmName(MatMulAlgorithm algorithm) { + switch (algorithm) { + case MatMulAlgorithm::SubgroupMatrix: + return "subgroup_matrix"; + case MatMulAlgorithm::Naive: + return "naive"; + case MatMulAlgorithm::Subgroup: + return "subgroup"; + case MatMulAlgorithm::Packed: + return "packed"; + case MatMulAlgorithm::PackedSplitK: + return "packed_split_k"; + } + return "unknown"; +} + +} // namespace webgpu +} // namespace onnxruntime \ No newline at end of file diff --git a/onnxruntime/core/providers/webgpu/math/matmul_algorithm.h b/onnxruntime/core/providers/webgpu/math/matmul_algorithm.h index 9740bfd451f8d..1118af425ab33 100644 --- a/onnxruntime/core/providers/webgpu/math/matmul_algorithm.h +++ b/onnxruntime/core/providers/webgpu/math/matmul_algorithm.h @@ -12,45 +12,14 @@ namespace webgpu { enum class MatMulAlgorithm { SubgroupMatrix, Naive, - IntelSubgroup, + Subgroup, Packed, PackedSplitK, }; -inline std::optional ParseMatMulAlgorithm(std::string_view name) { - if (name == "subgroup_matrix") { - return MatMulAlgorithm::SubgroupMatrix; - } - if (name == "naive") { - return MatMulAlgorithm::Naive; - } - if (name == "intel_subgroup") { - return MatMulAlgorithm::IntelSubgroup; - } - if (name == "packed") { - return MatMulAlgorithm::Packed; - } - if (name == "packed_split_k") { - return MatMulAlgorithm::PackedSplitK; - } - return std::nullopt; -} +std::optional ParseMatMulAlgorithm(std::string_view name); -inline std::string_view MatMulAlgorithmName(MatMulAlgorithm algorithm) { - switch (algorithm) { - case MatMulAlgorithm::SubgroupMatrix: - return "subgroup_matrix"; - case MatMulAlgorithm::Naive: - return "naive"; - case MatMulAlgorithm::IntelSubgroup: - return "intel_subgroup"; - case MatMulAlgorithm::Packed: - return "packed"; - case MatMulAlgorithm::PackedSplitK: - return "packed_split_k"; - } - return "unknown"; -} +std::string_view MatMulAlgorithmName(MatMulAlgorithm algorithm); } // namespace webgpu } // namespace onnxruntime diff --git a/onnxruntime/core/providers/webgpu/math/matmul_algorithm_scheduler.cc b/onnxruntime/core/providers/webgpu/math/matmul_algorithm_scheduler.cc new file mode 100644 index 0000000000000..6d02af7659f07 --- /dev/null +++ b/onnxruntime/core/providers/webgpu/math/matmul_algorithm_scheduler.cc @@ -0,0 +1,189 @@ +// Copyright (c) Microsoft Corporation. All rights reserved. +// Licensed under the MIT License. + +#include "core/providers/webgpu/math/matmul_algorithm_scheduler.h" + +#include +#include + +namespace onnxruntime { +namespace webgpu { + +bool IsMatMulPackedConfigurationValid( + const MatMulPackedConfiguration& configuration, + bool use_split_k) { + if (configuration.workgroup_size[0] == 0 || + configuration.workgroup_size[1] == 0 || + configuration.workgroup_size[2] != 1 || + configuration.elements_per_thread[0] == 0 || + configuration.elements_per_thread[1] == 0 || + configuration.elements_per_thread[2] != 1 || + configuration.tile_inner == 0) { + return false; + } + + return !use_split_k || + (configuration.split_dim_inner > 1 && + configuration.split_dim_inner % configuration.tile_inner == 0); +} + +std::optional TryGetMatMulPackedDispatchGroupCount( + uint64_t dimension, + uint32_t workgroup_size, + uint32_t elements_per_thread) { + if (workgroup_size == 0 || elements_per_thread == 0) { + return std::nullopt; + } + + const uint64_t elements_per_group = + static_cast(workgroup_size) * elements_per_thread; + const uint64_t group_count = + dimension == 0 ? 0 : 1 + (dimension - 1) / elements_per_group; + if (group_count > std::numeric_limits::max()) { + return std::nullopt; + } + + return static_cast(group_count); +} + +bool IsMatMulAlgorithmConfigurationCompatible(const MatMulExecutionPlan& plan) { + switch (plan.algorithm) { + case MatMulAlgorithm::SubgroupMatrix: + return std::holds_alternative(plan.configuration); + case MatMulAlgorithm::Naive: + return std::holds_alternative(plan.configuration); + case MatMulAlgorithm::Subgroup: + return std::holds_alternative(plan.configuration); + case MatMulAlgorithm::Packed: + case MatMulAlgorithm::PackedSplitK: + return std::holds_alternative(plan.configuration); + } + return false; +} + +bool MeetsMatMulAlgorithmPrerequisites( + MatMulAlgorithm algorithm, + const MatMulAlgorithmPrerequisites& prerequisites) { + switch (algorithm) { + case MatMulAlgorithm::SubgroupMatrix: + return prerequisites.can_use_subgroup_matrix; + case MatMulAlgorithm::Subgroup: + return prerequisites.has_subgroup_capability && + prerequisites.has_nonzero_k; + case MatMulAlgorithm::PackedSplitK: + return prerequisites.has_nonzero_k && + prerequisites.split_k_configured && + !prerequisites.deterministic_compute && + prerequisites.is_vec4 && + !prerequisites.has_fused_activation && + prerequisites.split_k_bias_layout_supported; + case MatMulAlgorithm::Packed: + return prerequisites.has_nonzero_k; + case MatMulAlgorithm::Naive: + return true; + } + return false; +} + +MatMulAlgorithmScheduler::MatMulAlgorithmScheduler(SplitKConfig split_k_config) + : split_k_config_{std::move(split_k_config)} {} + +MatMulAlgorithmScheduler::~MatMulAlgorithmScheduler() = default; + +MatMulAlgorithm MatMulAlgorithmScheduler::Select( + const MatMulAlgorithmSelectionParams& params, + std::optional forced_algorithm) const { + if (forced_algorithm.has_value()) { + return *forced_algorithm; + } + if (params.k == 0) { + return MatMulAlgorithm::Naive; + } + if (const auto vendor_algorithm = SelectVendorAlgorithm(params); vendor_algorithm.has_value()) { + return *vendor_algorithm; + } + return SelectCommonAlgorithm(params); +} + +MatMulExecutionPlan MatMulAlgorithmScheduler::CreateExecutionPlan( + const MatMulAlgorithmSelectionParams& params, + std::optional forced_algorithm) const { + const MatMulAlgorithm algorithm = Select(params, forced_algorithm); + std::optional configuration = + SelectVendorConfiguration(algorithm, params); + if (!configuration.has_value()) { + configuration = SelectCommonConfiguration(algorithm, params); + } + return MatMulExecutionPlan{algorithm, std::move(*configuration)}; +} + +std::optional MatMulAlgorithmScheduler::SelectVendorAlgorithm( + const MatMulAlgorithmSelectionParams& /*params*/) const { + return std::nullopt; +} + +std::optional MatMulAlgorithmScheduler::SelectVendorConfiguration( + MatMulAlgorithm /*algorithm*/, + const MatMulAlgorithmSelectionParams& /*params*/) const { + return std::nullopt; +} + +MatMulAlgorithm MatMulAlgorithmScheduler::SelectCommonAlgorithm( + const MatMulAlgorithmSelectionParams& params) const { + if (params.can_use_subgroup_matrix) { + return MatMulAlgorithm::SubgroupMatrix; + } + if (params.n < 8 && params.k < 8) { + return MatMulAlgorithm::Naive; + } + if (ShouldUseSplitK(params)) { + return MatMulAlgorithm::PackedSplitK; + } + return MatMulAlgorithm::Packed; +} + +MatMulAlgorithmConfiguration MatMulAlgorithmScheduler::SelectCommonConfiguration( + MatMulAlgorithm algorithm, + const MatMulAlgorithmSelectionParams& params) const { + switch (algorithm) { + case MatMulAlgorithm::SubgroupMatrix: + return MatMulSubgroupMatrixConfiguration{}; + case MatMulAlgorithm::Naive: + return MatMulNaiveConfiguration{}; + case MatMulAlgorithm::Subgroup: + return MatMulSubgroupConfiguration{params.subgroup_size}; + case MatMulAlgorithm::Packed: + case MatMulAlgorithm::PackedSplitK: { + MatMulPackedConfiguration configuration{}; + configuration.elements_per_thread = + params.packed_m <= 8 ? std::array{4, 1, 1} + : std::array{4, 4, 1}; + configuration.split_dim_inner = + algorithm == MatMulAlgorithm::PackedSplitK + ? split_k_config_.GetSplitDimInner() + : 1; + return configuration; + } + } + return MatMulNaiveConfiguration{}; +} + +bool MatMulAlgorithmScheduler::ShouldUseSplitK( + const MatMulAlgorithmSelectionParams& params) const { + if (params.deterministic_compute || params.has_fused_activation || + params.packed_m < 0 || params.n < 0 || params.k < 0) { + return false; + } + + return split_k_config_.UseSplitK( + params.is_vec4, + ActivationKind::None, + params.packed_batch_size, + static_cast(params.packed_m), + static_cast(params.n), + static_cast(params.k), + params.is_channels_last); +} + +} // namespace webgpu +} // namespace onnxruntime \ No newline at end of file diff --git a/onnxruntime/core/providers/webgpu/math/matmul_algorithm_scheduler.h b/onnxruntime/core/providers/webgpu/math/matmul_algorithm_scheduler.h index 303d097418386..337dc9d75690d 100644 --- a/onnxruntime/core/providers/webgpu/math/matmul_algorithm_scheduler.h +++ b/onnxruntime/core/providers/webgpu/math/matmul_algorithm_scheduler.h @@ -5,10 +5,8 @@ #include #include -#include #include #include -#include #include #include "core/common/common.h" @@ -18,6 +16,7 @@ namespace onnxruntime { namespace webgpu { +// Immutable problem facts used by automatic selection and execution tuning. struct MatMulAlgorithmSelectionParams { int64_t m = 0; int64_t n = 0; @@ -29,7 +28,8 @@ struct MatMulAlgorithmSelectionParams { int32_t a_data_type = 0; int32_t b_data_type = 0; bool can_use_subgroup_matrix = false; - bool has_intel_subgroup_capability = false; + bool has_subgroup_capability = false; + uint32_t subgroup_size = 0; bool is_vec4 = false; bool deterministic_compute = false; bool has_fused_activation = false; @@ -37,9 +37,12 @@ struct MatMulAlgorithmSelectionParams { bool is_channels_last = true; }; +// Algorithm-specific tuning selected together with the implementation. struct MatMulSubgroupMatrixConfiguration {}; struct MatMulNaiveConfiguration {}; -struct MatMulIntelSubgroupConfiguration {}; +struct MatMulSubgroupConfiguration { + uint32_t subgroup_size = 0; +}; struct MatMulPackedConfiguration { std::array workgroup_size{8, 8, 1}; @@ -48,73 +51,33 @@ struct MatMulPackedConfiguration { uint32_t split_dim_inner = 1; }; -inline bool IsMatMulPackedConfigurationValid( +bool IsMatMulPackedConfigurationValid( const MatMulPackedConfiguration& configuration, - bool use_split_k) { - if (configuration.workgroup_size[0] == 0 || - configuration.workgroup_size[1] == 0 || - configuration.workgroup_size[2] != 1 || - configuration.elements_per_thread[0] == 0 || - configuration.elements_per_thread[1] == 0 || - configuration.elements_per_thread[2] != 1 || - configuration.tile_inner == 0) { - return false; - } - - return !use_split_k || - (configuration.split_dim_inner > 1 && - configuration.split_dim_inner % configuration.tile_inner == 0); -} + bool use_split_k); -inline std::optional TryGetMatMulPackedDispatchGroupCount( +std::optional TryGetMatMulPackedDispatchGroupCount( uint64_t dimension, uint32_t workgroup_size, - uint32_t elements_per_thread) { - if (workgroup_size == 0 || elements_per_thread == 0) { - return std::nullopt; - } - - const uint64_t elements_per_group = - static_cast(workgroup_size) * elements_per_thread; - const uint64_t group_count = - dimension == 0 ? 0 : 1 + (dimension - 1) / elements_per_group; - if (group_count > std::numeric_limits::max()) { - return std::nullopt; - } - - return static_cast(group_count); -} + uint32_t elements_per_thread); using MatMulAlgorithmConfiguration = std::variant; +// Complete per-invocation decision consumed by the compute dispatcher. struct MatMulExecutionPlan { MatMulAlgorithm algorithm; MatMulAlgorithmConfiguration configuration; }; -inline bool IsMatMulAlgorithmConfigurationCompatible( - const MatMulExecutionPlan& plan) { - switch (plan.algorithm) { - case MatMulAlgorithm::SubgroupMatrix: - return std::holds_alternative(plan.configuration); - case MatMulAlgorithm::Naive: - return std::holds_alternative(plan.configuration); - case MatMulAlgorithm::IntelSubgroup: - return std::holds_alternative(plan.configuration); - case MatMulAlgorithm::Packed: - case MatMulAlgorithm::PackedSplitK: - return std::holds_alternative(plan.configuration); - } - return false; -} +bool IsMatMulAlgorithmConfigurationCompatible(const MatMulExecutionPlan& plan); +// Runtime correctness constraints validated immediately before dispatch. struct MatMulAlgorithmPrerequisites { bool can_use_subgroup_matrix = false; - bool has_intel_subgroup_capability = false; + bool has_subgroup_capability = false; bool has_nonzero_k = false; bool split_k_configured = false; bool deterministic_compute = false; @@ -123,131 +86,43 @@ struct MatMulAlgorithmPrerequisites { bool split_k_bias_layout_supported = true; }; -inline bool MeetsMatMulAlgorithmPrerequisites( +bool MeetsMatMulAlgorithmPrerequisites( MatMulAlgorithm algorithm, - const MatMulAlgorithmPrerequisites& prerequisites) { - switch (algorithm) { - case MatMulAlgorithm::SubgroupMatrix: - return prerequisites.can_use_subgroup_matrix; - case MatMulAlgorithm::IntelSubgroup: - return prerequisites.has_intel_subgroup_capability && - prerequisites.has_nonzero_k; - case MatMulAlgorithm::PackedSplitK: - return prerequisites.has_nonzero_k && - prerequisites.split_k_configured && - !prerequisites.deterministic_compute && - prerequisites.is_vec4 && - !prerequisites.has_fused_activation && - prerequisites.split_k_bias_layout_supported; - case MatMulAlgorithm::Packed: - return prerequisites.has_nonzero_k; - case MatMulAlgorithm::Naive: - return true; - } - return false; -} + const MatMulAlgorithmPrerequisites& prerequisites); +// Pure selection and tuning policy. Runtime validation remains in the dispatcher. class MatMulAlgorithmScheduler { public: - explicit MatMulAlgorithmScheduler(SplitKConfig split_k_config = {}) - : split_k_config_{std::move(split_k_config)} {} - virtual ~MatMulAlgorithmScheduler() = default; + explicit MatMulAlgorithmScheduler(SplitKConfig split_k_config = {}); + virtual ~MatMulAlgorithmScheduler(); ORT_DISALLOW_COPY_ASSIGNMENT_AND_MOVE(MatMulAlgorithmScheduler); MatMulAlgorithm Select( const MatMulAlgorithmSelectionParams& params, - std::optional forced_algorithm = std::nullopt) const { - if (forced_algorithm.has_value()) { - return *forced_algorithm; - } - if (params.k == 0) { - return MatMulAlgorithm::Naive; - } - if (const auto vendor_algorithm = SelectVendorAlgorithm(params); vendor_algorithm.has_value()) { - return *vendor_algorithm; - } - return SelectCommonAlgorithm(params); - } + std::optional forced_algorithm = std::nullopt) const; + // A forced plan may intentionally violate runtime prerequisites; the dispatcher validates it + // and reports an error naming the requested algorithm instead of silently changing selection. MatMulExecutionPlan CreateExecutionPlan( const MatMulAlgorithmSelectionParams& params, - std::optional forced_algorithm = std::nullopt) const { - const MatMulAlgorithm algorithm = Select(params, forced_algorithm); - std::optional configuration = - SelectVendorConfiguration(algorithm, params); - if (!configuration.has_value()) { - configuration = SelectCommonConfiguration(algorithm, params); - } - return MatMulExecutionPlan{algorithm, std::move(*configuration)}; - } + std::optional forced_algorithm = std::nullopt) const; protected: virtual std::optional SelectVendorAlgorithm( - const MatMulAlgorithmSelectionParams& /*params*/) const { - return std::nullopt; - } + const MatMulAlgorithmSelectionParams& params) const; virtual std::optional SelectVendorConfiguration( - MatMulAlgorithm /*algorithm*/, - const MatMulAlgorithmSelectionParams& /*params*/) const { - return std::nullopt; - } + MatMulAlgorithm algorithm, + const MatMulAlgorithmSelectionParams& params) const; private: - MatMulAlgorithm SelectCommonAlgorithm(const MatMulAlgorithmSelectionParams& params) const { - if (params.can_use_subgroup_matrix) { - return MatMulAlgorithm::SubgroupMatrix; - } - if (params.n < 8 && params.k < 8) { - return MatMulAlgorithm::Naive; - } - if (ShouldUseSplitK(params)) { - return MatMulAlgorithm::PackedSplitK; - } - return MatMulAlgorithm::Packed; - } + MatMulAlgorithm SelectCommonAlgorithm(const MatMulAlgorithmSelectionParams& params) const; MatMulAlgorithmConfiguration SelectCommonConfiguration( MatMulAlgorithm algorithm, - const MatMulAlgorithmSelectionParams& params) const { - switch (algorithm) { - case MatMulAlgorithm::SubgroupMatrix: - return MatMulSubgroupMatrixConfiguration{}; - case MatMulAlgorithm::Naive: - return MatMulNaiveConfiguration{}; - case MatMulAlgorithm::IntelSubgroup: - return MatMulIntelSubgroupConfiguration{}; - case MatMulAlgorithm::Packed: - case MatMulAlgorithm::PackedSplitK: { - MatMulPackedConfiguration configuration{}; - configuration.elements_per_thread = - params.packed_m <= 8 ? std::array{4, 1, 1} - : std::array{4, 4, 1}; - configuration.split_dim_inner = - algorithm == MatMulAlgorithm::PackedSplitK - ? split_k_config_.GetSplitDimInner() - : 1; - return configuration; - } - } - return MatMulNaiveConfiguration{}; - } - - bool ShouldUseSplitK(const MatMulAlgorithmSelectionParams& params) const { - if (params.deterministic_compute || params.has_fused_activation || - params.packed_m < 0 || params.n < 0 || params.k < 0) { - return false; - } + const MatMulAlgorithmSelectionParams& params) const; - return split_k_config_.UseSplitK( - params.is_vec4, - ActivationKind::None, - params.packed_batch_size, - static_cast(params.packed_m), - static_cast(params.n), - static_cast(params.k), - params.is_channels_last); - } + bool ShouldUseSplitK(const MatMulAlgorithmSelectionParams& params) const; SplitKConfig split_k_config_; }; diff --git a/onnxruntime/core/providers/webgpu/math/matmul_compute_dispatcher.h b/onnxruntime/core/providers/webgpu/math/matmul_compute_dispatcher.h index 9146807819f01..6d574411f6523 100644 --- a/onnxruntime/core/providers/webgpu/math/matmul_compute_dispatcher.h +++ b/onnxruntime/core/providers/webgpu/math/matmul_compute_dispatcher.h @@ -21,6 +21,7 @@ struct Activation; class ComputeContext; class ComputeContextBase; +// Stateful subgroup-matrix implementation owned independently from selection policy. class SubgroupMatrixMatMulImpl { public: SubgroupMatrixMatMulImpl() = default; @@ -40,6 +41,7 @@ class SubgroupMatrixMatMulImpl { bool b_is_constant) = 0; }; +// Selects, validates, and dispatches one concrete implementation per invocation. class MatMulComputeDispatcher { public: MatMulComputeDispatcher() = default; diff --git a/onnxruntime/core/providers/webgpu/vendor/intel/math/gemm_subgroup_utils.h b/onnxruntime/core/providers/webgpu/vendor/intel/math/gemm_subgroup_utils.h index 69e5d2907d7c2..c1d6139f16b0f 100644 --- a/onnxruntime/core/providers/webgpu/vendor/intel/math/gemm_subgroup_utils.h +++ b/onnxruntime/core/providers/webgpu/vendor/intel/math/gemm_subgroup_utils.h @@ -4,6 +4,7 @@ #pragma once #include +#include #include namespace onnxruntime { @@ -19,6 +20,10 @@ bool CanUseAVec4CooperativeLoad(std::string_view architecture, uint32_t dim_inner, int64_t elements_per_thread_y); +std::optional SelectMatMulSubgroupSize(uint32_t adapter_min_subgroup_size, + uint32_t adapter_max_subgroup_size, + bool has_subgroup_size_control); + } // namespace intel } // namespace webgpu } // namespace onnxruntime diff --git a/onnxruntime/core/providers/webgpu/vendor/intel/math/matmul.cc b/onnxruntime/core/providers/webgpu/vendor/intel/math/matmul.cc index 318ba4b3f35e9..e3b19088cd193 100644 --- a/onnxruntime/core/providers/webgpu/vendor/intel/math/matmul.cc +++ b/onnxruntime/core/providers/webgpu/vendor/intel/math/matmul.cc @@ -12,6 +12,12 @@ namespace onnxruntime { namespace webgpu { namespace intel { +namespace { + +constexpr std::array kSupportedMatMulSubgroupSizes{32, 16, 8}; + +} // namespace + Status MatMulSubgroupProgram::GenerateShaderCode(ShaderHelper& shader) const { const auto& a = shader.AddInput("a", ShaderUsage::UseUniform | ShaderUsage::UseIndicesTypeAlias | ShaderUsage::UseValueTypeAlias | ShaderUsage::UseElementTypeAlias); @@ -36,16 +42,46 @@ Status MatMulSubgroupProgram::GenerateShaderCode(ShaderHelper& shader) const { return Status::OK(); } -bool HasMatMulIntelCapability(const ComputeContext& context) { - return context.AdapterInfo().vendor == std::string_view{"intel"} && - context.HasFeature(wgpu::FeatureName::Subgroups); +std::optional SelectMatMulSubgroupSize(uint32_t adapter_min_subgroup_size, + uint32_t adapter_max_subgroup_size, + bool has_subgroup_size_control) { + if (adapter_min_subgroup_size == adapter_max_subgroup_size) { + for (const uint32_t size : kSupportedMatMulSubgroupSizes) { + if (adapter_min_subgroup_size == size) { + return size; + } + } + return std::nullopt; + } + + if (has_subgroup_size_control) { + for (const uint32_t size : kSupportedMatMulSubgroupSizes) { + if (adapter_min_subgroup_size <= size && size <= adapter_max_subgroup_size) { + return size; + } + } + } + return std::nullopt; } -Status ApplyMatMulIntel(ComputeContext& context, - const Activation& activation, - const std::vector& inputs, - Tensor* output, - bool is_channels_last) { +std::optional SelectMatMulSubgroupSize(const ComputeContext& context) { + if (!context.HasFeature(wgpu::FeatureName::Subgroups)) { + return std::nullopt; + } + + const auto& adapter_info = context.AdapterInfo(); + return SelectMatMulSubgroupSize( + adapter_info.subgroupMinSize, + adapter_info.subgroupMaxSize, + context.HasFeature(wgpu::FeatureName::SubgroupSizeControl)); +} + +Status ApplyMatMulSubgroup(ComputeContext& context, + const Activation& activation, + const std::vector& inputs, + Tensor* output, + bool is_channels_last, + uint32_t subgroup_size) { const auto* a = inputs[0]; const auto* b = inputs[1]; bool has_bias = inputs.size() > 2; @@ -131,6 +167,9 @@ Status ApplyMatMulIntel(ComputeContext& context, MatMulSubgroupProgram program{activation, has_bias, is_vec4, a_vec4, b_is_fp16, is_channels_last, elements_per_thread}; + if (context.HasFeature(wgpu::FeatureName::SubgroupSizeControl)) { + program.SetSubgroupSize(subgroup_size); + } program .CacheHint(activation.CacheKey(), absl::StrJoin(elements_per_thread, "-"), a_vec4, b_is_fp16, is_channels_last) diff --git a/onnxruntime/core/providers/webgpu/vendor/intel/math/matmul.h b/onnxruntime/core/providers/webgpu/vendor/intel/math/matmul.h index 691cce96e8ace..8fc1816d6baba 100644 --- a/onnxruntime/core/providers/webgpu/vendor/intel/math/matmul.h +++ b/onnxruntime/core/providers/webgpu/vendor/intel/math/matmul.h @@ -3,6 +3,9 @@ #pragma once +#include +#include + #include "core/providers/webgpu/webgpu_kernel.h" #include "core/providers/webgpu/shader_helper.h" #include "core/providers/webgpu/program.h" @@ -12,6 +15,8 @@ namespace onnxruntime { namespace webgpu { namespace intel { +// TODO: Move this common subgroup implementation out of vendor/intel. Keep only +// Intel-specific automatic-selection thresholds and tuning under the vendor directory. class MatMulSubgroupProgram final : public Program { public: MatMulSubgroupProgram(const Activation& activation, @@ -46,13 +51,14 @@ class MatMulSubgroupProgram final : public Program { const InlinedVector elements_per_thread_; }; -bool HasMatMulIntelCapability(const ComputeContext& context); +std::optional SelectMatMulSubgroupSize(const ComputeContext& context); -Status ApplyMatMulIntel(ComputeContext& context, - const Activation& activation, - const std::vector& inputs, - Tensor* output, - bool is_channels_last); +Status ApplyMatMulSubgroup(ComputeContext& context, + const Activation& activation, + const std::vector& inputs, + Tensor* output, + bool is_channels_last, + uint32_t subgroup_size); } // namespace intel } // namespace webgpu diff --git a/onnxruntime/core/providers/webgpu/vendor/intel/math/matmul_algorithm_scheduler.cc b/onnxruntime/core/providers/webgpu/vendor/intel/math/matmul_algorithm_scheduler.cc new file mode 100644 index 0000000000000..0f859c6a93d01 --- /dev/null +++ b/onnxruntime/core/providers/webgpu/vendor/intel/math/matmul_algorithm_scheduler.cc @@ -0,0 +1,29 @@ +// Copyright (c) Microsoft Corporation. All rights reserved. +// Licensed under the MIT License. + +#include "core/providers/webgpu/vendor/intel/math/matmul_algorithm_scheduler.h" + +#include + +namespace onnxruntime { +namespace webgpu { +namespace intel { + +IntelMatMulAlgorithmScheduler::IntelMatMulAlgorithmScheduler(SplitKConfig split_k_config) + : MatMulAlgorithmScheduler{std::move(split_k_config)} {} + +std::optional IntelMatMulAlgorithmScheduler::SelectVendorAlgorithm( + const MatMulAlgorithmSelectionParams& params) const { + if (params.can_use_subgroup_matrix) { + return MatMulAlgorithm::SubgroupMatrix; + } + if (params.has_subgroup_capability && + params.m >= 64 && params.n >= 512 && params.k >= 32) { + return MatMulAlgorithm::Subgroup; + } + return std::nullopt; +} + +} // namespace intel +} // namespace webgpu +} // namespace onnxruntime \ No newline at end of file diff --git a/onnxruntime/core/providers/webgpu/vendor/intel/math/matmul_algorithm_scheduler.h b/onnxruntime/core/providers/webgpu/vendor/intel/math/matmul_algorithm_scheduler.h index 745e620c901a8..477e22a34ed91 100644 --- a/onnxruntime/core/providers/webgpu/vendor/intel/math/matmul_algorithm_scheduler.h +++ b/onnxruntime/core/providers/webgpu/vendor/intel/math/matmul_algorithm_scheduler.h @@ -3,8 +3,6 @@ #pragma once -#include - #include "core/providers/webgpu/math/matmul_algorithm_scheduler.h" namespace onnxruntime { @@ -14,21 +12,11 @@ namespace intel { class IntelMatMulAlgorithmScheduler final : public MatMulAlgorithmScheduler { public: IntelMatMulAlgorithmScheduler() = default; - explicit IntelMatMulAlgorithmScheduler(SplitKConfig split_k_config) - : MatMulAlgorithmScheduler{std::move(split_k_config)} {} + explicit IntelMatMulAlgorithmScheduler(SplitKConfig split_k_config); protected: std::optional SelectVendorAlgorithm( - const MatMulAlgorithmSelectionParams& params) const override { - if (params.can_use_subgroup_matrix) { - return MatMulAlgorithm::SubgroupMatrix; - } - if (params.has_intel_subgroup_capability && - params.m >= 64 && params.n >= 512 && params.k >= 32) { - return MatMulAlgorithm::IntelSubgroup; - } - return std::nullopt; - } + const MatMulAlgorithmSelectionParams& params) const override; }; } // namespace intel diff --git a/onnxruntime/core/providers/webgpu/webgpu_provider_factory.cc b/onnxruntime/core/providers/webgpu/webgpu_provider_factory.cc index f8d68dc1c4a2b..9f5d3a6556ca6 100644 --- a/onnxruntime/core/providers/webgpu/webgpu_provider_factory.cc +++ b/onnxruntime/core/providers/webgpu/webgpu_provider_factory.cc @@ -129,7 +129,7 @@ WebGpuExecutionProviderConfig ParseEpConfig(const ConfigOptions& config_options) webgpu_ep_config.forced_matmul_algorithm = ParseMatMulAlgorithm(forced_matmul_algorithm_str); ORT_ENFORCE(webgpu_ep_config.forced_matmul_algorithm.has_value(), "Invalid forced MatMul algorithm: ", forced_matmul_algorithm_str, - ". Must be one of: subgroup_matrix, naive, intel_subgroup, packed, packed_split_k."); + ". Must be one of: subgroup_matrix, naive, subgroup, packed, packed_split_k."); } // parse force CPU node names diff --git a/onnxruntime/test/providers/webgpu/matmul_algorithm_scheduler_test.cc b/onnxruntime/test/providers/webgpu/matmul_algorithm_scheduler_test.cc index f58de6dd579aa..7313e22c040b3 100644 --- a/onnxruntime/test/providers/webgpu/matmul_algorithm_scheduler_test.cc +++ b/onnxruntime/test/providers/webgpu/matmul_algorithm_scheduler_test.cc @@ -99,7 +99,7 @@ TEST(MatMulAlgorithmParsingTest, RoundTripsEveryAlgorithmName) { constexpr TestCase test_cases[] = { {"subgroup_matrix", MatMulAlgorithm::SubgroupMatrix}, {"naive", MatMulAlgorithm::Naive}, - {"intel_subgroup", MatMulAlgorithm::IntelSubgroup}, + {"subgroup", MatMulAlgorithm::Subgroup}, {"packed", MatMulAlgorithm::Packed}, {"packed_split_k", MatMulAlgorithm::PackedSplitK}, }; @@ -268,6 +268,28 @@ TEST(MatMulAlgorithmConfigurationTest, PackedConfigurationKeepsBatchAxesUntiled) EXPECT_FALSE(IsMatMulPackedConfigurationValid(configuration, /*use_split_k=*/false)); } +TEST(MatMulAlgorithmConfigurationTest, SubgroupConfigurationCarriesSelectedSize) { + MatMulAlgorithmScheduler scheduler; + MatMulAlgorithmSelectionParams params{}; + params.subgroup_size = 16; + + const auto plan = scheduler.CreateExecutionPlan(params, MatMulAlgorithm::Subgroup); + const auto* configuration = std::get_if(&plan.configuration); + ASSERT_NE(configuration, nullptr); + EXPECT_EQ(configuration->subgroup_size, 16u); +} + +TEST(MatMulAlgorithmConfigurationTest, SubgroupSizeSelectionRequiresSupportedSize) { + EXPECT_EQ(intel::SelectMatMulSubgroupSize(8, 8, false), 8u); + EXPECT_EQ(intel::SelectMatMulSubgroupSize(16, 16, false), 16u); + EXPECT_EQ(intel::SelectMatMulSubgroupSize(32, 32, false), 32u); + EXPECT_FALSE(intel::SelectMatMulSubgroupSize(64, 64, false).has_value()); + EXPECT_FALSE(intel::SelectMatMulSubgroupSize(8, 32, false).has_value()); + EXPECT_EQ(intel::SelectMatMulSubgroupSize(8, 32, true), 32u); + EXPECT_EQ(intel::SelectMatMulSubgroupSize(8, 16, true), 16u); + EXPECT_FALSE(intel::SelectMatMulSubgroupSize(4, 4, true).has_value()); +} + TEST(MatMulAlgorithmConfigurationTest, SplitKConfigurationRequiresTileAlignedSplits) { MatMulPackedConfiguration configuration{}; configuration.tile_inner = 32; @@ -355,9 +377,9 @@ TEST(MatMulAlgorithmSchedulerTest, IntelSchedulerAppliesCurrentVendorRule) { params.m = 64; params.n = 512; params.k = 32; - params.has_intel_subgroup_capability = true; + params.has_subgroup_capability = true; - EXPECT_EQ(scheduler.Select(params), MatMulAlgorithm::IntelSubgroup); + EXPECT_EQ(scheduler.Select(params), MatMulAlgorithm::Subgroup); params.n = 511; EXPECT_EQ(scheduler.Select(params), MatMulAlgorithm::Packed); @@ -370,7 +392,7 @@ TEST(MatMulAlgorithmSchedulerTest, IntelSchedulerPreservesSubgroupMatrixPreceden params.n = 512; params.k = 32; params.can_use_subgroup_matrix = true; - params.has_intel_subgroup_capability = true; + params.has_subgroup_capability = true; EXPECT_EQ(scheduler.Select(params), MatMulAlgorithm::SubgroupMatrix); } @@ -437,27 +459,27 @@ TEST(MatMulAlgorithmPrerequisiteTest, IntelAVec4RequiresCompatibleRowsPerThread) TEST(MatMulAlgorithmPrerequisiteTest, IntelCapabilityDoesNotIncludeAutomaticThresholds) { MatMulAlgorithmPrerequisites prerequisites{}; - prerequisites.has_intel_subgroup_capability = true; + prerequisites.has_subgroup_capability = true; prerequisites.has_nonzero_k = true; - EXPECT_TRUE(MeetsMatMulAlgorithmPrerequisites(MatMulAlgorithm::IntelSubgroup, prerequisites)); + EXPECT_TRUE(MeetsMatMulAlgorithmPrerequisites(MatMulAlgorithm::Subgroup, prerequisites)); MatMulAlgorithmSelectionParams below_heuristic_threshold{}; below_heuristic_threshold.m = 1; below_heuristic_threshold.n = 1; below_heuristic_threshold.k = 1; - below_heuristic_threshold.has_intel_subgroup_capability = true; + below_heuristic_threshold.has_subgroup_capability = true; intel::IntelMatMulAlgorithmScheduler scheduler; - EXPECT_NE(scheduler.Select(below_heuristic_threshold), MatMulAlgorithm::IntelSubgroup); + EXPECT_NE(scheduler.Select(below_heuristic_threshold), MatMulAlgorithm::Subgroup); } -TEST(MatMulAlgorithmPrerequisiteTest, IntelSubgroupRejectsZeroContractionDimension) { +TEST(MatMulAlgorithmPrerequisiteTest, SubgroupRejectsZeroContractionDimension) { MatMulAlgorithmPrerequisites prerequisites{}; - prerequisites.has_intel_subgroup_capability = true; + prerequisites.has_subgroup_capability = true; - EXPECT_FALSE(MeetsMatMulAlgorithmPrerequisites(MatMulAlgorithm::IntelSubgroup, prerequisites)); + EXPECT_FALSE(MeetsMatMulAlgorithmPrerequisites(MatMulAlgorithm::Subgroup, prerequisites)); prerequisites.has_nonzero_k = true; - EXPECT_TRUE(MeetsMatMulAlgorithmPrerequisites(MatMulAlgorithm::IntelSubgroup, prerequisites)); + EXPECT_TRUE(MeetsMatMulAlgorithmPrerequisites(MatMulAlgorithm::Subgroup, prerequisites)); } } // namespace test diff --git a/onnxruntime/test/providers/webgpu/matmul_large_test.cc b/onnxruntime/test/providers/webgpu/matmul_large_test.cc index 58cc30722f8b6..ba8548ecb3431 100644 --- a/onnxruntime/test/providers/webgpu/matmul_large_test.cc +++ b/onnxruntime/test/providers/webgpu/matmul_large_test.cc @@ -17,6 +17,7 @@ #include "core/providers/webgpu/math/matmul_algorithm.h" #if !defined(ORT_USE_EP_API_ADAPTERS) #include "core/providers/webgpu/math/subgroup_matrix_config.h" +#include "core/providers/webgpu/vendor/intel/math/gemm_subgroup_utils.h" #include "core/providers/webgpu/webgpu_context.h" #endif #include "core/providers/webgpu/webgpu_provider_options.h" @@ -95,7 +96,7 @@ static std::optional GetForcedAlgorithmUnsupportedReason( #if defined(ORT_USE_EP_API_ADAPTERS) ORT_UNUSED_PARAMETER(ep); switch (algorithm) { - case webgpu::MatMulAlgorithm::IntelSubgroup: + case webgpu::MatMulAlgorithm::Subgroup: case webgpu::MatMulAlgorithm::PackedSplitK: case webgpu::MatMulAlgorithm::SubgroupMatrix: return "hardware-specific forced MatMul tests require direct adapter capability inspection."; @@ -108,14 +109,20 @@ static std::optional GetForcedAlgorithmUnsupportedReason( auto& context = webgpu::WebGpuContextFactory::GetContext(ep.GetDeviceId()); switch (algorithm) { - case webgpu::MatMulAlgorithm::IntelSubgroup: - if (context.AdapterInfo().vendor != std::string_view{"intel"}) { - return "intel_subgroup requires an Intel adapter."; - } + case webgpu::MatMulAlgorithm::Subgroup: { if (!context.DeviceHasFeature(wgpu::FeatureName::Subgroups)) { - return "intel_subgroup requires the WebGPU Subgroups feature."; + return "subgroup requires the WebGPU Subgroups feature."; + } + const auto& adapter_info = context.AdapterInfo(); + if (!webgpu::intel::SelectMatMulSubgroupSize( + adapter_info.subgroupMinSize, + adapter_info.subgroupMaxSize, + context.DeviceHasFeature(wgpu::FeatureName::SubgroupSizeControl)) + .has_value()) { + return "subgroup requires a supported subgroup size (8, 16, or 32)."; } break; + } case webgpu::MatMulAlgorithm::PackedSplitK: if (context.GetSplitKConfig().GetSplitDimInner() == 0) { return "packed_split_k is not configured for the selected adapter."; @@ -304,14 +311,14 @@ TEST(WebGpuMatMulAlgorithmTest, ForcedPackedRejectsZeroContractionDimension) { "MatMul algorithm packed"); } -TEST(WebGpuMatMulAlgorithmTest, ForcedIntelSubgroup) { - RunTestTyped({8, 32}, {32, 64}, false, webgpu::MatMulAlgorithm::IntelSubgroup); +TEST(WebGpuMatMulAlgorithmTest, ForcedSubgroup) { + RunTestTyped({8, 32}, {32, 64}, false, webgpu::MatMulAlgorithm::Subgroup); } -TEST(WebGpuMatMulAlgorithmTest, ForcedIntelSubgroupRejectsZeroContractionDimension) { - RunTestTyped({1, 0}, {0, 1}, false, webgpu::MatMulAlgorithm::IntelSubgroup, +TEST(WebGpuMatMulAlgorithmTest, ForcedSubgroupRejectsZeroContractionDimension) { + RunTestTyped({1, 0}, {0, 1}, false, webgpu::MatMulAlgorithm::Subgroup, OpTester::ExpectResult::kExpectFailure, - "MatMul algorithm intel_subgroup"); + "MatMul algorithm subgroup"); } TEST(WebGpuMatMulAlgorithmTest, ForcedPackedSplitK) { @@ -405,7 +412,7 @@ TEST(MatMul_Large, DISABLED_Broadcast4D) { } // Batched B (true bmm): A [..., M, K] x B [..., K, N] with matching batch. On the -// Intel subgroup path each (A, B) slice is dispatched on z. Covers small and +// Subgroup path each (A, B) slice is dispatched on z. Covers small and // larger batch counts with tile-aligned per-slice shapes. TEST(MatMul_Large, DISABLED_BatchedB) { RunBothTypes({2, 128, 64}, {2, 64, 1024}); @@ -458,7 +465,7 @@ TEST(MatMul_Large, DISABLED_ConstantWeightOddN) { // Broadcasted batch dims that are NOT identical but share the same batch // *product* (A=[2,1,...], B=[1,2,...] -> [2,2,...]; A=[1,4,...], B=[4,1,...] -> // [4,4,...]). A product-only batch check would wrongly route these onto the -// Intel subgroup path, which pairs slice i of A with slice i of B and copies A's +// Subgroup path, which pairs slice i of A with slice i of B and copies A's // shape to the output - producing the wrong output shape and mismatched pairing. // N is even and every per-slice shape is tile-aligned, so only the // identical-batch-dims guard (not the odd-N or partial-tile fallbacks) keeps them From 30f0fede9a48d35a9eeae9add3e46a57a024916d Mon Sep 17 00:00:00 2001 From: Jiajia Qin Date: Tue, 29 Sep 2026 16:17:57 +0800 Subject: [PATCH 25/25] Fix WebGPU MatMul scheduler formatting --- .../core/providers/webgpu/math/matmul_algorithm_scheduler.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/onnxruntime/core/providers/webgpu/math/matmul_algorithm_scheduler.h b/onnxruntime/core/providers/webgpu/math/matmul_algorithm_scheduler.h index 337dc9d75690d..05d102c5beb6b 100644 --- a/onnxruntime/core/providers/webgpu/math/matmul_algorithm_scheduler.h +++ b/onnxruntime/core/providers/webgpu/math/matmul_algorithm_scheduler.h @@ -63,7 +63,7 @@ std::optional TryGetMatMulPackedDispatchGroupCount( using MatMulAlgorithmConfiguration = std::variant; // Complete per-invocation decision consumed by the compute dispatcher.