CVE-2026-18015

Low chrome Logic Error 📄 Reporter bug report 🔧 Commit mapped

Overview

Low

Severity

—

CVSS

Exploited ITW

Fixed

Fix Status

ImpactInappropriate implementation in Tint

DescriptionInappropriate implementation in Tint

ComponentTint

Bug ClassLogic Error

Tracker523698428

Fix commited24c432dcb8 (dawn) +97/-8

CISA KEVNot listed

CreditedGoogle

Disclosed2026-07-29

Changed Functions

Function	Change	Notes
`TEST_P` src/dawn/tests/end2end/ShaderTests.cpp	modified
`if` src/tint/lang/msl/writer/raise/binary_polyfill.cc	modified
`for` src/tint/lang/msl/writer/raise/binary_polyfill.cc	modified
`TEST_F` src/tint/lang/msl/writer/raise/binary_polyfill_test.cc	modified

Files Changed

src/dawn/tests/end2end/ShaderTests.cpp
src/tint/lang/msl/writer/raise/binary_polyfill.cc
src/tint/lang/msl/writer/raise/binary_polyfill_test.cc

Patch

view commit

From ed24c432dcb80344927e36905a860285195d3e45 Mon Sep 17 00:00:00 2001
From: James Price <[email protected]>
Date: Mon, 15 Jun 2026 14:20:53 -0700
Subject: [PATCH] [msl] Apply u32 div/mod polyfill to vectors

The bug did not reproduce for vector operations, but we should
polyfill anyway just in case other optimizations scalarize the
operations and then trigger the bug that way.

Fixed: 523698428
Change-Id: I2f844c3e736130af63c3ab3d3fc17faa7dc3e246
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/316895
Auto-Submit: James Price <[email protected]>
Reviewed-by: dan sinclair <[email protected]>
Commit-Queue: James Price <[email protected]>
---

diff --git a/src/dawn/tests/end2end/ShaderTests.cpp b/src/dawn/tests/end2end/ShaderTests.cpp
index af2b702..f8d4553 100644
--- a/src/dawn/tests/end2end/ShaderTests.cpp
+++ b/src/dawn/tests/end2end/ShaderTests.cpp
@@ -1276,6 +1276,55 @@
     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_Vector) {
+    wgpu::ComputePipelineDescriptor cDesc;
+    cDesc.compute.module = utils::CreateShaderModule(device, R"(
+        @group(0) @binding(0)
+        var<storage, read_write> value: vec4u;
+
+        @compute @workgroup_size(1u)
+        fn main() {
+            let input = value;
+            value = ((input * input) >> vec4u(16u)) % 3u;
+        }
+    )");
+    wgpu::ComputePipeline pipeline = device.CreateComputePipeline(&cDesc);
+
+    wgpu::BufferDescriptor bufDesc;
+    bufDesc.size = 4 * sizeof(uint32_t);
+    bufDesc.usage =
+        wgpu::BufferUsage::Storage | wgpu::BufferUsage::CopySrc | wgpu::BufferUsage::CopyDst;
+    wgpu::Buffer buf = device.CreateBuffer(&bufDesc);
+
+    // Write 0x10004 to the buffer.
+    uint32_t inputValues[] = {0x10004, 0x10004, 0x10004, 0x10004};
+    queue.WriteBuffer(buf, 0, inputValues, 4 * sizeof(uint32_t));
+
+    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);
+    EXPECT_BUFFER_U32_EQ(2, buf, 4);
+    EXPECT_BUFFER_U32_EQ(2, buf, 8);
+    EXPECT_BUFFER_U32_EQ(2, buf, 12);
+}
+
 // 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.cc b/src/tint/lang/msl/writer/raise/binary_polyfill.cc
