Report description

WebGPU (Dawn/Tint Metal) SubstituteOverrides integer overflow causes threadgroup OOB read/write

Bug location

Where do you want to report your vulnerability?

Chrome VRP – Report security issues affecting the Chrome browser. See program rules

Which URL (or repository) have you found the vulnerability in?

https://dawn.googlesource.com/dawn/+/main/src/tint/lang/core/ir/transform/substitute_overrides.cc

The problem

Please describe the technical details of the vulnerability

• Chrome Stable 144.0.7559.133 (Official Build) (arm64), macOS Tahoe 26.2, Apple M1 (GPU family Metal-3)

A web-reachable out-of-bounds read/write in Metal threadgroup memory, caused by a uint32_t overflow in Tint’s SubstituteOverrides transform. The overflow creates a size/count mismatch: Dawn allocates only 32 bytes of threadgroup memory while the emitted MSL declares a threadgroup struct whose first member is an array describing ~4,294,967,300 bytes. In testing on Apple Silicon, pipeline creation succeeds and dispatch completes with the undersized allocation. Writes through the oversized array deterministically corrupt values read through other workgroup variables, providing a controlled read/write primitive within threadgroup storage.

In SubstituteOverrides::Run (substitute_overrides.cc, near line 215):

auto* new_ty = ty.Get<core::type::Array>(old_ty->ElemType(), new_cnt,
                                         num_elements * old_ty->ImplicitStride());

num_elements is uint32_t (line 213) and ImplicitStride() returns uint32_t (array.h:84). The multiplication is computed in 32-bit arithmetic and passed to the Array constructor as a 32-bit byte size.

When num_elements = 1073741825 and ImplicitStride() = 4:

• 1073741825 * 4 = 0x1_0000_0004 → truncated to uint32_t = 4
• Array::size_ (uint32_t, array.h:111) stores 4 instead of 4,294,967,300
• Array::count_ (array.h:110) points to ConstantArrayCount(1,073,741,825) (element count correct)

This allows Array::count_ and Array::size_ to diverge. As a result, allocation/validation paths that consult Size() undercount workgroup storage, while MSL codegen still emits the full declared extent from count_. The mismatch propagates through the MSL writer (printer.cc):

1. Undersized allocation - Workgroup variable emission (near line 433): ty->Size() pushes 4 → setThreadgroupMemoryLength:32 atIndex:0 at pipeline creation (ComputePipelineMTL.mm, near line 105)
2. Validation bypass - Workgroup storage calculation (near line 452): mem_ty->Size() returns 4 → workgroup_info.storage_size = 32 → passed via ShaderModuleMTL.mm (near line 393) to ValidateComputeStageWorkgroupSize (ShaderModule.cpp) → passes (32 < 16384)
3. MSL codegen - Uses count_ (correct 1,073,741,825) for the array declaration → emitted MSL type describes ~4,294,967,300 bytes

The emitted MSL and host-side allocation size mismatch creates out-of-bounds threadgroup accesses. The PoC demonstrates deterministic corruption consistent with overlap between packed workgroup variables.

Steps to Reproduce

1. Open tint_poc.html in Chrome Stable on an Apple Silicon Mac
2. Open DevTools Console and observe the output (shown below)
3. To see generated MSL: relaunch Chrome with --enable-dawn-features=dump_shaders and repeat step 1
4. Additional PoCs (each self-contained, no special flags):
   • tint_types_poc.html - type diversity and offset map
   • tint_controlflow_poc.html - control-flow steering
   • tint_disclosure_poc.html - bidirectional read/write

Reproduces on every run tested across page refreshes and new-tab lifecycles on Apple Silicon M1. Results do not depend on --enable-dawn-features=dump_shaders or any runtime flags.

PoC output (Chrome Stable 144, Apple Silicon M1)

The PoC runs the same WGSL shader twice - once with N=16 (CONTROL), once with N=1073741825 (TEST). Both write to overflow[1..4] and then read victim[0..3]:

CONTROL (N=16):         victim[] = 0xaaaa0000 0xaaaa0001 0xaaaa0002 0xaaaa0003
TEST    (N=1073741825): victim[] = 0x41414141 0x42424242 0x43434343 0x44444444

CONTROL: victim[] retains its init values - overflow[] and victim[] are separate (correct). TEST: victim[] contains the values written to overflow[], indicating that writes through the oversized array overlap the storage used for the second workgroup variable.

WGSL input (identical for both runs)

