From 81eda65489d8a1781ee4a076f8a3d2dceda03835 Mon Sep 17 00:00:00 2001
From: Peter McNeeley <petermcneeley@google.com>
Date: Fri, 03 Oct 2025 14:34:11 -0700
Subject: [PATCH] [tint] Polyfill case switch with if

Bug: 443906252
Change-Id: Ibec456e950717af8b6a2e3ff192291703026f7d9
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/262994
Reviewed-by: dan sinclair <dsinclair@chromium.org>
Reviewed-by: James Price <jrprice@google.com>
Commit-Queue: Peter McNeeley <petermcneeley@google.com>
---

diff --git a/src/dawn/native/Toggles.cpp b/src/dawn/native/Toggles.cpp
index f7cc6c4..0b2a0e8 100644
--- a/src/dawn/native/Toggles.cpp
+++ b/src/dawn/native/Toggles.cpp
@@ -566,6 +566,10 @@
       "and unpack4xU8() on D3D12 backends. Note that these functions are always polyfilled on all "
       "other backends right now.",
       "https://crbug.com/tint/1497", ToggleStage::Device}},
+    {Toggle::VulkanPolyfillSwitchWithIf,
+     {"vulkan_polyfill_switch_with_if",
+      "Polyfill switch statements with if/else statements on Vulkan.",
+      "https://crbug.com/443906252", ToggleStage::Device}},
     {Toggle::ExposeWGSLTestingFeatures,
      {"expose_wgsl_testing_features",
       "Make the Instance expose the ChromiumTesting* features for testing of "
diff --git a/src/dawn/native/Toggles.h b/src/dawn/native/Toggles.h
index c4c299a..090482b 100644
--- a/src/dawn/native/Toggles.h
+++ b/src/dawn/native/Toggles.h
@@ -138,6 +138,7 @@
     PolyfillPackUnpack4x8Norm,
     EnableSubgroupsIntelGen9,
     D3D12PolyFillPackUnpack4x8,
+    VulkanPolyfillSwitchWithIf,
     ExposeWGSLTestingFeatures,
     ExposeWGSLExperimentalFeatures,
     DisablePolyfillsOnIntegerDivisonAndModulo,
diff --git a/src/dawn/native/vulkan/PhysicalDeviceVk.cpp b/src/dawn/native/vulkan/PhysicalDeviceVk.cpp
index 8467b71..3ac8785 100644
--- a/src/dawn/native/vulkan/PhysicalDeviceVk.cpp
+++ b/src/dawn/native/vulkan/PhysicalDeviceVk.cpp
@@ -979,6 +979,9 @@
         // TODO(crbug.com/437807243): If newer driver version without bug is released then we can
         // gate this on driver version.
         deviceToggles->Default(Toggle::VulkanIncompletePipelineCacheWorkaround, true);
+
+        // crbug.com/443906252: Polyfill for case switch with large ranges.
+        deviceToggles->Default(Toggle::VulkanPolyfillSwitchWithIf, true);
     }
 
     if (IsAndroidARM()) {
diff --git a/src/dawn/native/vulkan/ShaderModuleVk.cpp b/src/dawn/native/vulkan/ShaderModuleVk.cpp
index 08c8e5e..468a630 100644
--- a/src/dawn/native/vulkan/ShaderModuleVk.cpp
+++ b/src/dawn/native/vulkan/ShaderModuleVk.cpp
@@ -293,6 +293,8 @@
         GetDevice()->IsToggleEnabled(Toggle::PolyFillPacked4x8DotProduct);
     req.tintOptions.polyfill_pack_unpack_4x8_norm =
         GetDevice()->IsToggleEnabled(Toggle::PolyfillPackUnpack4x8Norm);
+    req.tintOptions.polyfill_case_switch =
+        GetDevice()->IsToggleEnabled(Toggle::VulkanPolyfillSwitchWithIf);
     req.tintOptions.polyfill_subgroup_broadcast_f16 =
         GetDevice()->IsToggleEnabled(Toggle::EnableSubgroupsIntelGen9);
     req.tintOptions.disable_polyfill_integer_div_mod =
diff --git a/src/dawn/tests/end2end/PolyfillBuiltinSimpleTests.cpp b/src/dawn/tests/end2end/PolyfillBuiltinSimpleTests.cpp
index 53ca6ac..97a7b3d 100644
--- a/src/dawn/tests/end2end/PolyfillBuiltinSimpleTests.cpp
+++ b/src/dawn/tests/end2end/PolyfillBuiltinSimpleTests.cpp
@@ -229,6 +229,142 @@
     EXPECT_BUFFER_U32_RANGE_EQ(expected.data(), output, 0, expected.size());
 }
 
+TEST_P(PolyfillBuiltinSimpleTests, CaseSwitchToIf) {
+    std::string kShaderCode = R"(
+    struct Data { values: array<i32> };
+    @group(0) @binding(0) var<storage, read> input_data: Data;
+    @group(0) @binding(1) var<storage, read_write> output_data: Data;
+
+    @compute @workgroup_size(4)
+    fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {
+        var input_ = input_data.values[global_id.x];
+        var ret = 0i;
+        switch( input_ ) {
+            case 1: {
+                ret = 3;
+            }
+            case 2:{
+                ret = 7;
+            }
+            case -2147483648:{
+                ret = 71;
+            }
+            case 123, 87:{
+                ret = 11;
+            }
+            case -1:{
+                ret = 33;
+            }
+            default {
+                ret = 82;
+            }
+        }
+        output_data.values[global_id.x]  = ret;
+    }
+    )";
+
+    wgpu::ComputePipeline pipeline = CreateComputePipeline(kShaderCode);
+    uint32_t kDefaultVal = 0;
+    std::vector<uint32_t> init_input = {uint32_t(std::numeric_limits<int32_t>::lowest()),
+                                        uint32_t(-15), 17, 123};
+
+    wgpu::Buffer input = CreateBuffer(init_input);
+    wgpu::Buffer output = CreateBuffer(4, kDefaultVal);
+    wgpu::BindGroup bindGroup =
+        utils::MakeBindGroup(device, pipeline.GetBindGroupLayout(0), {{0, input}, {1, output}});
+
+    wgpu::CommandBuffer commands;
+    {
+        wgpu::CommandEncoder encoder = device.CreateCommandEncoder();
+        wgpu::ComputePassEncoder pass = encoder.BeginComputePass();
+        pass.SetPipeline(pipeline);
+        pass.SetBindGroup(0, bindGroup);
+        pass.DispatchWorkgroups(64);
+        pass.End();
+        commands = encoder.Finish();
+    }
+
+    queue.Submit(1, &commands);
+    std::vector<uint32_t> expected = {71, 82, 82, 11};
+
+    EXPECT_BUFFER_U32_RANGE_EQ(expected.data(), output, 0, expected.size());
+}
+
+TEST_P(PolyfillBuiltinSimpleTests, CaseSwitchToIfComplex) {
+    std::string kShaderCode = R"(
+    @group(0) @binding(0) var<storage, read> input_data: array<i32>;
+    @group(0) @binding(1) var<storage, read_write> output_data: array<i32>;
+
+    @compute @workgroup_size(4)
+    fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {
+        var input_ = input_data[global_id.x];
+        var ret = 0i;
+        switch( input_ ) {
+            case 1: {
+                ret = 3;
+            }
+            case -2:{
+                switch(input_){
+                    case 1: {
+                        ret = 3;
+                    }
+                    case -2:{
+                        ret = 4;
+                    }
+                    default{
+                        ret = 99;
+                    }
+                }
+                break;
+                ret = 7;
+            }
+            case -2147483648:{
+                if(input_ == 17){
+                    ret = 71;
+                    break;
+                }
+                ret = 13;
+            }
+            case 3, 5:{
+                if(input_ == 3){
+                    break;
+                }
+                ret = 11;
+            }
+            default {
+                ret = 82;
+            }
+        }
+        output_data[global_id.x]  = ret;
+    }
+    )";
+
+    wgpu::ComputePipeline pipeline = CreateComputePipeline(kShaderCode);
+    uint32_t kDefaultVal = 0;
+    std::vector<uint32_t> init_input = {uint32_t(std::numeric_limits<int32_t>::lowest()),
+                                        uint32_t(-2), 3, 5};
+    std::vector<uint32_t> expected = {13, 4, 0, 11};
+    wgpu::Buffer input = CreateBuffer(init_input);
+    wgpu::Buffer output = CreateBuffer(4, kDefaultVal);
+    wgpu::BindGroup bindGroup =
+        utils::MakeBindGroup(device, pipeline.GetBindGroupLayout(0), {{0, input}, {1, output}});
+
+    wgpu::CommandBuffer commands;
+    {
+        wgpu::CommandEncoder encoder = device.CreateCommandEncoder();
+        wgpu::ComputePassEncoder pass = encoder.BeginComputePass();
+        pass.SetPipeline(pipeline);

diff --git a/src/dawn/tests/end2end/PolyfillBuiltinSimpleTests.cpp b/src/dawn/tests/end2end/PolyfillBuiltinSimpleTests.cpp
index 53ca6ac..97a7b3d 100644
--- a/src/dawn/tests/end2end/PolyfillBuiltinSimpleTests.cpp
+++ b/src/dawn/tests/end2end/PolyfillBuiltinSimpleTests.cpp
@@ -229,6 +229,142 @@
     EXPECT_BUFFER_U32_RANGE_EQ(expected.data(), output, 0, expected.size());
 }
 