index 9f7e329..261e108 100644
--- a/src/tint/lang/msl/writer/raise/binary_polyfill.cc
+++ b/src/tint/lang/msl/writer/raise/binary_polyfill.cc
@@ -52,7 +52,7 @@
     void Process() {
         // Find the binary operators that need replacing.
         Vector<core::ir::CoreBinary*, 4> fmod_worklist;
-        Vector<core::ir::CoreBinary*, 4> umod_worklist;
+        Vector<core::ir::CoreBinary*, 4> udivmod_worklist;
         Vector<core::ir::CoreBinary*, 4> logical_bool_worklist;
 
         for (auto* inst : ir.Instructions()) {
@@ -62,9 +62,9 @@
                 if (op == core::BinaryOp::kModulo && lhs_type->IsFloatScalarOrVector()) {
                     fmod_worklist.Push(binary);
                 } else if ((op == core::BinaryOp::kModulo || op == core::BinaryOp::kDivide) &&
-                           lhs_type->Is<core::type::U32>()) {
+                           lhs_type->DeepestElement()->Is<core::type::U32>()) {
                     if (config.fix_u32_div_mod) {
-                        umod_worklist.Push(binary);
+                        udivmod_worklist.Push(binary);
                     }
                 } else if ((op == core::BinaryOp::kAnd || op == core::BinaryOp::kOr) &&
                            lhs_type->IsBoolScalarOrVector()) {
@@ -77,8 +77,8 @@
         for (auto* fmod : fmod_worklist) {
             FMod(fmod);
         }
-        for (auto* umod : umod_worklist) {
-            UMod(umod);
+        for (auto* udivmod : udivmod_worklist) {
+            UDivMod(udivmod);
         }
         for (auto* logical_bool : logical_bool_worklist) {
             LogicalBool(logical_bool);
@@ -94,9 +94,10 @@
         binary->Destroy();
     }
 
-    /// Add a volatile zero to unsigned modulo binary instructions to work around a driver bug.
-    /// @param binary the unsigned integer modulo binary instruction
-    void UMod(core::ir::CoreBinary* binary) {
+    /// Add a volatile zero to unsigned divide and modulo binary instructions to work around a
+    /// driver bug.
+    /// @param binary the unsigned integer divide or modulo binary instruction
+    void UDivMod(core::ir::CoreBinary* binary) {
         b.InsertBefore(binary, [&] {
             auto* zero = b.Call<msl::ir::BuiltinCall>(ty.u32(), msl::BuiltinFn::kVolatileZero);
             binary->SetOperand(0u, b.Add(binary->LHS(), zero)->Result());
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 ead1565..7cb451f 100644
--- a/src/tint/lang/msl/writer/raise/binary_polyfill_test.cc
+++ b/src/tint/lang/msl/writer/raise/binary_polyfill_test.cc
@@ -398,5 +398,44 @@
     EXPECT_EQ(expect, str());
 }
 
+TEST_F(MslWriter_BinaryPolyfillTest, UDiv_Vector_WithPolyfill) {
+    auto* lhs = b.FunctionParam<vec4u>("lhs");
+    auto* rhs = b.FunctionParam<vec4u>("rhs");
+    auto* func = b.Function("foo", ty.vec4u());
+    func->SetParams({lhs, rhs});
+    b.Append(func->Block(), [&] {
+        auto* result = b.Divide(lhs, rhs);
+        b.Return(func, result);
+    });
+
+    auto* src = R"(
+%foo = func(%lhs:vec4<u32>, %rhs:vec4<u32>):vec4<u32> {
+  $B1: {
+    %4:vec4<u32> = div %lhs, %rhs
+    ret %4
+  }
+}
+)";
+    EXPECT_EQ(src, str());
+
+    auto* expect = R"(
+%foo = func(%lhs:vec4<u32>, %rhs:vec4<u32>):vec4<u32> {
+  $B1: {
+    %4:u32 = msl.volatile_zero
+    %5:vec4<u32> = add %lhs, %4
+    %6:vec4<u32> = div %5, %rhs
+    ret %6
+  }
+}
+)";
+
+    BinaryPolyfillConfig config{
+        .fix_u32_div_mod = true,
+    };
+    Run(BinaryPolyfill, config);
+
+    EXPECT_EQ(expect, str());
+}
+
 }  // namespace
 }  // namespace tint::msl::writer::raise

Loading diff…

Regression Test / PoC

shipped with the fix

diff --git a/src/dawn/tests/end2end/ShaderTests.cpp b/src/dawn/tests/end2end/ShaderTests.cpp
index af2b702..f8d4553 100644
--- a/src/dawn/tests/end2end/ShaderTests.cpp
+++ b/src/dawn/tests/end2end/ShaderTests.cpp
@@ -1276,6 +1276,55 @@
     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_Vector) {
+    wgpu::ComputePipelineDescriptor cDesc;
+    cDesc.compute.module = utils::CreateShaderModule(device, R"(
+        @group(0) @binding(0)
+        var<storage, read_write> value: vec4u;
+
+        @compute @workgroup_size(1u)
+        fn main() {
+            let input = value;
+            value = ((input * input) >> vec4u(16u)) % 3u;
+        }
+    )");
+    wgpu::ComputePipeline pipeline = device.CreateComputePipeline(&cDesc);
+
+    wgpu::BufferDescriptor bufDesc;
+    bufDesc.size = 4 * sizeof(uint32_t);
+    bufDesc.usage =
+        wgpu::BufferUsage::Storage | wgpu::BufferUsage::CopySrc | wgpu::BufferUsage::CopyDst;
+    wgpu::Buffer buf = device.CreateBuffer(&bufDesc);
+
+    // Write 0x10004 to the buffer.
+    uint32_t inputValues[] = {0x10004, 0x10004, 0x10004, 0x10004};
+    queue.WriteBuffer(buf, 0, inputValues, 4 * sizeof(uint32_t));
+
+    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);
+    EXPECT_BUFFER_U32_EQ(2, buf, 4);
+    EXPECT_BUFFER_U32_EQ(2, buf, 8);
+    EXPECT_BUFFER_U32_EQ(2, buf, 12);
+}
+
 // 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 ead1565..7cb451f 100644
--- a/src/tint/lang/msl/writer/raise/binary_polyfill_test.cc
+++ b/src/tint/lang/msl/writer/raise/binary_polyfill_test.cc
@@ -398,5 +398,44 @@
     EXPECT_EQ(expect, str());
 }
 