override N: u32;
var<workgroup> overflow: array<u32, N>;
var<workgroup> victim: array<u32, 4>;
@group(0) @binding(0) var<storage, read_write> out: array<u32>;
@compute @workgroup_size(1)
fn main() {
    victim[0] = 0xAAAA0000u;  victim[1] = 0xAAAA0001u;
    victim[2] = 0xAAAA0002u;  victim[3] = 0xAAAA0003u;
    workgroupBarrier();
    out[0] = victim[0]; out[1] = victim[1]; out[2] = victim[2]; out[3] = victim[3];
    overflow[1] = 0x41414141u;  overflow[2] = 0x42424242u;
    overflow[3] = 0x43434343u;  overflow[4] = 0x44444444u;
    workgroupBarrier();
    out[4] = victim[0]; out[5] = victim[1]; out[6] = victim[2]; out[7] = victim[3];
}

Generated MSL (via --enable-dawn-features=dump_shaders)

The MSL backend packs the shader’s var<workgroup> variables into a single threadgroup struct, passed to the kernel as [[threadgroup(0)]]. Dawn sets the allocation size for this index via setThreadgroupMemoryLength:atIndex:0. The only difference between CONTROL and TEST is the array size in this struct:

CONTROL (N=16):

struct tint_struct_3 {
  tint_array<uint, 16> tint_member_6;          // overflow[] - 64 bytes
  tint_array<uint, 4> tint_member_7;            // victim[]   - 16 bytes
};
// kernel: void v_1(threadgroup tint_struct_3* v_18 [[threadgroup(0)]], ...)
// Dawn sets: setThreadgroupMemoryLength:80 atIndex:0. No overlap.

TEST (N=1073741825):

struct tint_struct_3 {
  tint_array<uint, 1073741825> tint_member_6;  // overflow[] - type describes ~4,294,967,300 bytes
  tint_array<uint, 4> tint_member_7;            // victim[]   - 16 bytes
};
// kernel: void v_1(threadgroup tint_struct_3* v_18 [[threadgroup(0)]], ...)
// Dawn sets: setThreadgroupMemoryLength:32 atIndex:0 - for a struct whose first member describes ~4,294,967,300 bytes.

Given member order and standard struct layout, tint_member_7 (victim) follows tint_member_6 (overflow). Because the allocation for tint_member_6 is based on the overflowed Size() of 4, stores to overflow[1..4] address the same bytes later read as victim[0..3]. The types PoC (tint_types_poc.html) confirms predictable alignment gaps - indices 2-3 are padding before vec4’s 16-byte boundary - reinforcing that the overlap follows standard packing rules rather than incidental behavior.

The TEST kernel also emits a zero-initialization loop iterating 1,073,741,825 times (CONTROL: 16):

if ((v_6 >= 1073741825u)) { break; }
(*v_2.tint_member)[v_6] = 0u;

The injected values are emitted as MSL literals in the generated shader. Because the indices (1-4) are in-range for the declared array extent of 1,073,741,825, index-based robustness checks (where present) do not prevent the access:

(*v_2.tint_member)[1u] = 1094795585u;   // 0x41414141
(*v_2.tint_member)[2u] = 1111638594u;   // 0x42424242
(*v_2.tint_member)[3u] = 1128481603u;   // 0x43434343
(*v_2.tint_member)[4u] = 1145324612u;   // 0x44444444

Capabilities demonstrated

The corruption primitive was characterized through three additional PoCs, each isolating a specific property. All results observed in testing on Apple Silicon M1, Chrome Stable 144.0.7559.133.

Type diversity (tint_types_poc.html): Each overflow[i] tagged 0xEE0000XX to discover the offset map. All 4 types corrupted (CONTROL: all intact):

  scalar (u32):    0xAAAA0001 → 0xEE000001  (overflow[1])
  vec4<u32>:       0xBBBB{01–04} → 0xEE0000{04–07}  (overflow[4–7])
  mat2x2<f32>:     0xCCCC{01–04} → 0xEE0000{08–0B}  (overflow[8–11])
  atomic<u32>:     0xDDDD0001 → 0xEE00000C  (overflow[12])

Gaps at indices 2-3 (padding before vec4’s 16-byte alignment) confirm the map is predictable from declaration order and type alignment rules.

Control-flow steering (tint_controlflow_poc.html): Three outcomes indicating corruption of values not written by the program after initialization/barrier (CONTROL: all pass correctly):

  Branch: guard[0] set to 0, branch requires !=0  →  TAKEN (0xFFFF0001)
  Loop:   guard[1] set to 3, loop bounded by guard[1]  →  ran 50 times
  Barrier: guard[3] set to 0xAAAAAAAA, never rewritten  →  reads 0x0000DEAD

