From 99622c4f7ef017a2f527e1e1d4afddcb93d04b35 Mon Sep 17 00:00:00 2001
From: Natalie Chouinard <[email protected]>
Date: Mon, 01 Jun 2026 11:46:17 -0700
Subject: [PATCH] [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
Change-Id: I658d0ddece1887d9211ad9ea2eb07a96ada0e3ce
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/311521
Commit-Queue: Natalie Chouinard <[email protected]>
Reviewed-by: dan sinclair <[email protected]>
---

diff --git a/src/dawn/native/Toggles.cpp b/src/dawn/native/Toggles.cpp
index 3c6be49..18e535d 100644
--- a/src/dawn/native/Toggles.cpp
+++ b/src/dawn/native/Toggles.cpp
@@ -749,6 +749,11 @@
       "Decompose workgroup memory variables into flat scalar arrays and rewrite accesses as "
       "element-wise loads/stores with bitcasts on selected Intel GPUs on D3D12",
       "https://crbug.com/341991439", ToggleStage::Device}},
+    {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 6c339d3..38a8a90 100644
--- a/src/dawn/native/Toggles.h
+++ b/src/dawn/native/Toggles.h
@@ -179,6 +179,7 @@
     BlobCacheHashValidation,
     DecomposeUniformBuffers,
     D3D12DecomposeWorkgroupAccess,
+    CollapseSubgroupMinMax,
     VulkanEnableF16OnNvidia,
     EnableRenderDocProcessInjection,
     VulkanUseDynamicRendering,
diff --git a/src/dawn/native/d3d12/PhysicalDeviceD3D12.cpp b/src/dawn/native/d3d12/PhysicalDeviceD3D12.cpp
index aa3285c..2bcf42c 100644
--- a/src/dawn/native/d3d12/PhysicalDeviceD3D12.cpp
+++ b/src/dawn/native/d3d12/PhysicalDeviceD3D12.cpp
@@ -849,6 +849,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 a8be334..8fad306 100644
--- a/src/dawn/native/d3d12/ShaderModuleD3D12.cpp
+++ b/src/dawn/native/d3d12/ShaderModuleD3D12.cpp
@@ -341,6 +341,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 a3e3f23..4dc5dd4 100644
--- a/src/dawn/native/metal/PhysicalDeviceMTL.mm
+++ b/src/dawn/native/metal/PhysicalDeviceMTL.mm
@@ -478,6 +478,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 2d16c2b..183f3d1 100644
--- a/src/dawn/native/metal/ShaderModuleMTL.mm
+++ b/src/dawn/native/metal/ShaderModuleMTL.mm
@@ -360,6 +360,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 7c6eee9..fc0c67e 100644
--- a/src/dawn/native/vulkan/PhysicalDeviceVk.cpp
+++ b/src/dawn/native/vulkan/PhysicalDeviceVk.cpp
@@ -1087,6 +1087,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
@@ -1411,6 +1419,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
+}
+
 bool PhysicalDevice::MayBeArmProprietary() const {
     if (!gpu_info::IsARM(GetVendorId())) {
         return false;
diff --git a/src/dawn/native/vulkan/PhysicalDeviceVk.h b/src/dawn/native/vulkan/PhysicalDeviceVk.h
index 9f305e8..8ab8b9d 100644
--- a/src/dawn/native/vulkan/PhysicalDeviceVk.h
+++ b/src/dawn/native/vulkan/PhysicalDeviceVk.h
@@ -68,6 +68,7 @@
     bool IsAndroidHuawei() const;
     bool IsPixel10() const;
     bool IsSwiftshader() const;
+    bool IsWindowsAMD() const;
 
     // Check using VkDriverId, which is available in Vk 1.2 or an extension.
     bool IsIntelMesa() const;
diff --git a/src/dawn/native/vulkan/ShaderModuleVk.cpp b/src/dawn/native/vulkan/ShaderModuleVk.cpp
index 55e56a5..5764ab0 100644
--- a/src/dawn/native/vulkan/ShaderModuleVk.cpp
+++ b/src/dawn/native/vulkan/ShaderModuleVk.cpp
@@ -337,6 +337,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 3a1032c..b673a5a 100644
--- a/src/dawn/tests/end2end/SubgroupsTests.cpp
+++ b/src/dawn/tests/end2end/SubgroupsTests.cpp
@@ -280,6 +280,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";

diff --git a/src/dawn/tests/end2end/SubgroupsTests.cpp b/src/dawn/tests/end2end/SubgroupsTests.cpp
index 3a1032c..b673a5a 100644
--- a/src/dawn/tests/end2end/SubgroupsTests.cpp
+++ b/src/dawn/tests/end2end/SubgroupsTests.cpp
@@ -280,6 +280,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/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