+TEST_F(MslWriter_BinaryPolyfillTest, UDiv_Vector_WithPolyfill) {
+    auto* lhs = b.FunctionParam<vec4u>("lhs");
+    auto* rhs = b.FunctionParam<vec4u>("rhs");
+    auto* func = b.Function("foo", ty.vec4u());
+    func->SetParams({lhs, rhs});
+    b.Append(func->Block(), [&] {
+        auto* result = b.Divide(lhs, rhs);
+        b.Return(func, result);
+    });
+
+    auto* src = R"(
+%foo = func(%lhs:vec4<u32>, %rhs:vec4<u32>):vec4<u32> {
+  $B1: {
+    %4:vec4<u32> = div %lhs, %rhs
+    ret %4
+  }
+}
+)";
+    EXPECT_EQ(src, str());
+
+    auto* expect = R"(
+%foo = func(%lhs:vec4<u32>, %rhs:vec4<u32>):vec4<u32> {
+  $B1: {
+    %4:u32 = msl.volatile_zero
+    %5:vec4<u32> = add %lhs, %4
+    %6:vec4<u32> = div %5, %rhs
+    ret %6
+  }
+}
+)";
+
+    BinaryPolyfillConfig config{
+        .fix_u32_div_mod = true,
+    };
+    Run(BinaryPolyfill, config);
+
+    EXPECT_EQ(expect, str());
+}
+
 }  // namespace
 }  // namespace tint::msl::writer::raise

Loading diff…

Original Bug Report

reported by [email protected]

Potential GPU sandbox escape via unmitigated Apple Silicon u32 vector div/mod in Tint

Project Fortify, an experimental security project, has identified the following potential security issue. If you’re a feature owner CC-ed on this bug, please do your best to review these reports. Please see https://chromium.googlesource.com/chromium/src/+/main/docs/security/ai-generated-security-bugs-faq.md for more information.

Overview: The MSL backend in Tint fails to apply the volatile_zero workaround for u32 division and modulo operations when performed on vector types. This allows vector operations to be miscompiled by a known Apple Silicon Metal compiler bug, potentially bypassing WebGPU robustness bounds checks and leading to arbitrary out-of-bounds GPU memory access.

