From 61ee1edd166afb895db95b644c3abde6ea830573 Mon Sep 17 00:00:00 2001 From: Julian Ng-Thow-Hing Date: Tue, 28 Jul 2026 16:58:58 -0700 Subject: [PATCH] Update [ghstack-poisoned] --- .../webgpu/runtime/WebGPUShaderRegistry.cpp | 4 +- backends/webgpu/runtime/ops/amax/Reduce.cpp | 2 +- backends/webgpu/runtime/ops/amin/Reduce.cpp | 2 +- backends/webgpu/runtime/ops/amin/amin.wgsl | 49 -------- .../runtime/ops/{amax => extrema}/amax_wgsl.h | 2 +- .../runtime/ops/{amin => extrema}/amin_wgsl.h | 2 +- .../{amax/amax.wgsl => extrema/extrema.wgsl} | 4 +- .../webgpu/runtime/ops/extrema/extrema.yaml | 13 +++ backends/webgpu/test/op_tests/cases.py | 55 +++++---- backends/webgpu/test/ops/test_reduce.py | 108 +++++++++++++++++- backends/webgpu/test/test_wgsl_codegen.py | 58 +++++++++- 11 files changed, 216 insertions(+), 83 deletions(-) delete mode 100644 backends/webgpu/runtime/ops/amin/amin.wgsl rename backends/webgpu/runtime/ops/{amax => extrema}/amax_wgsl.h (97%) rename backends/webgpu/runtime/ops/{amin => extrema}/amin_wgsl.h (97%) rename backends/webgpu/runtime/ops/{amax/amax.wgsl => extrema/extrema.wgsl} (94%) create mode 100644 backends/webgpu/runtime/ops/extrema/extrema.yaml diff --git a/backends/webgpu/runtime/WebGPUShaderRegistry.cpp b/backends/webgpu/runtime/WebGPUShaderRegistry.cpp index 1cf4b952a5f..b475a6cd999 100644 --- a/backends/webgpu/runtime/WebGPUShaderRegistry.cpp +++ b/backends/webgpu/runtime/WebGPUShaderRegistry.cpp @@ -13,8 +13,6 @@ #include #include #include -#include -#include #include #include #include @@ -48,6 +46,8 @@ #include #include #include +#include +#include #include #include #include diff --git a/backends/webgpu/runtime/ops/amax/Reduce.cpp b/backends/webgpu/runtime/ops/amax/Reduce.cpp index 14675d97e96..c1dff62e035 100644 --- a/backends/webgpu/runtime/ops/amax/Reduce.cpp +++ b/backends/webgpu/runtime/ops/amax/Reduce.cpp @@ -9,7 +9,7 @@ #include #include #include -#include +#include #include diff --git a/backends/webgpu/runtime/ops/amin/Reduce.cpp b/backends/webgpu/runtime/ops/amin/Reduce.cpp index fbe574fdf0b..24a5ebeb826 100644 --- a/backends/webgpu/runtime/ops/amin/Reduce.cpp +++ b/backends/webgpu/runtime/ops/amin/Reduce.cpp @@ -9,7 +9,7 @@ #include #include #include -#include +#include #include diff --git a/backends/webgpu/runtime/ops/amin/amin.wgsl b/backends/webgpu/runtime/ops/amin/amin.wgsl deleted file mode 100644 index 4778800ab3d..00000000000 --- a/backends/webgpu/runtime/ops/amin/amin.wgsl +++ /dev/null @@ -1,49 +0,0 @@ -@group(0) @binding(0) var input: array; -@group(0) @binding(1) var output: array; - -struct Params { - num_rows: u32, - reduce_size: u32, -} -@group(0) @binding(2) var params: Params; - -override wg_size: u32 = 256u; - -// Cooperative shared-memory reduction; mirrors Vulkan reduce.glsl (a group of -// threads co-operates per reduction row, partials aggregated in shared memory). -// Fixed upper bound (>= any clamped wg_size); only [0, wg_size) is used. -var partials: array; - -@compute @workgroup_size(wg_size) -fn main( - @builtin(workgroup_id) wid: vec3, - @builtin(local_invocation_id) lid: vec3, - @builtin(num_workgroups) num_workgroups: vec3) { - // One workgroup per reduction row; 2D-fold lifts the 65535 grid cap. - let row = wid.x + wid.y * num_workgroups.x; - if (row >= params.num_rows) { - return; - } - let base = row * params.reduce_size; - - // Each thread reduces a strided slice of the row into a partial. Seed with - // the row's first element (always valid; reduce_size >= 1) so threads that - // own no element contribute a real value, not an out-of-range identity. - var acc = input[base]; - var i = lid.x; - while (i < params.reduce_size) { - acc = min(acc, input[base + i]); - i = i + wg_size; - } - partials[lid.x] = acc; - workgroupBarrier(); - - // Thread 0 aggregates the wg_size partials (mirrors Vulkan's group aggregate). - if (lid.x == 0u) { - var m = partials[0]; - for (var t = 1u; t < wg_size; t = t + 1u) { - m = min(m, partials[t]); - } - output[row] = m; - } -} diff --git a/backends/webgpu/runtime/ops/amax/amax_wgsl.h b/backends/webgpu/runtime/ops/extrema/amax_wgsl.h similarity index 97% rename from backends/webgpu/runtime/ops/amax/amax_wgsl.h rename to backends/webgpu/runtime/ops/extrema/amax_wgsl.h index 48ec8f20e27..0b7a4d2238e 100644 --- a/backends/webgpu/runtime/ops/amax/amax_wgsl.h +++ b/backends/webgpu/runtime/ops/extrema/amax_wgsl.h @@ -12,7 +12,7 @@ namespace executorch::backends::webgpu { -// @generated from amax.wgsl - DO NOT EDIT. +// @generated from extrema.wgsl - DO NOT EDIT. // wgsl-sha256: 35fc059d7c72caa17f9cb1128823ecfd8f75be4ce24b6cd4f9629a97b52f64c0 inline constexpr const char* kAmaxWGSL = R"( @group(0) @binding(0) var input: array; diff --git a/backends/webgpu/runtime/ops/amin/amin_wgsl.h b/backends/webgpu/runtime/ops/extrema/amin_wgsl.h similarity index 97% rename from backends/webgpu/runtime/ops/amin/amin_wgsl.h rename to backends/webgpu/runtime/ops/extrema/amin_wgsl.h index 40a97c67a63..8b8bf5456b1 100644 --- a/backends/webgpu/runtime/ops/amin/amin_wgsl.h +++ b/backends/webgpu/runtime/ops/extrema/amin_wgsl.h @@ -12,7 +12,7 @@ namespace executorch::backends::webgpu { -// @generated from amin.wgsl - DO NOT EDIT. +// @generated from extrema.wgsl - DO NOT EDIT. // wgsl-sha256: 8cb6035ae4d34eb2a6cc973d93d9847905722e967239c96033fccfe3a1943cb2 inline constexpr const char* kAminWGSL = R"( @group(0) @binding(0) var input: array; diff --git a/backends/webgpu/runtime/ops/amax/amax.wgsl b/backends/webgpu/runtime/ops/extrema/extrema.wgsl similarity index 94% rename from backends/webgpu/runtime/ops/amax/amax.wgsl rename to backends/webgpu/runtime/ops/extrema/extrema.wgsl index 2f23b9c38fa..6c55d60d9b7 100644 --- a/backends/webgpu/runtime/ops/amax/amax.wgsl +++ b/backends/webgpu/runtime/ops/extrema/extrema.wgsl @@ -32,7 +32,7 @@ fn main( var acc = input[base]; var i = lid.x; while (i < params.reduce_size) { - acc = max(acc, input[base + i]); + acc = ${REDUCE_FN}(acc, input[base + i]); i = i + wg_size; } partials[lid.x] = acc; @@ -42,7 +42,7 @@ fn main( if (lid.x == 0u) { var m = partials[0]; for (var t = 1u; t < wg_size; t = t + 1u) { - m = max(m, partials[t]); + m = ${REDUCE_FN}(m, partials[t]); } output[row] = m; } diff --git a/backends/webgpu/runtime/ops/extrema/extrema.yaml b/backends/webgpu/runtime/ops/extrema/extrema.yaml new file mode 100644 index 00000000000..a85da849c60 --- /dev/null +++ b/backends/webgpu/runtime/ops/extrema/extrema.yaml @@ -0,0 +1,13 @@ +# Copyright (c) Meta Platforms, Inc. and affiliates. +# All rights reserved. +# +# This source code is licensed under the BSD-style license found in the +# LICENSE file in the root directory of this source tree. + +extrema: + parameter_names_with_default_values: + REDUCE_FN: max + shader_variants: + - NAME: amax + - NAME: amin + REDUCE_FN: min diff --git a/backends/webgpu/test/op_tests/cases.py b/backends/webgpu/test/op_tests/cases.py index be46d35b520..aa8f000bceb 100644 --- a/backends/webgpu/test/op_tests/cases.py +++ b/backends/webgpu/test/op_tests/cases.py @@ -112,7 +112,15 @@ DequantizeConstModule, QuantizeModule, ) -from executorch.backends.webgpu.test.ops.test_reduce import AmaxModule, AminModule +from executorch.backends.webgpu.test.ops.test_reduce import ( + amax_sign_trap_input, + amax_tie_input, + AmaxModule, + amin_sign_trap_input, + amin_tie_input, + AminModule, + EXTREMA_CONFIGS, +) from executorch.backends.webgpu.test.ops.test_repeat import RepeatModule from executorch.backends.webgpu.test.ops.test_rms_norm import ( _CASES, @@ -489,35 +497,42 @@ def _floor_divide_suite() -> WebGPUTestSuite: ) -def _reduce_suite(module_cls) -> WebGPUTestSuite: - # Last-dim reduction; both keepdim variants over a 2d and a 3d shape. - return WebGPUTestSuite( - module_factory=lambda keepdim: module_cls(keepdim), - cases=[ - Case(name="keepdim_2d", construct={"keepdim": True}, inputs=((M1, M2),)), - Case(name="nodim_2d", construct={"keepdim": False}, inputs=((M1, M2),)), - Case( - name="keepdim_3d", - construct={"keepdim": True}, - inputs=((S, S1, S2),), - ), +def _reduce_suite(module_cls, op: str) -> WebGPUTestSuite: + generators = { + ("amax", "sign_trap"): amax_sign_trap_input, + ("amax", "tie"): amax_tie_input, + ("amin", "sign_trap"): amin_sign_trap_input, + ("amin", "tie"): amin_tie_input, + } + cases = [] + for name, shape, dim, keepdim, input_class in EXTREMA_CONFIGS: + inputs = (shape,) + kwargs = {} + if input_class != "default": + inputs = (InputSpec(shape=shape, gen=generators[(op, input_class)]),) + kwargs = {"atol": 0.0, "rtol": 0.0} + cases.append( Case( - name="nodim_3d", - construct={"keepdim": False}, - inputs=((S, S1, S2),), - ), - ], + name=name, + construct={"keepdim": keepdim, "dim": dim}, + inputs=inputs, + **kwargs, + ) + ) + return WebGPUTestSuite( + module_factory=module_cls, + cases=cases, ) @register_op_test("amax") def _amax_suite() -> WebGPUTestSuite: - return _reduce_suite(AmaxModule) + return _reduce_suite(AmaxModule, "amax") @register_op_test("amin") def _amin_suite() -> WebGPUTestSuite: - return _reduce_suite(AminModule) + return _reduce_suite(AminModule, "amin") @register_op_test("flip") diff --git a/backends/webgpu/test/ops/test_reduce.py b/backends/webgpu/test/ops/test_reduce.py index c7b91d8156c..cdfd7eed96c 100644 --- a/backends/webgpu/test/ops/test_reduce.py +++ b/backends/webgpu/test/ops/test_reduce.py @@ -19,12 +19,14 @@ from __future__ import annotations +import math import unittest import torch from executorch.backends.vulkan.partitioner.vulkan_partitioner import VulkanPartitioner from executorch.exir import to_edge_transform_and_lower +from executorch.exir.backend.utils import get_delegates, get_non_lowered_nodes class ReduceModule(torch.nn.Module): @@ -61,6 +63,70 @@ def _det_input(shape) -> torch.Tensor: return ((flat % 17) - 8).div(16.0).reshape(shape) +# Shared structural and manifest-driven extrema case authority. +EXTREMA_CONFIGS = ( + ("keepdim_2d", (37, 41), -1, True, "default"), + ("nodim_2d", (37, 41), -1, False, "default"), + ("keepdim_3d", (5, 7, 11), -1, True, "default"), + ("nodim_3d", (5, 7, 11), -1, False, "default"), + ("sign_trap_63_drop", (3, 63), -1, False, "sign_trap"), + ("tie_64_keep", (2, 64), -1, True, "tie"), + ("tie_65_posdim_drop", (2, 65), 1, False, "tie"), + ("sign_trap_255_keep", (2, 255), -1, True, "sign_trap"), + ("tie_256_drop", (2, 256), -1, False, "tie"), + ("tie_257_posdim_keep", (2, 257), 1, True, "tie"), +) + + +def _sign_trap_input(shape, *, for_max: bool) -> torch.Tensor: + magnitude = ((torch.arange(math.prod(shape), dtype=torch.float32) % 31) + 1).div( + 8.0 + ) + rows = magnitude.reshape(-1, shape[-1]) + for row_index, row in enumerate(rows): + row.add_(float(row_index)) + values = -magnitude if for_max else magnitude + return values.reshape(shape) + + +def amax_sign_trap_input(shape) -> torch.Tensor: + return _sign_trap_input(shape, for_max=True) + + +def amin_sign_trap_input(shape) -> torch.Tensor: + return _sign_trap_input(shape, for_max=False) + + +def _tie_input(shape, *, for_max: bool) -> torch.Tensor: + values = ((torch.arange(math.prod(shape), dtype=torch.float32) % 29) - 14).div(8.0) + rows = values.reshape(-1, shape[-1]) + for row_index, row in enumerate(rows): + extreme = 16.0 + float(row_index) + if not for_max: + extreme = -extreme + row[-2] = extreme + row[-1] = extreme + return values.reshape(shape) + + +def amax_tie_input(shape) -> torch.Tensor: + return _tie_input(shape, for_max=True) + + +def amin_tie_input(shape) -> torch.Tensor: + return _tie_input(shape, for_max=False) + + +def _extrema_input(op: str, input_class: str, shape) -> torch.Tensor: + if input_class == "default": + return _det_input(shape) + if input_class == "sign_trap": + return ( + amax_sign_trap_input(shape) if op == "amax" else amin_sign_trap_input(shape) + ) + return amax_tie_input(shape) if op == "amax" else amin_tie_input(shape) + + def _export(m: torch.nn.Module, x: torch.Tensor): ep = torch.export.export(m, (x,)) return to_edge_transform_and_lower( @@ -128,21 +194,55 @@ def export_reduce_model( class AmaxModule(torch.nn.Module): - def __init__(self, keepdim: bool) -> None: + def __init__(self, keepdim: bool, dim: int = -1) -> None: super().__init__() self.keepdim = keepdim + self.dim = dim def forward(self, x: torch.Tensor) -> torch.Tensor: - return torch.amax(x, dim=-1, keepdim=self.keepdim) + return torch.amax(x, dim=self.dim, keepdim=self.keepdim) class AminModule(torch.nn.Module): - def __init__(self, keepdim: bool) -> None: + def __init__(self, keepdim: bool, dim: int = -1) -> None: super().__init__() self.keepdim = keepdim + self.dim = dim def forward(self, x: torch.Tensor) -> torch.Tensor: - return torch.amin(x, dim=-1, keepdim=self.keepdim) + return torch.amin(x, dim=self.dim, keepdim=self.keepdim) + + +class TestExtrema(unittest.TestCase): + def test_config_contract(self) -> None: + self.assertEqual( + EXTREMA_CONFIGS, + ( + ("keepdim_2d", (37, 41), -1, True, "default"), + ("nodim_2d", (37, 41), -1, False, "default"), + ("keepdim_3d", (5, 7, 11), -1, True, "default"), + ("nodim_3d", (5, 7, 11), -1, False, "default"), + ("sign_trap_63_drop", (3, 63), -1, False, "sign_trap"), + ("tie_64_keep", (2, 64), -1, True, "tie"), + ("tie_65_posdim_drop", (2, 65), 1, False, "tie"), + ("sign_trap_255_keep", (2, 255), -1, True, "sign_trap"), + ("tie_256_drop", (2, 256), -1, False, "tie"), + ("tie_257_posdim_keep", (2, 257), 1, True, "tie"), + ), + ) + + def test_exports_fully_delegated(self) -> None: + for op, module_cls in (("amax", AmaxModule), ("amin", AminModule)): + for name, shape, dim, keepdim, input_class in EXTREMA_CONFIGS: + with self.subTest(op=op, config=name): + x = _extrema_input(op, input_class, shape) + ep = torch.export.export(module_cls(keepdim, dim).eval(), (x,)) + edge = to_edge_transform_and_lower( + ep, partitioner=[VulkanPartitioner()] + ) + graph = edge.exported_program().graph_module.graph + self.assertEqual(len(get_delegates(graph)), 1) + self.assertEqual(get_non_lowered_nodes(graph), []) if __name__ == "__main__": diff --git a/backends/webgpu/test/test_wgsl_codegen.py b/backends/webgpu/test/test_wgsl_codegen.py index f7bd2f5797e..7305142c4d1 100644 --- a/backends/webgpu/test/test_wgsl_codegen.py +++ b/backends/webgpu/test/test_wgsl_codegen.py @@ -230,7 +230,7 @@ def test_generated_output_manifest_digest(self) -> None: self.assertEqual(len(outputs), 134) self.assertEqual( digest.hexdigest(), - "ab15be30e7cfa2cb2f6fa7743d3b9f03535f5cc0b88d64d9b6bad8f777efda25", + "a3f6324b224c049f239beea3c38e1dd532e364f8d44c595eb5fcc84dba4a1a83", ) def test_rope_hf_reconstructs_full_2d_grid_stride(self) -> None: @@ -969,7 +969,61 @@ def test_to_copy_convert_template_roundtrip_byte_identical(self) -> None: ) self.assertEqual( hashlib.sha256(g.registry_path().read_bytes()).hexdigest(), - "ce1777820bffe77e7cdda312f86a3fd41a090f8f282a6add66666446d06b1608", + "74f972fce4077f12a52dfcf67a0d20ebeea47748283ced9e5c0bcffd659fef74", + ) + + def test_extrema_template_roundtrip_byte_identical(self) -> None: + extrema_dir = g.BACKEND_ROOT / "runtime/ops/extrema" + template_path = extrema_dir / "extrema.wgsl" + spec = g.parse_template_spec(template_path.with_suffix(".yaml")) + variants = {params["NAME"]: params for params in spec[template_path.stem]} + expected = { + "amax": ( + "max", + "35fc059d7c72caa17f9cb1128823ecfd8f75be4ce24b6cd4f9629a97b52f64c0", + ), + "amin": ( + "min", + "8cb6035ae4d34eb2a6cc973d93d9847905722e967239c96033fccfe3a1943cb2", + ), + } + self.assertEqual(set(variants), set(expected)) + template = template_path.read_text() + + for name, (reduce_fn, expected_hash) in expected.items(): + params = variants[name] + self.assertEqual(params["REDUCE_FN"], reduce_fn) + expanded = g.preprocess(template, {**g.WGSL_HELPERS, **params}) + self.assertEqual(g.wgsl_sha256(expanded), expected_hash) + + header_path = extrema_dir / f"{name}_wgsl.h" + header = header_path.read_text() + body = header.split('R"(', 1)[1].split(')";', 1)[0][1:] + self.assertEqual(body, expanded) + self.assertEqual(g.embedded_sha256(header), expected_hash) + self.assertEqual(g.parse_workgroup_size(body), (256, 1, 1)) + + entries = {entry.name: entry for entry in g.registry_entries()} + for name in expected: + self.assertEqual( + entries[name].include, + f"runtime/ops/extrema/{name}_wgsl.h", + ) + self.assertEqual(entries[name].symbol, g.symbol_base(name)) + + handler_hashes = { + "amax": "57f929b9f3087dc32403c3587884ce2ed4be2d03c4e80ff7035428b52e7e0e51", + "amin": "5dc947d4781a67df953b9c5970c4ab2119317b29657f983a9d308dfdf123dede", + } + for name, expected_hash in handler_hashes.items(): + handler = g.BACKEND_ROOT / f"runtime/ops/{name}/Reduce.cpp" + self.assertEqual( + hashlib.sha256(handler.read_bytes()).hexdigest(), expected_hash + ) + + self.assertEqual( + hashlib.sha256(g.registry_path().read_bytes()).hexdigest(), + "74f972fce4077f12a52dfcf67a0d20ebeea47748283ced9e5c0bcffd659fef74", ) def test_rms_norm_half_variant_is_type_correct(self) -> None: