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/tint/lang/core/ir/switch.h b/src/tint/lang/core/ir/switch.h
index b6c44d9..55cf61f 100644
--- a/src/tint/lang/core/ir/switch.h
+++ b/src/tint/lang/core/ir/switch.h
@@ -106,6 +106,13 @@
     /// @returns the switch cases
     Vector<Case, 4>& Cases() { return cases_; }
 
+    /// @returns the switch cases by moving them
+    Vector<Case, 4> TakeCases() {
+        auto rtn = std::move(cases_);
+        cases_.Clear();
+        return rtn;
+    }
+
     /// @returns the switch cases
     VectorRef<Case> Cases() const { return cases_; }
 
diff --git a/src/tint/lang/spirv/writer/common/options.h b/src/tint/lang/spirv/writer/common/options.h
index 521faf3..c4a98b0 100644
--- a/src/tint/lang/spirv/writer/common/options.h
+++ b/src/tint/lang/spirv/writer/common/options.h
@@ -118,6 +118,9 @@
     /// `unpack4x8unorm` builtins
     bool polyfill_pack_unpack_4x8_norm = false;
 
+    /// Set to `true` to generate a polyfill for switch statements using if/else statements.
+    bool polyfill_case_switch = false;
+
     /// Set to `true` to generate a polyfill clamp of `id` param of subgroupShuffle to within the
     /// spec max subgroup size.
     bool subgroup_shuffle_clamped = false;
@@ -163,6 +166,7 @@
                  pass_matrix_by_pointer,
                  polyfill_dot_4x8_packed,
                  polyfill_pack_unpack_4x8_norm,
+                 polyfill_case_switch,
                  subgroup_shuffle_clamped,
                  polyfill_subgroup_broadcast_f16,
                  disable_polyfill_integer_div_mod,
diff --git a/src/tint/lang/spirv/writer/printer/printer.cc b/src/tint/lang/spirv/writer/printer/printer.cc
index b237f7b..dff0eac 100644
--- a/src/tint/lang/spirv/writer/printer/printer.cc
+++ b/src/tint/lang/spirv/writer/printer/printer.cc
@@ -1355,8 +1355,11 @@
                 }
                 break;
             }
-            default:
-                TINT_IR_UNIMPLEMENTED(ir_) << binary->Op();
+            default: {
+                TINT_IR_UNIMPLEMENTED(ir_)
+                    << "\"" << binary->Op() << "\"  (enum value=" << static_cast<int>(binary->Op())
+                    << ")";
+            }
         }
 
         // Emit the instruction.