Affected files:

third_party/dawn/src/tint/lang/msl/writer/raise/binary_polyfill.cc

Estimated timestamp from git blame: Unknown (Google3 checkout)

Description

Dawn enables a workaround toggle (MetalFixU32DivMod) on Apple Silicon devices to mitigate a known Metal compiler miscompilation affecting u32 division and modulo operations. This workaround, implemented in the MSL backend’s BinaryPolyfill transform, injects a volatile_zero to prevent the compiler from incorrectly optimizing these operations.

However, the implementation in third_party/dawn/src/tint/lang/msl/writer/raise/binary_polyfill.cc only applies this workaround to scalar u32 types:

} else if ((op == core::BinaryOp::kModulo || op == core::BinaryOp::kDivide) &&
           lhs_type->Is<core::type::U32>()) {
    if (config.fix_u32_div_mod) {
        umod_worklist.Push(binary);
    }
}

The check lhs_type->Is<core::type::U32>() strictly matches scalar types and returns false for core::type::Vector types (e.g., vec3<u32>).

When Tint’s core IR BinaryPolyfill transform processes a vector modulo operation, it moves it into a helper function (e.g., tint_mod_v3u32) that implements modulo manually using a vector division instruction (kDivide): result = lhs - ((lhs / rhs_or_one) * rhs_or_one). Because the MSL-specific transform ignores vectors, the vector division instruction inside this helper function is emitted as raw MSL without the critical volatile_zero protection.

Impact and Exploitation

This oversight potentially allows an attacker to trigger the underlying Metal compiler miscompilation from standard WebGPU content, leading to arbitrary memory access within the GPU process.

Potential Attack Steps: Note: These are suggested steps based on source code analysis; our tooling agent cannot execute code to provide a working PoC.

Craft a Malicious Shader: An attacker writes a WGSL shader that performs a vector modulo operation where the right-hand side is derived from a buffer’s length (e.g., let index = (a % vec3<u32>(arrayLength(&my_buffer))).x;).
Access Memory: The shader uses this index to read or write to my_buffer[index].
BCE Eliminates Clamps: Tint’s core Robustness transform inserts a safety clamp (e.g., min(index, length - 1)). However, when Apple’s Metal compiler processes the generated MSL, its optimizer recognizes the lhs - ((lhs / rhs) * rhs) pattern as a modulo. The Bounds Check Elimination (BCE) pass incorrectly assumes the modulo result is strictly less than the array length, and strips the WebGPU-inserted min() clamp, treating it as redundant.
Miscompilation Triggers OOB: Because the volatile_zero workaround was omitted for the vector type, the Metal backend miscompiles the raw 32-bit vector division. This produces a massive, incorrect integer value. Because the safety clamp was removed by BCE, this massive value is used directly for raw pointer arithmetic, leading to out-of-bounds reads or writes in the shared GPU virtual address space.

This could allow an attacker to compromise the sandboxed GPU process by overwriting internal structures or reading sensitive textures/buffers belonging to other origins.

Recommended Fix

In third_party/dawn/src/tint/lang/msl/writer/raise/binary_polyfill.cc, update the check to match both scalars and vectors:

} else if ((op == core::BinaryOp::kModulo || op == core::BinaryOp::kDivide) &&
           lhs_type->IsUnsignedIntegerScalarOrVector()) {

Update the UMod helper method within the same file to properly construct a volatile_zero that matches the width of the target operand (i.e., splatting the scalar zero into a vector if lhs is a vector type).

Evaluated with Chrome root at commit: 65b3256311f3ab6fb9870eaa522de7e6dd2663bb

Results so far have been promising, but there can be wrong deductions. Feel free to adjust as follows:

If you are familiar with the severity guidelines, you may adjust the severity.
If this is a false positive, and there’s no work to be done, please close as WAI.
If there is work to do here but not a vulnerability, please change the issue type to Task/Bug/FR.

Data from false positives will be used to improve accuracy over time. And please feel free to reach out to me directly if you have concerns or feedback on the project.

View on issue tracker

References

Vendor advisory Fix commit Bug tracker NVD

On This Page