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: Dawn’s Metal backend uses divergent logic when calculating and validating workgroup memory sizes for WGSL shaders. By heavily padding a struct via the @align attribute, an attacker can pass Dawn’s size validation checks but force the backend to allocate an out-of-bounds threadgroup memory length. This results in a potential Metal API precondition violation that triggers undefined driver behavior in the sandboxed GPU process.

Affected files:

• third_party/dawn/src/tint/lang/msl/writer/printer/printer.cc
• third_party/dawn/src/dawn/native/metal/ComputePipelineMTL.mm
• third_party/dawn/src/tint/lang/msl/writer/raise/module_scope_vars.cc
• third_party/dawn/src/tint/lang/core/type/manager.cc
• third_party/dawn/src/tint/lang/wgsl/resolver/resolver.cc
• third_party/dawn/src/dawn/native/ShaderModule.cpp

Estimated timestamp from git blame: 2024-12-11

1. Summary of the Issue (Meant for Human Triage)

In Dawn’s Metal backend, a discrepancy in how workgroup memory sizes are calculated allows attacker-controlled WGSL shaders to potentially bypass front-end validation limits. Tint’s MSL backend computes two divergent values for workgroup storage:

1. workgroup_info.storage_size: Computed by iterating over individual variables and summing their sizes. This smaller value is validated against the WebGPU device’s limit (maxComputeWorkgroupStorageSize).
2. workgroup_allocations[0]: Represents the actual size of a generated wrapper struct that aggregates all var<workgroup> variables. This larger value is passed directly to the Metal driver via [MTLComputeCommandEncoder setThreadgroupMemoryLength:atIndex:].

Because WGSL allows unbounded @align(N) attributes, Tint’s struct layout engine introduces substantial inter-member padding when packing variables into the wrapper struct. This padding is included in the final struct Size(), but is completely ignored by the validated storage_size. An attacker can supply a shader that passes the front-end validation check, but forces the GPU process to call setThreadgroupMemoryLength with a value exceeding the hardware driver limit (MTLDevice.maxThreadgroupMemoryLength).

This violates Apple’s API preconditions and causes undefined driver behavior or validation layer crashes. However, because Tint omits explicit MSL padding for kWorkgroup structs, the emitted MSL shader operates on a dense layout and does not perform out-of-bounds memory accesses. Consequently, this is a spec-invalid driver call without memory corruption, meriting a Medium (S2) severity in the sandboxed GPU process.

2. Proof-of-Concept & Detailed Execution Flow

Note: These are potential execution steps as our tooling agent has verified this statically but cannot dynamically execute a payload.

Step 1: Attacker WGSL Payload Construction An attacker constructs a WGSL compute shader containing a custom struct with a massive @align(N) attribute, and interleaves standard variables with this struct:

struct S { @align(16384) x : f32 };
var<workgroup> a : f32;
var<workgroup> b : S;
var<workgroup> c : f32;
@compute @workgroup_size(1) fn main() { _ = a + b.x + c; }

Step 2: WGSL Resolution and Unbounded @align Dawn delegates WGSL compilation to Tint. In third_party/dawn/src/tint/lang/wgsl/resolver/resolver.cc:4300-4329, Tint parses the @align(16384) attribute. Tint enforces that the alignment is a power of two (tint::IsPowerOfTwo), but applies no explicit upper-bound cap. The value 16384 is permitted.

Step 3: Wrapper Struct Layout Calculation To comply with MSL’s limits on threadgroup buffers, Tint’s ModuleScopeVars transform bundles the workgroup variables (a, b, c) into a single wrapper struct. In third_party/dawn/src/tint/lang/core/type/manager.cc:414-433, the layout manager computes the memory layout. When it processes b (alignment 16384), it executes offset = tint::RoundUp(align, current_size). This injects a massive 16380-byte padding gap between a and b. The final wrapper struct size (ty->Size()) is computed as 49152.

Step 4: Divergent Storage Size Computation During MSL generation (third_party/dawn/src/tint/lang/msl/writer/printer/printer.cc:453-481), Tint populates the driver allocation size:

auto& allocations = result_.workgroup_allocations;
allocations.push_back(ty->Size()); // Stores 49152

However, it independently calculates workgroup_info.storage_size using a divergent formula that iterates through the original individual variables, completely ignoring the layout manager’s inter-member offset padding:

for (auto& mem : ty->As<core::type::Struct>()->Members()) {
    auto mem_ty = mem->Type();
    uint64_t align = mem_ty->Align();
    uint64_t size = mem_ty->Size();
    result_.workgroup_info.storage_size +=
        tint::RoundUp(static_cast<uint64_t>(16u), tint::RoundUp(align, size));
}

This strictly calculates 16 + 16384 + 16 = 16416 bytes.