diff --git a/src/tint/lang/spirv/writer/raise/BUILD.cmake b/src/tint/lang/spirv/writer/raise/BUILD.cmake
index 0580e4c..8150d9e 100644
--- a/src/tint/lang/spirv/writer/raise/BUILD.cmake
+++ b/src/tint/lang/spirv/writer/raise/BUILD.cmake
@@ -43,6 +43,8 @@
 tint_add_target(tint_lang_spirv_writer_raise lib
   lang/spirv/writer/raise/builtin_polyfill.cc
   lang/spirv/writer/raise/builtin_polyfill.h
+  lang/spirv/writer/raise/case_switch_to_if_else.cc
+  lang/spirv/writer/raise/case_switch_to_if_else.h
   lang/spirv/writer/raise/expand_implicit_splats.cc
   lang/spirv/writer/raise/expand_implicit_splats.h
   lang/spirv/writer/raise/fork_explicit_layout_types.cc
@@ -79,6 +81,7 @@
   tint_lang_spirv_intrinsic
   tint_lang_spirv_ir
   tint_lang_spirv_type
+  tint_lang_wgsl_ast
   tint_utils
   tint_utils_containers
   tint_utils_diagnostic
@@ -116,6 +119,7 @@
 ################################################################################
 tint_add_target(tint_lang_spirv_writer_raise_test test
   lang/spirv/writer/raise/builtin_polyfill_test.cc
+  lang/spirv/writer/raise/case_switch_to_if_else_test.cc
   lang/spirv/writer/raise/expand_implicit_splats_test.cc
   lang/spirv/writer/raise/fork_explicit_layout_types_test.cc
   lang/spirv/writer/raise/handle_matrix_arithmetic_test.cc
diff --git a/src/tint/lang/spirv/writer/raise/case_switch_to_if_else.cc b/src/tint/lang/spirv/writer/raise/case_switch_to_if_else.cc
new file mode 100644
index 0000000..4026054
--- /dev/null
+++ b/src/tint/lang/spirv/writer/raise/case_switch_to_if_else.cc
@@ -0,0 +1,168 @@
+// 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 <algorithm>
+#include <cstdint>
+#include <limits>
+#include <utility>
+
+#include "src/tint/lang/core/ir/builder.h"
+#include "src/tint/lang/core/ir/control_instruction.h"
+#include "src/tint/lang/core/ir/exit_switch.h"
+#include "src/tint/lang/core/ir/if.h"
+#include "src/tint/lang/core/ir/instruction.h"
+#include "src/tint/lang/core/ir/module.h"
+#include "src/tint/lang/core/ir/switch.h"
+#include "src/tint/lang/core/ir/traverse.h"
+#include "src/tint/lang/core/ir/validator.h"
+#include "src/tint/lang/core/ir/value.h"
+#include "src/tint/lang/core/unary_op.h"
+#include "src/tint/lang/wgsl/ast/case_selector.h"
+#include "src/tint/utils/containers/transform.h"
+
+using namespace tint::core::fluent_types;  // NOLINT
+
+namespace tint::spirv::writer::raise {
+
+namespace {
+
+const core::ir::Capabilities kCaseSwitchToIfElseCapabilities{
+    core::ir::Capability::kAllowDuplicateBindings,
+    core::ir::Capability::kAllowAnyInputAttachmentIndexType,
+    core::ir::Capability::kAllowNonCoreTypes,
+};
+
+/// PIMPL state for the transform, for a single function.
+struct State {
+    /// The IR module.
+    core::ir::Module& ir;
+
+    /// The IR builder.
+    core::ir::Builder b{ir};
+
+    /// The type manager.
+    core::type::Manager& ty{ir.Types()};
+
+    /// Process the IR
+    void Process() {
+        Vector<core::ir::Switch*, 4> worklist;
+        for (auto* inst : ir.Instructions()) {
+            if (auto* s = inst->As<core::ir::Switch>()) {
+                // Even though the switch param could be signed we use u32 as that is likely the
+                // internal representation of the compiler.
+                uint32_t max_sel_case = 0u;
+                uint32_t min_sel_case = std::numeric_limits<uint32_t>().max();
+                for (auto& c : s->Cases()) {
+                    for (auto& sel : c.selectors) {
+                        if (!sel.IsDefault() && sel.val && sel.val->Value()) {
+                            auto val = sel.val->Value()->ValueAs<u32>().value;
+                            max_sel_case = std::max(max_sel_case, val);
+                            min_sel_case = std::min(min_sel_case, val);
+                        }
+                    }
+                }
+
+                // Our concern is around handling of signed range calculations (vs unsigned). Any
+                // range that gets close we will polyfill.
+                const uint32_t kSignedRangeLimit =
+                    static_cast<uint32_t>(std::numeric_limits<int32_t>().max() - 1);
+                if ((max_sel_case - min_sel_case) >= kSignedRangeLimit) {
+                    worklist.Push(s);
+                }
+            }
+        }
+
+        for (auto* s : worklist) {
+            auto* switch_cond = s->Condition();
+            Vector<core::ir::Value*, 4> conditions;
+            core::ir::Switch::Case* default_case = nullptr;
+            // We take the cases here because we're going to attach the case block to an `if`
+            // statement and the switch will now be replaced with a single default case block.
+            auto cases = s->TakeCases();
+            auto* def = b.DefaultCase(s);
+            b.Append(def, [&] {
+                for (auto& c : cases) {
+                    // Default block is required by spec. It will need to be treated special.
+                    // It is possible that default case will also have non default selectors.
+                    // These additional selectors are superfluous as they will just form one
+                    // default.
+                    bool found_default = false;
+                    for (auto& sel : c.selectors) {
+                        if (sel.IsDefault()) {
+                            default_case = &c;
+                            found_default = true;
+                            break;
+                        }
+                    }
+                    if (found_default) {
+                        continue;
+                    }
+
+                    core::ir::Value* case_cond = nullptr;
+                    for (auto& sel : c.selectors) {
+                        auto* curr_selector =
+                            b.Equal(ty.bool_(), switch_cond, sel.val->As<core::ir::Value>());
+                        if (case_cond) {
+                            case_cond = b.Or(ty.bool_(), curr_selector, case_cond)->Result();
+                        } else {
+                            case_cond = curr_selector->Result();
+                        }
+                    }
+                    conditions.Push(case_cond);
+                    auto* if_cond = b.If(case_cond);
+                    if_cond->SetTrue(c.block);
+                }
+
+                TINT_ASSERT(default_case);
+                // Special handling required for default case. All non-default cases have exited the
+                // switch by this point so the only possibility that remains is the default.
+                auto* if_cond = b.If(b.Constant(true));
+                if_cond->SetTrue(default_case->block);
+            });
+
+            b.Append(s->Cases()[0].block, [&] { b.Unreachable(); });
+        }
+    }
+};
+
+}  // namespace
+
+Result<SuccessType> CaseSwitchToIfElse(core::ir::Module& ir) {
+    auto result =
+        ValidateAndDumpIfNeeded(ir, "spirv.CaseSwitchToIfElse", kCaseSwitchToIfElseCapabilities);
+    if (result != Success) {
+        return result;
+    }
+
+    State{ir}.Process();
+
+    return Success;
+}
+
+}  // namespace tint::spirv::writer::raise
diff --git a/src/tint/lang/spirv/writer/raise/case_switch_to_if_else.h b/src/tint/lang/spirv/writer/raise/case_switch_to_if_else.h
new file mode 100644
index 0000000..ae846ff
--- /dev/null
+++ b/src/tint/lang/spirv/writer/raise/case_switch_to_if_else.h
@@ -0,0 +1,48 @@
+// 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.
+
+#ifndef SRC_TINT_LANG_SPIRV_WRITER_RAISE_CASE_SWITCH_TO_IF_ELSE_H_
+#define SRC_TINT_LANG_SPIRV_WRITER_RAISE_CASE_SWITCH_TO_IF_ELSE_H_
+
+#include "src/tint/utils/result.h"
+
+// Forward declarations.
+namespace tint::core::ir {
+class Module;
+}
+
+namespace tint::spirv::writer::raise {
+
+/// CaseSwitchToIfElse is a transform that converts case switch statements to a series of if
+/// statements with the default being an emulated else condition.
+/// @param module the module to transform
+/// @returns success or failure
+Result<SuccessType> CaseSwitchToIfElse(core::ir::Module& module);
+
+}  // namespace tint::spirv::writer::raise
+
+#endif  // SRC_TINT_LANG_SPIRV_WRITER_RAISE_CASE_SWITCH_TO_IF_ELSE_H_
diff --git a/src/tint/lang/spirv/writer/raise/raise.cc b/src/tint/lang/spirv/writer/raise/raise.cc
index 1cbca64..3e4ef75 100644
--- a/src/tint/lang/spirv/writer/raise/raise.cc
+++ b/src/tint/lang/spirv/writer/raise/raise.cc
@@ -52,6 +52,7 @@
 #include "src/tint/lang/core/type/f32.h"
 #include "src/tint/lang/spirv/writer/common/option_helpers.h"
 #include "src/tint/lang/spirv/writer/raise/builtin_polyfill.h"
+#include "src/tint/lang/spirv/writer/raise/case_switch_to_if_else.h"
 #include "src/tint/lang/spirv/writer/raise/expand_implicit_splats.h"
 #include "src/tint/lang/spirv/writer/raise/fork_explicit_layout_types.h"
 #include "src/tint/lang/spirv/writer/raise/handle_matrix_arithmetic.h"
@@ -204,6 +205,9 @@
     // kAllowAnyInputAttachmentIndexType required after ExpandImplicitSplats
     RUN_TRANSFORM(raise::HandleMatrixArithmetic, module);
     RUN_TRANSFORM(raise::MergeReturn, module);
+    if (options.polyfill_case_switch) {
+        RUN_TRANSFORM(raise::CaseSwitchToIfElse, module);
+    }
     RUN_TRANSFORM(raise::RemoveUnreachableInLoopContinuing, module);
     RUN_TRANSFORM(
         raise::ShaderIO, module,
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) {
... (file diff truncated)

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)