[M149][tint] Collapse nested subgroupMin/Max ops on AMD On some AMD GPUs on Windows, nested subgroupMin/Max operations can trigger a crash. This change collapses multiple nested subgroupMin/Max calls to a single call, the innermost subgroupMin/Max operation. This is safe because the result of a subgroupMin/Max call is necessarily subgroup uniform. This bug exists in both the Vulkan and D3D12 drivers on Windows, and on Mac. On Windows it is fixed in later versions, but because it's a harmless optimizing transform it's applied broadly to all versions on affected backend platforms. Bug: 508265321 Fixed: 519445781 Change-Id: I658d0ddece1887d9211ad9ea2eb07a96ada0e3ce Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/311521 Commit-Queue: Natalie Chouinard <chouinard@google.com> Reviewed-by: dan sinclair <dsinclair@chromium.org> (cherry picked from commit 99622c4f7ef017a2f527e1e1d4afddcb93d04b35) Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/313895 Reviewed-by: James Price <jrprice@google.com>
diff --git a/src/dawn/native/Toggles.cpp b/src/dawn/native/Toggles.cpp index 16ba507..8935ef6 100644 --- a/src/dawn/native/Toggles.cpp +++ b/src/dawn/native/Toggles.cpp
@@ -753,6 +753,11 @@ {"decompose_uniform_buffers", "Decompose uniform buffers into arrays of vec4<u32> on backends for Vulkan and OpenGL.", "https://crbug.com/448452698", ToggleStage::Adapter}}, + {Toggle::CollapseSubgroupMinMax, + {"collapse_subgroup_min_max", + "Collapse redundant subgroup min and max operations (e.g., subgroupMin(subgroupMin(x))) into " + "a single operation. This works around a driver crash on some AMD GPUs.", + "https://crbug.com/508265321", ToggleStage::Device}}, {Toggle::VulkanEnableF16OnNvidia, {"vulkan_enable_f16_on_nvidia", "Enables F16 on Nvidia GPUs with Vulkan", "https://crbug.com/42251215", ToggleStage::Adapter}},
diff --git a/src/dawn/native/Toggles.h b/src/dawn/native/Toggles.h index 723b9b2..36d5ce5 100644 --- a/src/dawn/native/Toggles.h +++ b/src/dawn/native/Toggles.h
@@ -179,6 +179,7 @@ EnableShaderPrint, BlobCacheHashValidation, DecomposeUniformBuffers, + CollapseSubgroupMinMax, VulkanEnableF16OnNvidia, EnableRenderDocProcessInjection, VulkanUseDynamicRendering,
diff --git a/src/dawn/native/d3d12/PhysicalDeviceD3D12.cpp b/src/dawn/native/d3d12/PhysicalDeviceD3D12.cpp index 44787d3..6ee25b2 100644 --- a/src/dawn/native/d3d12/PhysicalDeviceD3D12.cpp +++ b/src/dawn/native/d3d12/PhysicalDeviceD3D12.cpp
@@ -839,6 +839,14 @@ deviceToggles->Default(Toggle::D3D12ForceClearCopyableDepthStencilTextureOnCreation, false); } + // Collapse redundant subgroup min and max operations to workaround a driver crash on older AMD + // GPUs. Should only affect AMD Windows Driver versions < 31.0.22000.0, but because this is a + // harmless "optimizing" workaround go ahead enable for all versions. See: + // https://crbug.com/508265321. + if (gpu_info::IsAMD(vendorId)) { + deviceToggles->Default(Toggle::CollapseSubgroupMinMax, true); + } + // Currently this toggle is only needed on Intel Gen9 and Gen9.5 GPUs. // See http://crbug.com/dawn/1579 for more information. if (gpu_info::IsIntelGen9(vendorId, deviceId)) {
diff --git a/src/dawn/native/d3d12/ShaderModuleD3D12.cpp b/src/dawn/native/d3d12/ShaderModuleD3D12.cpp index f7167ed..b6bb3ba 100644 --- a/src/dawn/native/d3d12/ShaderModuleD3D12.cpp +++ b/src/dawn/native/d3d12/ShaderModuleD3D12.cpp
@@ -330,6 +330,8 @@ device->IsToggleEnabled(Toggle::D3D12PolyfillReflectVec2F32); req.hlsl.tintOptions.workarounds.polyfill_subgroup_broadcast_f16 = device->IsToggleEnabled(Toggle::EnableSubgroupsIntelGen9); + req.hlsl.tintOptions.workarounds.collapse_subgroup_min_max = + device->IsToggleEnabled(Toggle::CollapseSubgroupMinMax); req.hlsl.tintOptions.extensions.polyfill_dot_4x8_packed = device->IsToggleEnabled(Toggle::PolyFillPacked4x8DotProduct);
diff --git a/src/dawn/native/metal/PhysicalDeviceMTL.mm b/src/dawn/native/metal/PhysicalDeviceMTL.mm index f5d5d42..ef27b26 100644 --- a/src/dawn/native/metal/PhysicalDeviceMTL.mm +++ b/src/dawn/native/metal/PhysicalDeviceMTL.mm
@@ -482,6 +482,8 @@ deviceToggles->Default(Toggle::MetalPolyfillTanhF16, true); // chromium:407109056: Floating point clamp is slightly inaccurate for subnormal values. deviceToggles->Default(Toggle::MetalPolyfillClampFloat, true); + // crbug.com/508265321: Nested subgroupMin/Max operations cause a crash in the AMD driver. + deviceToggles->Default(Toggle::CollapseSubgroupMinMax, true); } // On some Intel GPUs vertex only render pipeline get wrong depth result if no fragment
diff --git a/src/dawn/native/metal/ShaderModuleMTL.mm b/src/dawn/native/metal/ShaderModuleMTL.mm index 5ab2e9c..b3bf472 100644 --- a/src/dawn/native/metal/ShaderModuleMTL.mm +++ b/src/dawn/native/metal/ShaderModuleMTL.mm
@@ -357,6 +357,8 @@ device->IsToggleEnabled(Toggle::MetalPolyfillTanhF16); req.tintOptions.workarounds.replace_workgroup_bool_with_u32 = device->IsToggleEnabled(Toggle::MetalReplaceWorkgroupBoolWithU32); + req.tintOptions.workarounds.collapse_subgroup_min_max = + device->IsToggleEnabled(Toggle::CollapseSubgroupMinMax); req.tintOptions.extensions.disable_demote_to_helper = device->IsToggleEnabled(Toggle::DisableDemoteToHelper);
diff --git a/src/dawn/native/vulkan/PhysicalDeviceVk.cpp b/src/dawn/native/vulkan/PhysicalDeviceVk.cpp index fc5642f..93011e1 100644 --- a/src/dawn/native/vulkan/PhysicalDeviceVk.cpp +++ b/src/dawn/native/vulkan/PhysicalDeviceVk.cpp
@@ -1086,6 +1086,14 @@ deviceToggles->Default(Toggle::IgnoreImportedAHardwareBufferVulkanImageSize, true); } + // Collapse redundant subgroup min and max operations to workaround a driver crash on some AMD + // GPUs. Should only affect AMD Windows Driver versions < 31.0.22000.0, but because this is a + // harmless "optimizing" workaround go ahead enable for all versions. See: + // https://crbug.com/508265321. + if (IsWindowsAMD()) { + deviceToggles->Default(Toggle::CollapseSubgroupMinMax, true); + } + if (IsSwiftshader()) { // Swiftshader doesn't handle propagating decorations for descriptors through // OpCompositeExtract which happens when a binding_array is indexed "by value" instead of @@ -1404,6 +1412,14 @@ return gpu_info::IsGoogleSwiftshader(GetVendorId(), GetDeviceId()); } +bool PhysicalDevice::IsWindowsAMD() const { +#if DAWN_PLATFORM_IS(WINDOWS) + return gpu_info::IsAMD(GetVendorId()); +#else + return false; +#endif +} + std::optional<uint32_t> PhysicalDevice::FindDefaultComputeSubgroupSize() const { if (!mDeviceInfo.HasExt(DeviceExt::SubgroupSizeControl)) { return std::nullopt;
diff --git a/src/dawn/native/vulkan/PhysicalDeviceVk.h b/src/dawn/native/vulkan/PhysicalDeviceVk.h index 582b949..9598dc3 100644 --- a/src/dawn/native/vulkan/PhysicalDeviceVk.h +++ b/src/dawn/native/vulkan/PhysicalDeviceVk.h
@@ -66,6 +66,7 @@ bool IsAndroidSamsung() const; bool IsAndroidImgTec() const; bool IsPixel10() const; + bool IsWindowsAMD() const; bool IsIntelMesa() const; bool IsAmdMesa() const; bool IsAndroidHuawei() const;
diff --git a/src/dawn/native/vulkan/ShaderModuleVk.cpp b/src/dawn/native/vulkan/ShaderModuleVk.cpp index 534f0d1..1fd0900 100644 --- a/src/dawn/native/vulkan/ShaderModuleVk.cpp +++ b/src/dawn/native/vulkan/ShaderModuleVk.cpp
@@ -319,6 +319,8 @@ GetDevice()->IsToggleEnabled(Toggle::VulkanDirectVariableAccessTransformHandle); req.tintOptions.workarounds.polyfill_subgroup_broadcast_f16 = GetDevice()->IsToggleEnabled(Toggle::EnableSubgroupsIntelGen9); + req.tintOptions.workarounds.collapse_subgroup_min_max = + GetDevice()->IsToggleEnabled(Toggle::CollapseSubgroupMinMax); req.tintOptions.workarounds.cooperative_matrix_stride_is_matrix_elements = GetDevice()->IsToggleEnabled(Toggle::VulkanCooperativeMatrixStrideIsMatrixElements);
diff --git a/src/dawn/tests/end2end/SubgroupsTests.cpp b/src/dawn/tests/end2end/SubgroupsTests.cpp index edf4155..8a2a39a 100644 --- a/src/dawn/tests/end2end/SubgroupsTests.cpp +++ b/src/dawn/tests/end2end/SubgroupsTests.cpp
@@ -284,6 +284,53 @@ } } +// Regression test for a crash in the AMD driver when using nested subgroupMin/Max operations. +// See crbug.com/508265321. +TEST_P(SubgroupsShaderTests, NestedSubgroupMinMax) { + DAWN_TEST_UNSUPPORTED_IF(!IsSubgroupsEnabledInWGSL()); + + wgpu::ShaderModule module = utils::CreateShaderModule(device, R"( + enable subgroups; + + @group(0) @binding(0) var<storage, read> in: i32; + @group(0) @binding(1) var<storage, read_write> out: i32; + + @compute @workgroup_size(1) + fn main() { + let t = subgroupMax(in); + let r = t; + out = subgroupMin(r); + } + )"); + + wgpu::ComputePipelineDescriptor descriptor; + descriptor.layout = nullptr; + descriptor.compute.module = module; + descriptor.compute.entryPoint = "main"; + + wgpu::ComputePipeline pipeline = device.CreateComputePipeline(&descriptor); + + wgpu::Buffer inputBuffer = + utils::CreateBufferFromData(device, wgpu::BufferUsage::Storage, {42}); + wgpu::Buffer outputBuffer = utils::CreateBufferFromData( + device, wgpu::BufferUsage::Storage | wgpu::BufferUsage::CopySrc, {0}); + + wgpu::BindGroup bindGroup = utils::MakeBindGroup(device, pipeline.GetBindGroupLayout(0), + {{0, inputBuffer}, {1, outputBuffer}}); + + wgpu::CommandEncoder encoder = device.CreateCommandEncoder(); + wgpu::ComputePassEncoder pass = encoder.BeginComputePass(); + pass.SetPipeline(pipeline); + pass.SetBindGroup(0, bindGroup); + pass.DispatchWorkgroups(1); + pass.End(); + + wgpu::CommandBuffer commands = encoder.Finish(); + queue.Submit(1, &commands); + + EXPECT_BUFFER_U32_EQ(42u, outputBuffer, 0); +} + // DawnTestBase::CreateDeviceImpl always enables allow_unsafe_apis toggle. DAWN_INSTANTIATE_TEST(SubgroupsShaderTests, D3D12Backend(),
diff --git a/src/tint/lang/core/ir/transform/BUILD.bazel b/src/tint/lang/core/ir/transform/BUILD.bazel index ecdb32f..d025c2b 100644 --- a/src/tint/lang/core/ir/transform/BUILD.bazel +++ b/src/tint/lang/core/ir/transform/BUILD.bazel
@@ -49,6 +49,7 @@ "builtin_polyfill.cc", "builtin_scalarize.cc", "change_immediate_to_uniform.cc", + "collapse_subgroup_min_max.cc", "combine_access_instructions.cc", "conversion_polyfill.cc", "dead_code_elimination.cc", @@ -87,6 +88,7 @@ "builtin_polyfill.h", "builtin_scalarize.h", "change_immediate_to_uniform.h", + "collapse_subgroup_min_max.h", "combine_access_instructions.h", "conversion_polyfill.h", "dead_code_elimination.h", @@ -154,6 +156,7 @@ "builtin_polyfill_test.cc", "builtin_scalarize_test.cc", "change_immediate_to_uniform_test.cc", + "collapse_subgroup_min_max_test.cc", "combine_access_instructions_test.cc", "conversion_polyfill_test.cc", "dead_code_elimination_test.cc",
diff --git a/src/tint/lang/core/ir/transform/BUILD.cmake b/src/tint/lang/core/ir/transform/BUILD.cmake index c3ffc4f..384ce25 100644 --- a/src/tint/lang/core/ir/transform/BUILD.cmake +++ b/src/tint/lang/core/ir/transform/BUILD.cmake
@@ -57,6 +57,8 @@ lang/core/ir/transform/builtin_scalarize.h lang/core/ir/transform/change_immediate_to_uniform.cc lang/core/ir/transform/change_immediate_to_uniform.h + lang/core/ir/transform/collapse_subgroup_min_max.cc + lang/core/ir/transform/collapse_subgroup_min_max.h lang/core/ir/transform/combine_access_instructions.cc lang/core/ir/transform/combine_access_instructions.h lang/core/ir/transform/conversion_polyfill.cc @@ -154,6 +156,7 @@ lang/core/ir/transform/builtin_polyfill_test.cc lang/core/ir/transform/builtin_scalarize_test.cc lang/core/ir/transform/change_immediate_to_uniform_test.cc + lang/core/ir/transform/collapse_subgroup_min_max_test.cc lang/core/ir/transform/combine_access_instructions_test.cc lang/core/ir/transform/conversion_polyfill_test.cc lang/core/ir/transform/dead_code_elimination_test.cc
diff --git a/src/tint/lang/core/ir/transform/BUILD.gn b/src/tint/lang/core/ir/transform/BUILD.gn index 2b35934..8d6fb41 100644 --- a/src/tint/lang/core/ir/transform/BUILD.gn +++ b/src/tint/lang/core/ir/transform/BUILD.gn
@@ -63,6 +63,8 @@ "builtin_scalarize.h", "change_immediate_to_uniform.cc", "change_immediate_to_uniform.h", + "collapse_subgroup_min_max.cc", + "collapse_subgroup_min_max.h", "combine_access_instructions.cc", "combine_access_instructions.h", "conversion_polyfill.cc", @@ -154,6 +156,7 @@ "builtin_polyfill_test.cc", "builtin_scalarize_test.cc", "change_immediate_to_uniform_test.cc", + "collapse_subgroup_min_max_test.cc", "combine_access_instructions_test.cc", "conversion_polyfill_test.cc", "dead_code_elimination_test.cc",
diff --git a/src/tint/lang/core/ir/transform/collapse_subgroup_min_max.cc b/src/tint/lang/core/ir/transform/collapse_subgroup_min_max.cc new file mode 100644 index 0000000..ac11d36 --- /dev/null +++ b/src/tint/lang/core/ir/transform/collapse_subgroup_min_max.cc
@@ -0,0 +1,86 @@ +// Copyright 2026 The Dawn & Tint Authors +// +// Redistribution and use in source and binary forms, with or without +// modification, are permitted provided that the following conditions are met: +// +// 1. Redistributions of source code must retain the above copyright notice, this +// list of conditions and the following disclaimer. +// +// 2. Redistributions in binary form must reproduce the above copyright notice, +// this list of conditions and the following disclaimer in the documentation +// and/or other materials provided with the distribution. +// +// 3. Neither the name of the copyright holder nor the names of its +// contributors may be used to endorse or promote products derived from +// this software without specific prior written permission. +// +// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" +// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE +// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE +// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE +// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL +// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR +// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER +// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, +// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE +// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. + +#include "src/tint/lang/core/ir/transform/collapse_subgroup_min_max.h" + +#include "src/tint/lang/core/ir/builder.h" +#include "src/tint/lang/core/ir/module.h" +#include "src/tint/lang/core/ir/validator.h" + +namespace tint::core::ir::transform { + +namespace { + +bool IsCollapsibleSubgroupOp(core::BuiltinFn func) { + return (func == core::BuiltinFn::kSubgroupMin) || (func == core::BuiltinFn::kSubgroupMax); +} + +// Helper to check if a value is the result of a collapsible subgroup operation, +// possibly through "let" instructions. +bool IsCollapsibleSubgroupValue(core::ir::Value* value) { + while (auto* res = value->As<core::ir::InstructionResult>()) { + auto* inst = res->Instruction(); + if (auto* let = inst->As<core::ir::Let>()) { + value = let->Value(); + continue; + } + if (auto* call = inst->As<core::ir::CoreBuiltinCall>()) { + return IsCollapsibleSubgroupOp(call->Func()); + } + break; + } + return false; +} + +void Process(Module& ir) { + // Add all nested subgroupMin/Max calls that can be collapsed to a worklist. + Vector<core::ir::CoreBuiltinCall*, 16> worklist; + for (auto* inst : ir.Instructions()) { + if (auto* call = inst->As<core::ir::CoreBuiltinCall>()) { + if (IsCollapsibleSubgroupOp(call->Func())) { + if (IsCollapsibleSubgroupValue(call->Args()[0])) { + worklist.Push(call); + } + } + } + } + + // Replace outer call with inner result. + for (auto* call : worklist) { + call->Result(0)->ReplaceAllUsesWith(call->Args()[0]); + call->Destroy(); + } +} + +} // namespace + +Result<SuccessType> CollapseSubgroupMinMax(Module& ir) { + Process(ir); + return Success; +} + +} // namespace tint::core::ir::transform
diff --git a/src/tint/lang/core/ir/transform/collapse_subgroup_min_max.h b/src/tint/lang/core/ir/transform/collapse_subgroup_min_max.h new file mode 100644 index 0000000..352697d --- /dev/null +++ b/src/tint/lang/core/ir/transform/collapse_subgroup_min_max.h
@@ -0,0 +1,49 @@ +// Copyright 2026 The Dawn & Tint Authors +// +// Redistribution and use in source and binary forms, with or without +// modification, are permitted provided that the following conditions are met: +// +// 1. Redistributions of source code must retain the above copyright notice, this +// list of conditions and the following disclaimer. +// +// 2. Redistributions in binary form must reproduce the above copyright notice, +// this list of conditions and the following disclaimer in the documentation +// and/or other materials provided with the distribution. +// +// 3. Neither the name of the copyright holder nor the names of its +// contributors may be used to endorse or promote products derived from +// this software without specific prior written permission. +// +// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" +// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE +// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE +// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE +// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL +// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR +// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER +// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, +// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE +// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. + +#ifndef SRC_TINT_LANG_CORE_IR_TRANSFORM_COLLAPSE_SUBGROUP_MIN_MAX_H_ +#define SRC_TINT_LANG_CORE_IR_TRANSFORM_COLLAPSE_SUBGROUP_MIN_MAX_H_ + +#include "src/tint/utils/result.h" + +// Forward declarations. +namespace tint::core::ir { +class Module; +} + +namespace tint::core::ir::transform { + +/// CollapseSubgroupMinMax is a transform that replaces calls to subgroupMin and +/// subgroupMax builtins with their argument if that argument is also a subgroupMin/Max call which +/// is then by definition subgroup uniform. +/// @param module the module to transform +/// @returns success or failure +Result<SuccessType> CollapseSubgroupMinMax(Module& module); + +} // namespace tint::core::ir::transform + +#endif // SRC_TINT_LANG_CORE_IR_TRANSFORM_COLLAPSE_SUBGROUP_MIN_MAX_H_
diff --git a/src/tint/lang/core/ir/transform/collapse_subgroup_min_max_test.cc b/src/tint/lang/core/ir/transform/collapse_subgroup_min_max_test.cc new file mode 100644 index 0000000..56dd24c --- /dev/null +++ b/src/tint/lang/core/ir/transform/collapse_subgroup_min_max_test.cc
@@ -0,0 +1,191 @@ +// Copyright 2026 The Dawn & Tint Authors +// +// Redistribution and use in source and binary forms, with or without +// modification, are permitted provided that the following conditions are met: +// +// 1. Redistributions of source code must retain the above copyright notice, this +// list of conditions and the following disclaimer. +// +// 2. Redistributions in binary form must reproduce the above copyright notice, +// this list of conditions and the following disclaimer in the documentation +// and/or other materials provided with the distribution. +// +// 3. Neither the name of the copyright holder nor the names of its +// contributors may be used to endorse or promote products derived from +// this software without specific prior written permission. +// +// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" +// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE +// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE +// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE +// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL +// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR +// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER +// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, +// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE +// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. + +#include "src/tint/lang/core/ir/transform/collapse_subgroup_min_max.h" + +#include <utility> + +#include "src/tint/lang/core/ir/transform/helper_test.h" + +namespace tint::core::ir::transform { +namespace { + +using namespace tint::core::fluent_types; // NOLINT +using namespace tint::core::number_suffixes; // NOLINT + +class IR_CollapseSubgroupMinMaxTest : public TransformTest { + public: + IR_CollapseSubgroupMinMaxTest() {} +}; + +TEST_F(IR_CollapseSubgroupMinMaxTest, SubgroupMin_SubgroupMin) { + auto* u = b.FunctionParam("u", ty.i32()); + auto* func = b.Function("foo", ty.i32()); + func->SetParams({u}); + + b.Append(func->Block(), [&] { + auto* first = b.Call<i32>(core::BuiltinFn::kSubgroupMin, u); + auto* second = b.Call<i32>(core::BuiltinFn::kSubgroupMin, first); + b.Return(func, second); + }); + + auto* src = R"( +%foo = func(%u:i32):i32 { + $B1: { + %3:i32 = subgroupMin %u + %4:i32 = subgroupMin %3 + ret %4 + } +} +)"; + EXPECT_EQ(src, str()); + + auto* expect = R"( +%foo = func(%u:i32):i32 { + $B1: { + %3:i32 = subgroupMin %u + ret %3 + } +} +)"; + + Run(CollapseSubgroupMinMax); + EXPECT_EQ(expect, str()); +} + +TEST_F(IR_CollapseSubgroupMinMaxTest, SubgroupMin_SubgroupMax) { + auto* u = b.FunctionParam("u", ty.i32()); + auto* func = b.Function("foo", ty.void_()); + func->SetParams({u}); + + b.Append(func->Block(), [&] { + auto* first = b.Call<i32>(core::BuiltinFn::kSubgroupMax, u); + b.Call<i32>(core::BuiltinFn::kSubgroupMin, first); + b.Return(func); + }); + + auto* src = R"( +%foo = func(%u:i32):void { + $B1: { + %3:i32 = subgroupMax %u + %4:i32 = subgroupMin %3 + ret + } +} +)"; + EXPECT_EQ(src, str()); + + auto* expect = R"( +%foo = func(%u:i32):void { + $B1: { + %3:i32 = subgroupMax %u + ret + } +} +)"; + + Run(CollapseSubgroupMinMax); + EXPECT_EQ(expect, str()); +} + +TEST_F(IR_CollapseSubgroupMinMaxTest, SubgroupMax_SubgroupMin_SubgroupMax) { + auto* u = b.FunctionParam("u", ty.i32()); + auto* func = b.Function("foo", ty.void_()); + func->SetParams({u}); + + b.Append(func->Block(), [&] { + auto* first = b.Call<i32>(core::BuiltinFn::kSubgroupMax, u); + auto* second = b.Call<i32>(core::BuiltinFn::kSubgroupMin, first); + b.Call<i32>(core::BuiltinFn::kSubgroupMax, second); + b.Return(func); + }); + + auto* src = R"( +%foo = func(%u:i32):void { + $B1: { + %3:i32 = subgroupMax %u + %4:i32 = subgroupMin %3 + %5:i32 = subgroupMax %4 + ret + } +} +)"; + EXPECT_EQ(src, str()); + + auto* expect = R"( +%foo = func(%u:i32):void { + $B1: { + %3:i32 = subgroupMax %u + ret + } +} +)"; + + Run(CollapseSubgroupMinMax); + EXPECT_EQ(expect, str()); +} + +TEST_F(IR_CollapseSubgroupMinMaxTest, SubgroupMin_Let) { + auto* u = b.FunctionParam("u", ty.i32()); + auto* func = b.Function("foo", ty.void_()); + func->SetParams({u}); + + b.Append(func->Block(), [&] { + auto* first = b.Call<i32>(core::BuiltinFn::kSubgroupMin, u); + auto* let = b.Let(first); + b.Call<i32>(core::BuiltinFn::kSubgroupMin, let); + b.Return(func); + }); + + auto* src = R"( +%foo = func(%u:i32):void { + $B1: { + %3:i32 = subgroupMin %u + %4:i32 = let %3 + %5:i32 = subgroupMin %4 + ret + } +} +)"; + EXPECT_EQ(src, str()); + + auto* expect = R"( +%foo = func(%u:i32):void { + $B1: { + %3:i32 = subgroupMin %u + %4:i32 = let %3 + ret + } +} +)"; + + Run(CollapseSubgroupMinMax); + EXPECT_EQ(expect, str()); +} + +} // namespace +} // namespace tint::core::ir::transform
diff --git a/src/tint/lang/hlsl/writer/common/options.h b/src/tint/lang/hlsl/writer/common/options.h index 1530eec..5bb92d9 100644 --- a/src/tint/lang/hlsl/writer/common/options.h +++ b/src/tint/lang/hlsl/writer/common/options.h
@@ -147,10 +147,14 @@ /// Set to `true` to generate polyfill for `subgroupBroadcast(f16)` bool polyfill_subgroup_broadcast_f16 = false; + /// Set to `true` to collapse redundant subgroup min and max operations + bool collapse_subgroup_min_max = false; + TINT_REFLECT(Workarounds, scalarize_max_min_clamp, polyfill_reflect_vec2_f32, - polyfill_subgroup_broadcast_f16); + polyfill_subgroup_broadcast_f16, + collapse_subgroup_min_max); bool operator==(const Workarounds&) const = default; };
diff --git a/src/tint/lang/hlsl/writer/raise/raise.cc b/src/tint/lang/hlsl/writer/raise/raise.cc index 9466aca..e62466f 100644 --- a/src/tint/lang/hlsl/writer/raise/raise.cc +++ b/src/tint/lang/hlsl/writer/raise/raise.cc
@@ -39,6 +39,7 @@ #include "src/tint/lang/core/ir/transform/builtin_polyfill.h" #include "src/tint/lang/core/ir/transform/builtin_scalarize.h" #include "src/tint/lang/core/ir/transform/change_immediate_to_uniform.h" +#include "src/tint/lang/core/ir/transform/collapse_subgroup_min_max.h" #include "src/tint/lang/core/ir/transform/conversion_polyfill.h" #include "src/tint/lang/core/ir/transform/decompose_access.h" #include "src/tint/lang/core/ir/transform/demote_to_helper.h" @@ -350,6 +351,10 @@ TINT_CHECK_RESULT(raise::PixelLocal(module, config)); } + if (options.workarounds.collapse_subgroup_min_max) { + TINT_CHECK_RESULT(core::ir::transform::CollapseSubgroupMinMax(module)); + } + TINT_CHECK_RESULT(raise::BinaryPolyfill(module)); // Avoid potential UB (aka signed overflow) by performing unsigned integer arithmetic.
diff --git a/src/tint/lang/hlsl/writer/writer_fuzz.cc b/src/tint/lang/hlsl/writer/writer_fuzz.cc index 8bf2c8b..6968f56 100644 --- a/src/tint/lang/hlsl/writer/writer_fuzz.cc +++ b/src/tint/lang/hlsl/writer/writer_fuzz.cc
@@ -63,6 +63,7 @@ std::optional<uint32_t> num_workgroups_start_offset; std::vector<BindingPoint> ignored_by_robustness_transform; SubstituteOverridesConfig substitute_overrides_config; + bool collapse_subgroup_min_max; /// Reflect the fields of this class so that it can be used by tint::ForeachField() TINT_REFLECT(FuzzedOptions, @@ -83,7 +84,8 @@ first_instance_offset, num_workgroups_start_offset, ignored_by_robustness_transform, - substitute_overrides_config); + substitute_overrides_config, + collapse_subgroup_min_max); }; Result<SuccessType> IRFuzzer(core::ir::Module& module, @@ -123,6 +125,7 @@ options.workarounds.polyfill_reflect_vec2_f32 = fuzzed_options.polyfill_reflect_vec2_f32; options.workarounds.polyfill_subgroup_broadcast_f16 = fuzzed_options.polyfill_subgroup_broadcast_f16; + options.workarounds.collapse_subgroup_min_max = fuzzed_options.collapse_subgroup_min_max; options.extensions.polyfill_dot_4x8_packed = fuzzed_options.polyfill_dot_4x8_packed; options.extensions.polyfill_pack_unpack_4x8 = fuzzed_options.polyfill_pack_unpack_4x8; options.compiler = fuzzed_options.compiler;
diff --git a/src/tint/lang/msl/writer/common/options.h b/src/tint/lang/msl/writer/common/options.h index bdc4895..dbb0c3b 100644 --- a/src/tint/lang/msl/writer/common/options.h +++ b/src/tint/lang/msl/writer/common/options.h
@@ -124,6 +124,9 @@ /// Set to `true` to replace bool types in workgroup storage with u32. bool replace_workgroup_bool_with_u32 = false; + /// Set to `true` to collapse nested subgroupMin and subgroupMax operations. + bool collapse_subgroup_min_max = false; + TINT_REFLECT(Workarounds, scalarize_max_min_clamp, disable_module_constant_f16, @@ -132,7 +135,8 @@ polyfill_unpack_2x16_snorm, polyfill_unpack_2x16_unorm, polyfill_tanh_f16, - replace_workgroup_bool_with_u32); + replace_workgroup_bool_with_u32, + collapse_subgroup_min_max); TINT_REFLECT_HASH_CODE(Workarounds); bool operator==(const Workarounds&) const = default;
diff --git a/src/tint/lang/msl/writer/raise/raise.cc b/src/tint/lang/msl/writer/raise/raise.cc index 2bb6e77..c839dc4 100644 --- a/src/tint/lang/msl/writer/raise/raise.cc +++ b/src/tint/lang/msl/writer/raise/raise.cc
@@ -39,6 +39,7 @@ #include "src/tint/lang/core/ir/transform/builtin_polyfill.h" #include "src/tint/lang/core/ir/transform/builtin_scalarize.h" #include "src/tint/lang/core/ir/transform/change_immediate_to_uniform.h" +#include "src/tint/lang/core/ir/transform/collapse_subgroup_min_max.h" #include "src/tint/lang/core/ir/transform/conversion_polyfill.h" #include "src/tint/lang/core/ir/transform/demote_to_helper.h" #include "src/tint/lang/core/ir/transform/multiplanar_external_texture.h" @@ -84,6 +85,10 @@ TINT_CHECK_RESULT(raise::ValidateSubgroupMatrix(module)); + if (options.workarounds.collapse_subgroup_min_max) { + TINT_CHECK_RESULT(core::ir::transform::CollapseSubgroupMinMax(module)); + } + RaiseResult raise_result; // VertexPulling must come before BindingRemapper and Robustness.
diff --git a/src/tint/lang/msl/writer/writer_fuzz.cc b/src/tint/lang/msl/writer/writer_fuzz.cc index b3535af..2044002 100644 --- a/src/tint/lang/msl/writer/writer_fuzz.cc +++ b/src/tint/lang/msl/writer/writer_fuzz.cc
@@ -62,6 +62,7 @@ SubstituteOverridesConfig substitute_overrides_config; bool polyfill_tanh_f16; bool replace_workgroup_bool_with_u32; + bool collapse_subgroup_min_max; /// Reflect the fields of this class so that it can be used by tint::ForeachField() TINT_REFLECT(FuzzedOptions, @@ -83,7 +84,8 @@ group_to_argument_buffer_info, substitute_overrides_config, polyfill_tanh_f16, - replace_workgroup_bool_with_u32); + replace_workgroup_bool_with_u32, + collapse_subgroup_min_max); TINT_REFLECT_HASH_CODE(FuzzedOptions); }; @@ -131,6 +133,7 @@ options.workarounds.polyfill_tanh_f16 = fuzzed_options.polyfill_tanh_f16; options.workarounds.replace_workgroup_bool_with_u32 = fuzzed_options.replace_workgroup_bool_with_u32; + options.workarounds.collapse_subgroup_min_max = fuzzed_options.collapse_subgroup_min_max; options.fixed_sample_mask = fuzzed_options.fixed_sample_mask; options.pixel_local_attachments = fuzzed_options.pixel_local_attachments; options.vertex_pulling_config = fuzzed_options.vertex_pulling_config;
diff --git a/src/tint/lang/spirv/writer/common/options.h b/src/tint/lang/spirv/writer/common/options.h index afde49f..8c4d22d 100644 --- a/src/tint/lang/spirv/writer/common/options.h +++ b/src/tint/lang/spirv/writer/common/options.h
@@ -124,6 +124,9 @@ /// instructions as matrix elements instead of a source/dest pointee elements. bool cooperative_matrix_stride_is_matrix_elements = false; + /// Set to `true` to collapse redundant subgroup min and max operations + bool collapse_subgroup_min_max = false; + TINT_REFLECT(Workarounds, polyfill_case_switch, scalarize_max_min_clamp, @@ -139,7 +142,8 @@ polyfill_length_scalar_float, polyfill_distance_scalar_float, polyfill_saturate_as_min_max_f16, - cooperative_matrix_stride_is_matrix_elements); + cooperative_matrix_stride_is_matrix_elements, + collapse_subgroup_min_max); }; /// Any options which are controlled by the presence/absence of a vulkan extension.
diff --git a/src/tint/lang/spirv/writer/raise/raise.cc b/src/tint/lang/spirv/writer/raise/raise.cc index b30de70..ca5cc33 100644 --- a/src/tint/lang/spirv/writer/raise/raise.cc +++ b/src/tint/lang/spirv/writer/raise/raise.cc
@@ -34,6 +34,7 @@ #include "src/tint/lang/core/ir/transform/block_decorated_structs.h" #include "src/tint/lang/core/ir/transform/builtin_polyfill.h" #include "src/tint/lang/core/ir/transform/builtin_scalarize.h" +#include "src/tint/lang/core/ir/transform/collapse_subgroup_min_max.h" #include "src/tint/lang/core/ir/transform/combine_access_instructions.h" #include "src/tint/lang/core/ir/transform/conversion_polyfill.h" #include "src/tint/lang/core/ir/transform/decompose_access.h" @@ -205,6 +206,10 @@ TINT_CHECK_RESULT(core::ir::transform::DemoteToHelper(module)); } + if (options.workarounds.collapse_subgroup_min_max) { + TINT_CHECK_RESULT(core::ir::transform::CollapseSubgroupMinMax(module)); + } + raise::PolyfillConfig config = { .use_vulkan_memory_model = options.extensions.use_vulkan_memory_model, .version = options.spirv_version,
diff --git a/src/tint/lang/spirv/writer/writer_fuzz.cc b/src/tint/lang/spirv/writer/writer_fuzz.cc index e8ef81a..b0e685b 100644 --- a/src/tint/lang/spirv/writer/writer_fuzz.cc +++ b/src/tint/lang/spirv/writer/writer_fuzz.cc
@@ -86,6 +86,7 @@ bool cooperative_matrix_stride_is_matrix_elements; bool polyfill_length_scalar_float; bool polyfill_distance_scalar_float; + bool collapse_subgroup_min_max; /// Reflect the fields of this class so that it can be used by tint::ForeachField() TINT_REFLECT(FuzzedOptions, @@ -121,7 +122,8 @@ multisampled_framebuffer_fetch, cooperative_matrix_stride_is_matrix_elements, polyfill_length_scalar_float, - polyfill_distance_scalar_float); + polyfill_distance_scalar_float, + collapse_subgroup_min_max); TINT_REFLECT_HASH_CODE(FuzzedOptions); }; @@ -340,6 +342,7 @@ fuzzed_options.polyfill_distance_scalar_float; options.workarounds.cooperative_matrix_stride_is_matrix_elements = fuzzed_options.cooperative_matrix_stride_is_matrix_elements; + options.workarounds.collapse_subgroup_min_max = fuzzed_options.collapse_subgroup_min_max; options.multisampled_framebuffer_fetch = fuzzed_options.multisampled_framebuffer_fetch; TINT_CHECK_RESULT_UNWRAP(output, Generate(module, options));