+TEST_P(PolyfillBuiltinSimpleTests, CaseSwitchToIf) {
+    std::string kShaderCode = R"(
+    struct Data { values: array<i32> };
+    @group(0) @binding(0) var<storage, read> input_data: Data;
+    @group(0) @binding(1) var<storage, read_write> output_data: Data;
+
+    @compute @workgroup_size(4)
+    fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {
+        var input_ = input_data.values[global_id.x];
+        var ret = 0i;
+        switch( input_ ) {
+            case 1: {
+                ret = 3;
+            }
+            case 2:{
+                ret = 7;
+            }
+            case -2147483648:{
+                ret = 71;
+            }
+            case 123, 87:{
+                ret = 11;
+            }
+            case -1:{
+                ret = 33;
+            }
+            default {
+                ret = 82;
+            }
+        }
+        output_data.values[global_id.x]  = ret;
+    }
+    )";
+
+    wgpu::ComputePipeline pipeline = CreateComputePipeline(kShaderCode);
+    uint32_t kDefaultVal = 0;
+    std::vector<uint32_t> init_input = {uint32_t(std::numeric_limits<int32_t>::lowest()),
+                                        uint32_t(-15), 17, 123};
+
+    wgpu::Buffer input = CreateBuffer(init_input);
+    wgpu::Buffer output = CreateBuffer(4, kDefaultVal);
+    wgpu::BindGroup bindGroup =
+        utils::MakeBindGroup(device, pipeline.GetBindGroupLayout(0), {{0, input}, {1, output}});
+
+    wgpu::CommandBuffer commands;
+    {
+        wgpu::CommandEncoder encoder = device.CreateCommandEncoder();
+        wgpu::ComputePassEncoder pass = encoder.BeginComputePass();
+        pass.SetPipeline(pipeline);
+        pass.SetBindGroup(0, bindGroup);
+        pass.DispatchWorkgroups(64);
+        pass.End();
+        commands = encoder.Finish();
+    }
+
+    queue.Submit(1, &commands);
+    std::vector<uint32_t> expected = {71, 82, 82, 11};
+
+    EXPECT_BUFFER_U32_RANGE_EQ(expected.data(), output, 0, expected.size());
+}
+
+TEST_P(PolyfillBuiltinSimpleTests, CaseSwitchToIfComplex) {
+    std::string kShaderCode = R"(
+    @group(0) @binding(0) var<storage, read> input_data: array<i32>;
+    @group(0) @binding(1) var<storage, read_write> output_data: array<i32>;
+
+    @compute @workgroup_size(4)
+    fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {
+        var input_ = input_data[global_id.x];
+        var ret = 0i;
+        switch( input_ ) {
+            case 1: {
+                ret = 3;
+            }
+            case -2:{
+                switch(input_){
+                    case 1: {
+                        ret = 3;
+                    }
+                    case -2:{
+                        ret = 4;
+                    }
+                    default{
+                        ret = 99;
+                    }
+                }
+                break;
+                ret = 7;
+            }
+            case -2147483648:{
+                if(input_ == 17){
+                    ret = 71;
+                    break;
+                }
+                ret = 13;
+            }
+            case 3, 5:{
+                if(input_ == 3){
+                    break;
+                }
+                ret = 11;
+            }
+            default {
+                ret = 82;
+            }
+        }
+        output_data[global_id.x]  = ret;
+    }
+    )";
+
+    wgpu::ComputePipeline pipeline = CreateComputePipeline(kShaderCode);
+    uint32_t kDefaultVal = 0;
+    std::vector<uint32_t> init_input = {uint32_t(std::numeric_limits<int32_t>::lowest()),
+                                        uint32_t(-2), 3, 5};
+    std::vector<uint32_t> expected = {13, 4, 0, 11};
+    wgpu::Buffer input = CreateBuffer(init_input);
+    wgpu::Buffer output = CreateBuffer(4, kDefaultVal);
+    wgpu::BindGroup bindGroup =
+        utils::MakeBindGroup(device, pipeline.GetBindGroupLayout(0), {{0, input}, {1, output}});
+
+    wgpu::CommandBuffer commands;
+    {
+        wgpu::CommandEncoder encoder = device.CreateCommandEncoder();
+        wgpu::ComputePassEncoder pass = encoder.BeginComputePass();
+        pass.SetPipeline(pipeline);
+        pass.SetBindGroup(0, bindGroup);
+        pass.DispatchWorkgroups(64);
+        pass.End();
+        commands = encoder.Finish();
+    }
+
+    queue.Submit(1, &commands);
+
+    EXPECT_BUFFER_U32_RANGE_EQ(expected.data(), output, 0, expected.size());
+}
+
 DAWN_INSTANTIATE_TEST(PolyfillBuiltinSimpleTests,
                       D3D12Backend(),
                       D3D11Backend(),
@@ -237,6 +373,7 @@
                       D3D12Backend({"scalarize_max_min_clamp"}),
                       MetalBackend({"scalarize_max_min_clamp"}),
                       VulkanBackend({"scalarize_max_min_clamp"}),
+                      VulkanBackend({"vulkan_polyfill_switch_with_if"}),
                       D3D11Backend({"scalarize_max_min_clamp"}),
                       OpenGLESBackend());