Step 5: Validation Bypass Control returns to Dawn’s Metal ShaderModule (third_party/dawn/src/dawn/native/ShaderModule.cpp:1405-1412). Dawn compares workgroupInfo.storage_size (16416) against limits.maxComputeWorkgroupStorageSize (e.g., 32768 on an Apple4 GPU). Since 16416 is less than 32768, the DAWN_INVALID_IF check passes successfully.

Step 6: Driver API Contract Violation (The Sink) When the attacker invokes pass.dispatchWorkgroups(), the GPU process executes ComputePipeline::Encode (third_party/dawn/src/dawn/native/metal/ComputePipelineMTL.mm:97-107). The backend loops over mWorkgroupAllocations and directly calls:

[encoder setThreadgroupMemoryLength:rounded atIndex:i]; // length = 49152

Apple’s documentation strictly defines that this length must be less than or equal to MTLDevice.maxThreadgroupMemoryLength (32768). Passing 49152 violates the API precondition, triggering an assertion in the Metal Validation Layer or undefined behavior in the release OS driver.

Mitigating Factor (No Shader OOB Execution) In third_party/dawn/src/tint/lang/msl/writer/printer/printer.cc:1648-1815, EmitStructType omits generating explicit tint_pad padding bytes for structs unless they are in host_shareable_structs_. FindHostShareableStructs explicitly skips kWorkgroup address spaces. As a result, the generated MSL uses a tightly packed natural C++ layout. The compiled shader only accesses the first ~12 bytes of the threadgroup memory. Even if the driver under-allocates memory due to the precondition violation, no out-of-bounds reads or writes occur from the shader execution side.

Suggested Fix The storage_size calculation in third_party/dawn/src/tint/lang/msl/writer/printer/printer.cc should be updated to match the final bundled wrapper struct’s size (ty->Size()). Alternatively, Dawn’s validation in ShaderModule.cpp should validate the values within workgroup_allocations against maxComputeWorkgroupStorageSize before passing them to the Metal driver.

3. Technical Verification Details (Automated Audit Logs)

> The vulnerability report is 100% accurate and correctly identifies a validation gap in Dawn’s Metal backend. Tint computes workgroup_info.storage_size using a per-variable sum (validated against maxComputeWorkgroupStorageSize in ShaderModule.cpp), but the Metal backend uses workgroup_allocations, populated with the wrapper struct’s Size(). Because WGSL allows unbounded @align(N) attributes, ty->Size() can be arbitrarily inflated by inter-member padding. This allows an attacker to bypass Dawn’s validation and pass an oversized length to [MTLComputeCommandEncoder setThreadgroupMemoryLength:atIndex:] in ComputePipelineMTL.mm. > Per the KB ‘GPU/media driver-sink triage’ guide, length is a bounds-bearing parameter. However, because Tint omits explicit padding for non-host-shareable structs (like kWorkgroup), the emitted MSL shader uses a natural layout and does not perform out-of-bounds memory accesses. Thus, the issue is a spec-invalid driver call without memory corruption. According to the severity guidelines, a GPU validating-layer gap where a spec-invalid value reaches the driver but no OOB write is demonstrated (or where drivers cleanly tolerate it) is rated Medium (S2). Furthermore, the Metal backend runs in the GPU process, which is sandboxed on macOS/iOS, eliminating any path to S0.

Codebase Investigation Traces:

• Divergent calculations: Verified in third_party/dawn/src/tint/lang/msl/writer/printer/printer.cc:460-489. ty->Size() is appended to allocations (line 466), while result_.workgroup_info.storage_size rounds and sums individual members (line 486).
• @align bounds missing: Verified in third_party/dawn/src/tint/lang/wgsl/resolver/resolver.cc:4300-4329. The validation uses tint::IsPowerOfTwo(value), but lacks an explicit upper bound (the only limits are the maximum possible u32 value and the 0xffffffff size limits on struct offset/size lines 4413/4433).
• Device Limits: third_party/dawn/src/dawn/native/metal/PhysicalDeviceMTL.mm:849 sets maxTotalThreadgroupMemory up to 32768u on modern Apple hardware. Line 947 initializes limits->v1.maxComputeWorkgroupStorageSize to this same value.
• Padding Mitigation Proof: Investigated FindHostShareableStructs (printer.cc:235-281). Parameters are filtered by core::IsHostShareable(ptr->AddressSpace()). According to third_party/dawn/src/tint/lang/core/enums.h:127-131, kWorkgroup is explicitly excluded from returning true. Thus, EmitStructType does not generate tint_pad bytes for the wrapper struct, confirming natural C++ layout and preventing out-of-bounds shader execution.

Evaluated with Chrome root at commit: b96d2ec58f4f5f92b540a723966b199d6e9951b4

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.