The attacker controls the injected values (50 and 0xDEAD) via overflow[] writes.

Disclosure (tint_disclosure_poc.html): Secrets computed at runtime, then recovered via overflow[] reads (CONTROL: overflow reads return zero, secrets intact):

  Disclosure: overflow[1..4] = 0x53450301..04  (= secret[0..3] values)
  Injection:  secret[0..3]  = 0xBEEF0001..04  (= overflow write values)

4/4 secrets disclosed; 4/4 injections confirmed. The primitive is a full read/write within threadgroup storage.

Determinism and scope: This yields a deterministic, parameterizable out-of-bounds threadgroup read/write primitive (corruption + disclosure) within a single dispatch. No observable effects outside threadgroup/tile SRAM in targeted probes.

MTL_SHADER_VALIDATION - diagnostic confirmation

With MTL_SHADER_VALIDATION=1, Apple’s Metal Shader Validation layer flags the threadgroup out-of-bounds access during execution; on our test machine this triggers a GPU fault that causes WindowServer to terminate and logs the user out. This is included as diagnostic confirmation of the OOB; the validation layer is not enabled in normal Chrome configurations. Production configurations may instead drop the dispatch or lose the GPU device.

Affected Code

Overflow:

1. substitute_overrides.cc:~215 - num_elements * ImplicitStride() overflows: 1073741825 * 4 = 4 (uint32_t)
2. array.h:111 - Array::size_ stores 4; Array::count_ stores 1,073,741,825 (correct)

Allocation path (uses size_ → wrong): 3. printer.cc:~595 - mem_ty->Size() returns 4 → workgroup_info.storage_size = 32 4. ShaderModule.cpp - ValidateComputeStageWorkgroupSize: 32 < 16384 → passes 5. ComputePipelineMTL.mm:~105 - setThreadgroupMemoryLength:32 atIndex:0

Codegen path (uses count_ → correct): 6. printer.cc:~1060 - arr->ConstantCount() returns 1,073,741,825 → emits tint_array<uint, 1073741825>

Result: 32 bytes allocated for a struct whose first member describes ~4,294,967,300 bytes.

Fix

In Tint (SubstituteOverrides::Run): check for overflow before constructing the Array type. The error must cause shader compilation / pipeline creation to fail (not continue with a partially-transformed module):

uint32_t num_elements = new_value->Value()->ValueAs<uint32_t>();
uint64_t array_size = static_cast<uint64_t>(num_elements) * old_ty->ImplicitStride();
if (array_size > std::numeric_limits<uint32_t>::max()) {
    b.Diagnostics().AddError(source) << "override-sized array exceeds maximum byte size";
    return;
}

As defense-in-depth, widening Array::size_ to uint64_t would cause Dawn’s existing DAWN_INVALID_IF(workgroupStorageSize > limits.v.maxComputeWorkgroupStorageSize) check in ValidateComputeStageWorkgroupSize to correctly reject the pipeline, since it would see the true size (~4,294,967,300) instead of the truncated value (4).

Bisect

• Last good: b35ac660e5c9c8510646ad292e030c13883dab9a - “[ir] Implement Substitute Overrides” (Nov 12 2024). Creates the transform but explicitly defers array support: “the array type override information still needs to be substituted.”
• First bad: 7a4fb42c574409f70d3e5b29ccc8b8992674dcd4 - “[ir] Handle array types in substitute overrides” (Nov 13 2024). Introduces uint32_t num_elements * old_ty->Stride() with no overflow check.
• Later refactor: afef3f952f7cfd254891c3e22b202fed7082793f - “Cleanup stride in arrays” (Aug 26 2025). Changes Stride() to ImplicitStride() but does not add an overflow check.

Impact analysis

• Demonstrated OOB read/write within threadgroup storage (corruption + disclosure). See capabilities table below.
• Demonstrated control-flow influence: branch condition, loop bound, and barrier invariant corrupted by injected values. See control-flow steering PoC below.
• No device-memory effects observed in targeted probes.

The cause

What version of Chrome have you found the security issue in?

144.0.7559.133 Stable

Is the security issue related to a crash?

Yes, it is related to a crash.

Choose the type of vulnerability

Memory Corruption (in a sandboxed process)

How would you like to be publicly acknowledged for your report?

cinzinga