diff --git a/src/tint/lang/spirv/writer/raise/case_switch_to_if_else_test.cc b/src/tint/lang/spirv/writer/raise/case_switch_to_if_else_test.cc
new file mode 100644
index 0000000..738d69f
--- /dev/null
+++ b/src/tint/lang/spirv/writer/raise/case_switch_to_if_else_test.cc
@@ -0,0 +1,1375 @@
+// Copyright 2025 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/spirv/writer/raise/case_switch_to_if_else.h"
+
+#include <utility>
+
+#include "src/tint/lang/core/ir/transform/helper_test.h"
+
+namespace tint::spirv::writer::raise {
+namespace {
+
+using namespace tint::core::fluent_types;     // NOLINT
+using namespace tint::core::number_suffixes;  // NOLINT
+
+using SpirvWriter_CaseSwitchToIfElseTest = core::ir::transform::TransformTest;
+
+TEST_F(SpirvWriter_CaseSwitchToIfElseTest, BasicSwitch) {
+    auto* cond = b.FunctionParam("param0", ty.i32());
+    auto* func = b.Function("foo", ty.void_());
+    func->SetParams({cond});
+    b.Append(func->Block(), [&] {
+        auto* s = b.Switch(cond);
+        b.Append(b.Case(s, {b.Constant(-1_i)}), [&] {  //
+            b.Return(func);
+        });
+        b.Append(b.Case(s, {b.Constant(2_i)}), [&] {  //
+            b.Return(func);
+        });
+        b.Append(b.DefaultCase(s), [&] {  //
+            b.Return(func);
+        });
+        b.Unreachable();
+    });
+
+    auto* src = R"(
+%foo = func(%param0:i32):void {
+  $B1: {
+    switch %param0 [c: (-1i, $B2), c: (2i, $B3), c: (default, $B4)] {  # switch_1
+      $B2: {  # case
+        ret
+      }
+      $B3: {  # case
+        ret
+      }
+      $B4: {  # case
+        ret
+      }
+    }
+    unreachable
+  }
+}
+)";
+    EXPECT_EQ(src, str());
+
+    auto* expect = R"(
+%foo = func(%param0:i32):void {
+  $B1: {
+    switch %param0 [c: (default, $B2)] {  # switch_1
+      $B2: {  # case
+        %3:bool = eq %param0, -1i
+        if %3 [t: $B3] {  # if_1
+          $B3: {  # true
+            ret
+          }
+        }
+        %4:bool = eq %param0, 2i
+        if %4 [t: $B4] {  # if_2
+          $B4: {  # true
+            ret
+          }
+        }
+        if true [t: $B5] {  # if_3
+          $B5: {  # true
+            ret
+          }
+        }
+        unreachable
+      }
+    }
+    unreachable
+  }
+}
+)";
+
+    Run(CaseSwitchToIfElse);
+
+    EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvWriter_CaseSwitchToIfElseTest, ReorderedBasicSwitch) {
+    auto* cond = b.FunctionParam("param0", ty.i32());
+    auto* func = b.Function("foo", ty.void_());
+    func->SetParams({cond});
+    b.Append(func->Block(), [&] {
+        auto* s = b.Switch(cond);
+        b.Append(b.DefaultCase(s), [&] {  //
+            b.Return(func);
+        });
+        b.Append(b.Case(s, {b.Constant(2_i)}), [&] {  //
+            b.Return(func);
+        });
+
+        b.Append(b.Case(s, {b.Constant(-1_i)}), [&] {  //
+            b.Return(func);
+        });
+
+        b.Unreachable();
+    });
+
+    auto* src = R"(
+%foo = func(%param0:i32):void {
+  $B1: {
+    switch %param0 [c: (default, $B2), c: (2i, $B3), c: (-1i, $B4)] {  # switch_1
+      $B2: {  # case
+        ret
... (truncated)