PatchGap
Chrome · Dawn › GL

CVE-2026-14420

OOB in Dawn (GL)
Critical watch CVSS 9.6 chrome OOB 📄 Reporter bug report 🔧 Commit mapped 🧬 Front-end validation exempts a record the backend still uses

Overview

Critical
Severity
9.6
CVSS
No
Exploited ITW
Fixed
Fix Status
ImpactOut of bounds read and write in Dawn
DescriptionOut of bounds read and write in Dawn
ComponentDawn
Sub-componentGL
Bug ClassOOB
Action tierwatch — corruption primitive but gated (uncommon) — not default-reachable
Computed by PatchGap's deterministic exploitability ruler from the fixing commit. patch-now = RCE-class & attacker-reachable in a default config, or exploited in the wild; watch = real security signal below that bar; noise = no security relevance.
Tracker517031505
Fix commit73591422f504 (dawn) +94/-1
CWECWE-125, CWE-787 (Out-of-bounds read, Out-of-bounds write)
CVSS vectorCVSS:3.1/AV:N/AC:L/PR:N/UI:R/S:C/C:H/I:H/A:H
CISA KEVNot listed
CreditedGoogle
Reported2026-05-27 (per vendor advisory)
Fix landed2026-06-02 (public git)
Disclosed2026-06-30
Patch gap28 days — fix public in source before the advisory

Background

`ArrayLengthFromUniform`
a Dawn/Tint GL-backend mechanism that supplies arrayLength() for runtime-sized storage buffers by loading each SSBO’s bound byte-size from a UBO slot (enabled by Toggle::GLUseArrayLengthFromUniform / gl_use_array_length_from_uniform) instead of a native GLSL length query.
`bindpoint_to_size_index`
the Tint option map keyed by a {group, binding} bindpoint that names the UBO slot holding that buffer’s size; if two bindpoints resolve to the same key, one entry silently overwrites the other and arrayLength() returns the wrong buffer’s size.
`SingleShaderStage` / `visibility`
each bind group layout entry carries a visibility mask (Vertex/Fragment/Compute); only stages present in the mask should contribute to that binding’s codegen for the stage currently being compiled.
Robustness clamping
Dawn’s Robustness pass rewrites a storage-buffer access B[idx] as B[min(idx, arrayLength(&B)-1)], so the clamp is only as tight as the arrayLength() value it trusts.

Root Cause Analysis

GenerateArrayLengthFromuniformData() in ShaderModuleGL.cpp walked every buffer binding in the PipelineLayout and, for each binding whose type was Storage or kInternalStorageBufferBinding, assigned it an ArrayLengthFromUniform size-index entry — with no check that the binding was visible to the stage being compiled. A storage buffer visible only to another stage (e.g. a Fragment-only buffer while the Compute stage is compiled) therefore still received an entry, and it kept its original pre-remap {group, binding} key while a genuinely stage-visible buffer received its post-remap key. The regression test makes these coincide: buffer A (Fragment-only, group 0, @binding 1) keeps its pre-remap key {0, 1}, and buffer B (Compute, group 1, remapped to glIndex_B == 1) gets post-remap key {0, 1}, so A’s stale entry aliases B’s slot in bindpoint_to_size_index. The visible buffer B then loaded A’s much larger bound size as its own arrayLength(), breaking the invariant that a buffer’s reported length matches its own bound sub-range. Because Robustness clamps indices against that inflated length, a clamped access on B could reach far past its 64-byte bound range, yielding OOB read and write in the GPU process.

The fix threads SingleShaderStage stage into the function and continues past any binding whose visibility mask lacks StageBit(stage), so a stage-hidden buffer no longer contributes an entry that can collide.

Key insight
The collision runs the direction the test spells out: the visible buffer’s key is the one that gets remapped (B’s post-remap {0, glIndex_B}), and it lands on the hidden buffer’s stale pre-remap key (A’s {0, 1}) — so the hidden buffer’s unfiltered entry hijacks the visible buffer’s arrayLength(), not the reverse.

Attack Path

  1. Craft the aliasing layout Build a pipeline layout with a large storage buffer A visible only to Fragment (group 0, @binding 1) and a storage buffer B visible to Compute (group 1, @binding 0), so A’s pre-remap key {0, 1} equals B’s post-remap key {0, glIndex_B = 1}.
  2. Bind the victim to a narrow sub-range Bind B to only a 64-byte sub-range of a 16 KiB allocation (true arrayLength() = 16) while binding A at its full 1 MiB size, so the two reported lengths differ sharply.
  3. Trigger the collision at compile time Compile the compute pipeline; GenerateArrayLengthFromuniformData() emits a size-index entry for the Fragment-only buffer A, and after remapping A’s key overwrites B’s entry in bindpoint_to_size_index.
  4. Inflate the clamp bound At dispatch, arrayLength(&B) returns A’s large bound size instead of 16, so Robustness clamping no longer constrains B’s index to its 64-byte range.
  5. Perform the OOB access Write B[1000] (byte 4000); the clamp passes it through, landing past B’s bound sub-range within the larger allocation or adjacent memory — OOB read/write in the GPU process. Chaining this into information disclosure or a sandbox escape is not shown by the patch (inferred).

Impact Assessment

On the OpenGL/OpenGL ES backend with Toggle::GLUseArrayLengthFromUniform active, WebGPU content that arranges the described key aliasing obtains out-of-bounds read and write in the GPU-process address space, bounded by the inflated arrayLength() and surrounding allocations. The advisory classifies this as a critical out-of-bounds read and write in Dawn; it is not recorded as exploited in the wild (ITW: False) or in CISA KEV (KEV: False). Whether the GPU-process corruption can be escalated to a sandbox escape or broader compromise is not demonstrated by the patch (inferred).

Changed Functions

FunctionChangeNotes
if
src/dawn/native/opengl/ShaderModuleGL.cpp
 modified
OpArrayLengthVisibilityCollisionTest
src/dawn/tests/end2end/OpArrayLengthTests.cpp
 modified
TEST_P
src/dawn/tests/end2end/OpArrayLengthTests.cpp
 modified

Files Changed

  • src/dawn/native/opengl/ShaderModuleGL.cpp
  • src/dawn/tests/end2end/OpArrayLengthTests.cpp

Audit Directions

  • Stage-visibility filtering in per-stage codegen
    Audit every place Dawn iterates all PipelineLayout bindings to emit stage-specific data (size indices, remapper tables) and confirm each filters by bindingInfo.visibility & StageBit(stage) before emitting anything for that stage.
  • Post-remap key aliasing
    Review Tint binding-remapper consumers such as bindpoint_to_size_index for keys that can collide once unrelated bindings (kept at pre-remap keys) coexist with remapped ones, so that one entry cannot overwrite another.
  • `arrayLength()` provenance vs. Robustness
    Trace all sources feeding arrayLength() / bindpoint_to_size_index across backends and verify a buffer can never be handed another buffer’s bound size, since Robustness clamping trusts that value as the sole bound.

Patch

view commit
From 73591422f504902c597a67beaeb8e22101f030bd Mon Sep 17 00:00:00 2001
From: Stephen White <senorblanco@chromium.org>
Date: Tue, 02 Jun 2026 13:24:30 -0700
Subject: [PATCH] GL: add missing stage visibility check

In GenerateArrayLengthFromUniformData(), skip bindings that are
not visible to the current stage.

Bug: 517031505

Change-Id: Ida25ca9233a6880461f67a177e9d2257c1727094
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/313155
Reviewed-by: Geoff Lang <geofflang@chromium.org>
Commit-Queue: Stephen White <senorblanco@chromium.org>
---

diff --git a/src/dawn/native/opengl/ShaderModuleGL.cpp b/src/dawn/native/opengl/ShaderModuleGL.cpp
index 1a8687b..f5e304a 100644
--- a/src/dawn/native/opengl/ShaderModuleGL.cpp
+++ b/src/dawn/native/opengl/ShaderModuleGL.cpp
@@ -273,6 +273,7 @@
 bool GenerateArrayLengthFromuniformData(
     const BindingInfoArray& moduleBindingInfo,
     const PipelineLayout* layout,
+    SingleShaderStage stage,
     tint::glsl::writer::ArrayLengthFromUniformOptions& options) {
     const PipelineLayout::BindingIndexInfo& indexInfo = layout->GetBindingIndexInfo();
 
@@ -282,6 +283,11 @@
         for (BindingIndex binding : bgl->GetBufferIndices()) {
             const BindingInfo& bindingInfo = bgl->GetBindingInfo(binding);
 
+            // Skip bindings that aren't visible to this stage.
+            if (!(bindingInfo.visibility & StageBit(stage))) {
+                continue;
+            }
+
             switch (std::get<BufferBindingInfo>(bindingInfo.bindingLayout).type) {
                 case wgpu::BufferBindingType::Storage:
                 case kInternalStorageBufferBinding:
@@ -409,7 +415,7 @@
 
     if (GetDevice()->IsToggleEnabled(Toggle::GLUseArrayLengthFromUniform)) {
         *needsSSBOLengthUniformBuffer = GenerateArrayLengthFromuniformData(
-            moduleBindingInfo, layout, req.tintOptions.array_length_from_uniform);
+            moduleBindingInfo, layout, stage, req.tintOptions.array_length_from_uniform);
         if (*needsSSBOLengthUniformBuffer) {
             req.tintOptions.use_array_length_from_uniform = true;
             req.tintOptions.array_length_from_uniform.ubo_binding = {
diff --git a/src/dawn/tests/end2end/OpArrayLengthTests.cpp b/src/dawn/tests/end2end/OpArrayLengthTests.cpp
index ea36ca3..a96a58d 100644
--- a/src/dawn/tests/end2end/OpArrayLengthTests.cpp
+++ b/src/dawn/tests/end2end/OpArrayLengthTests.cpp
@@ -303,6 +303,93 @@
                       VulkanBackend(),
                       WebGPUBackend());
 
+// Regression test for stage-visibility filtering in
+// GenerateArrayLengthFromuniformData (ShaderModuleGL.cpp). A storage
+// buffer in the layout that is *not* visible to the stage being
+// compiled should not be given an entry in Tint's
+// bindpoint_to_size_index nor in the remapper_data. Otherwise,
+// ArrayLengthFromUniform loads the wrong UBO slot for arrayLength(),
+// defeating Robustness clamping.
+class OpArrayLengthVisibilityCollisionTest : public DawnTest {
+  protected:
+    void GetRequiredLimits(const dawn::utils::ComboLimits& supported,
+                           dawn::utils::ComboLimits& required) override {
+        supported.UnlinkedCopyTo(&required);
+    }
+};
+
+TEST_P(OpArrayLengthVisibilityCollisionTest, ComputeWithFragmentOnlyBufferInLayout) {
+    DAWN_TEST_UNSUPPORTED_IF(GetSupportedLimits().maxStorageBuffersInFragmentStage < 1);
+
+    // Buffer A: large, bound at full size to a FRAGMENT-only storage slot.
+    constexpr uint32_t kBufASize = 1u << 20;  // 1 MiB → arrayLength<u32> = 262144
+    wgpu::BufferDescriptor descA;
+    descA.size = kBufASize;
+    descA.usage = wgpu::BufferUsage::Storage;
+    wgpu::Buffer bufA = device.CreateBuffer(&descA);
+
+    // Buffer B: large allocation, but only a 64-byte sub-range will be bound to
+    // the COMPUTE-visible storage slot. With the collision the shader receives
+    // A's bound size as B's arrayLength, so the Robustness clamp on B[idx]
+    // permits writes far past the 64-byte bound range.
+    constexpr uint32_t kBufBSize = 16384;  // 16 KiB underlying allocation
+    constexpr uint32_t kBufBBound = 64;    // 64 bytes bound → arrayLength = 16
+    constexpr uint32_t kOobIndex = 1000;   // Target B[1000] (byte 4000) – OOB
+    constexpr uint32_t kOobMarker = 0xDEAD4A11u;
+    std::vector<uint32_t> initialData(kBufBSize / 4, 0u);
+    wgpu::Buffer bufB = utils::CreateBufferFromData(
+        device, initialData.data(), initialData.size() * sizeof(uint32_t),
+        wgpu::BufferUsage::Storage | wgpu::BufferUsage::CopySrc | wgpu::BufferUsage::CopyDst);
+
+    // PipelineLayoutGL assigns ssboIndex by iterating groups in order, so:
+    //   group 0 (A, FRAGMENT-only) → ssboIndex 0
+    //   group 1 (B, COMPUTE)       → ssboIndex 1  → glIndex_B = 1
+    // Choose A's WGSL @binding == glIndex_B (= 1) so that A's pre-remap key
+    // {0,1} equals B's post-remap key {0, glIndex_B}.
+    wgpu::BindGroupLayout bglA = utils::MakeBindGroupLayout(
+        device, {{1, wgpu::ShaderStage::Fragment, wgpu::BufferBindingType::ReadOnlyStorage}});
+    wgpu::BindGroupLayout bglB = utils::MakeBindGroupLayout(
+        device, {{0, wgpu::ShaderStage::Compute, wgpu::BufferBindingType::Storage}});
+
+    std::string shaderSource = R"(
+        @group(1) @binding(0) var<storage, read_write> B : array<u32>;
+        @compute @workgroup_size(1) fn main() {
+            B[0] = arrayLength(&B);
+            // Robustness wraps this as B[min(idx, arrayLength(&B)-1)].
+            // With the bug arrayLength(&B) is huge, so the write lands at
+            // index 1000 – past the 64-byte bound range.
+            B[)" + std::to_string(kOobIndex) +
+                               R"(u] = )" + std::to_string(kOobMarker) + R"(u;
+        })";
+
+    wgpu::ComputePipelineDescriptor pipelineDesc;
+    pipelineDesc.layout = utils::MakePipelineLayout(device, {bglA, bglB});
+    pipelineDesc.compute.module = utils::CreateShaderModule(device, shaderSource);
+    pipelineDesc.compute.entryPoint = "main";
+    wgpu::ComputePipeline pipeline = device.CreateComputePipeline(&pipelineDesc);
+
+    wgpu::BindGroup bgA = utils::MakeBindGroup(device, bglA, {{1, bufA, 0, kBufASize}});
+    wgpu::BindGroup bgB = utils::MakeBindGroup(device, bglB, {{0, bufB, 0, kBufBBound}});
+
+    wgpu::CommandEncoder encoder = device.CreateCommandEncoder();
+    wgpu::ComputePassEncoder pass = encoder.BeginComputePass();
+    pass.SetPipeline(pipeline);
+    pass.SetBindGroup(0, bgA);
+    pass.SetBindGroup(1, bgB);
+    pass.DispatchWorkgroups(1);
+    pass.End();
+    wgpu::CommandBuffer commands = encoder.Finish();
+    queue.Submit(1, &commands);
+
+    EXPECT_BUFFER_U32_EQ(kBufBBound / 4, bufB, 0);
+    EXPECT_BUFFER_U32_EQ(0u, bufB, kOobIndex * sizeof(uint32_t));
+}
+
+DAWN_INSTANTIATE_TEST(OpArrayLengthVisibilityCollisionTest,
+                      OpenGLESBackend(),
+                      OpenGLESBackend({"gl_use_array_length_from_uniform"}),
+                      VulkanBackend());
+
 enum class TieredLimits {
     No,
     Yes,
Loading diff…

Regression Test / PoC

shipped with the fix
diff --git a/src/dawn/tests/end2end/OpArrayLengthTests.cpp b/src/dawn/tests/end2end/OpArrayLengthTests.cpp
index ea36ca3..a96a58d 100644
--- a/src/dawn/tests/end2end/OpArrayLengthTests.cpp
+++ b/src/dawn/tests/end2end/OpArrayLengthTests.cpp
@@ -303,6 +303,93 @@
                       VulkanBackend(),
                       WebGPUBackend());
 
+// Regression test for stage-visibility filtering in
+// GenerateArrayLengthFromuniformData (ShaderModuleGL.cpp). A storage
+// buffer in the layout that is *not* visible to the stage being
+// compiled should not be given an entry in Tint's
+// bindpoint_to_size_index nor in the remapper_data. Otherwise,
+// ArrayLengthFromUniform loads the wrong UBO slot for arrayLength(),
+// defeating Robustness clamping.
+class OpArrayLengthVisibilityCollisionTest : public DawnTest {
+  protected:
+    void GetRequiredLimits(const dawn::utils::ComboLimits& supported,
+                           dawn::utils::ComboLimits& required) override {
+        supported.UnlinkedCopyTo(&required);
+    }
+};
+
+TEST_P(OpArrayLengthVisibilityCollisionTest, ComputeWithFragmentOnlyBufferInLayout) {
+    DAWN_TEST_UNSUPPORTED_IF(GetSupportedLimits().maxStorageBuffersInFragmentStage < 1);
+
+    // Buffer A: large, bound at full size to a FRAGMENT-only storage slot.
+    constexpr uint32_t kBufASize = 1u << 20;  // 1 MiB → arrayLength<u32> = 262144
+    wgpu::BufferDescriptor descA;
+    descA.size = kBufASize;
+    descA.usage = wgpu::BufferUsage::Storage;
+    wgpu::Buffer bufA = device.CreateBuffer(&descA);
+
+    // Buffer B: large allocation, but only a 64-byte sub-range will be bound to
+    // the COMPUTE-visible storage slot. With the collision the shader receives
+    // A's bound size as B's arrayLength, so the Robustness clamp on B[idx]
+    // permits writes far past the 64-byte bound range.
+    constexpr uint32_t kBufBSize = 16384;  // 16 KiB underlying allocation
+    constexpr uint32_t kBufBBound = 64;    // 64 bytes bound → arrayLength = 16
+    constexpr uint32_t kOobIndex = 1000;   // Target B[1000] (byte 4000) – OOB
+    constexpr uint32_t kOobMarker = 0xDEAD4A11u;
+    std::vector<uint32_t> initialData(kBufBSize / 4, 0u);
+    wgpu::Buffer bufB = utils::CreateBufferFromData(
+        device, initialData.data(), initialData.size() * sizeof(uint32_t),
+        wgpu::BufferUsage::Storage | wgpu::BufferUsage::CopySrc | wgpu::BufferUsage::CopyDst);
+
+    // PipelineLayoutGL assigns ssboIndex by iterating groups in order, so:
+    //   group 0 (A, FRAGMENT-only) → ssboIndex 0
+    //   group 1 (B, COMPUTE)       → ssboIndex 1  → glIndex_B = 1
+    // Choose A's WGSL @binding == glIndex_B (= 1) so that A's pre-remap key
+    // {0,1} equals B's post-remap key {0, glIndex_B}.
+    wgpu::BindGroupLayout bglA = utils::MakeBindGroupLayout(
+        device, {{1, wgpu::ShaderStage::Fragment, wgpu::BufferBindingType::ReadOnlyStorage}});
+    wgpu::BindGroupLayout bglB = utils::MakeBindGroupLayout(
+        device, {{0, wgpu::ShaderStage::Compute, wgpu::BufferBindingType::Storage}});
+
+    std::string shaderSource = R"(
+        @group(1) @binding(0) var<storage, read_write> B : array<u32>;
+        @compute @workgroup_size(1) fn main() {
+            B[0] = arrayLength(&B);
+            // Robustness wraps this as B[min(idx, arrayLength(&B)-1)].
+            // With the bug arrayLength(&B) is huge, so the write lands at
+            // index 1000 – past the 64-byte bound range.
+            B[)" + std::to_string(kOobIndex) +
+                               R"(u] = )" + std::to_string(kOobMarker) + R"(u;
+        })";
+
+    wgpu::ComputePipelineDescriptor pipelineDesc;
+    pipelineDesc.layout = utils::MakePipelineLayout(device, {bglA, bglB});
+    pipelineDesc.compute.module = utils::CreateShaderModule(device, shaderSource);
+    pipelineDesc.compute.entryPoint = "main";
+    wgpu::ComputePipeline pipeline = device.CreateComputePipeline(&pipelineDesc);
+
+    wgpu::BindGroup bgA = utils::MakeBindGroup(device, bglA, {{1, bufA, 0, kBufASize}});
+    wgpu::BindGroup bgB = utils::MakeBindGroup(device, bglB, {{0, bufB, 0, kBufBBound}});
+
+    wgpu::CommandEncoder encoder = device.CreateCommandEncoder();
+    wgpu::ComputePassEncoder pass = encoder.BeginComputePass();
+    pass.SetPipeline(pipeline);
+    pass.SetBindGroup(0, bgA);
+    pass.SetBindGroup(1, bgB);
+    pass.DispatchWorkgroups(1);
+    pass.End();
+    wgpu::CommandBuffer commands = encoder.Finish();
+    queue.Submit(1, &commands);
+
+    EXPECT_BUFFER_U32_EQ(kBufBBound / 4, bufB, 0);
+    EXPECT_BUFFER_U32_EQ(0u, bufB, kOobIndex * sizeof(uint32_t));
+}
+
+DAWN_INSTANTIATE_TEST(OpArrayLengthVisibilityCollisionTest,
+                      OpenGLESBackend(),
+                      OpenGLESBackend({"gl_use_array_length_from_uniform"}),
+                      VulkanBackend());
+
 enum class TieredLimits {
     No,
     Yes,
Loading diff…

Original Bug Report

reported by vm...@google.com

Bypass of Dawn's SSBO robustness checks via stage-filtering mismatch in OpenGL backend

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: A potential logic flaw in Dawn’s OpenGL/GLES backend stage-filtering can cause a mapping collision during Tint’s compilation phase. This mismatch allows an inactive shader storage buffer’s size lookup configuration to overwrite an active buffer’s configuration, potentially causing robustness checks to use an incorrect (larger) buffer size. An attacker could exploit this behavior to perform out-of-bounds memory reads and writes on storage buffers.

Affected files:

  • third_party/dawn/src/dawn/native/opengl/ShaderModuleGL.cpp
  • third_party/dawn/src/tint/lang/glsl/writer/raise/raise.cc

Estimated timestamp from git blame: 2025-04-10

Description

There is a potential stage-filtering inconsistency in Dawn’s OpenGL backend during the translation of WGSL to GLSL. When generating the layout metadata for shader compilation, Dawn constructs remapping tables and uniform-based array length tables. However, the logic that populates these tables handles stage-visibility differently:

  1. Active Binding Remapping: GenerateBindingRemapping (defined in third_party/dawn/src/dawn/native/TintUtils.h) correctly checks stage-visibility using StageBit(stage). Non-stage-visible bindings are skipped:

    for (const auto& [bindingNumber, apiBindingIndex] : bgl->GetBindingMap()) {
    if (!(bgl->GetAPIBindingInfo(apiBindingIndex).visibility & StageBit(stage))) {
        continue;
    }

  2. Array Length Uniform Options: GenerateArrayLengthFromuniformData (defined in third_party/dawn/src/dawn/native/opengl/ShaderModuleGL.cpp) does not check stage-visibility. It populates bindpoint_to_size_index for all storage buffers declared in the pipeline layout, regardless of whether they are active or visible in the current stage:

    for (BindingIndex binding : bgl->GetBufferIndices()) {
    const BindingInfo& bindingInfo = bgl->GetBindingInfo(binding);
    switch (std::get<BufferBindingInfo>(bindingInfo.bindingLayout).type) {
        case wgpu::BufferBindingType::Storage:
        // ...
        tint::BindingPoint srcBindingPoint = {uint32_t(group), uint32_t(bindingInfo.binding)};
        options.bindpoint_to_size_index.emplace(srcBindingPoint, uint32_t(ssboIndex));

Potential Exploit Mechanism

An attacker could potentially configure a pipeline layout containing:

  • Buffer A (inactive in the compute stage) at Bind Group 0, Binding N (pre-remap: {0, N}).
  • Buffer B (active in the compute stage) at Bind Group 1, Binding M (pre-remap: {1, M}).

During compilation of a compute shader that only utilizes Buffer B:

  1. Remapping Stage: GenerateBindingRemapping filters out Buffer A. It remaps Buffer B from {1, M} to a flat 1D binding point (e.g., {0, N}). Consequently, the remapper data used by Tint (remapper_data) maps {1, M} -> {0, N}. It contains no entries for Buffer A’s original binding point {0, N}.

  2. Array Length Options Stage: GenerateArrayLengthFromuniformData populates options for both buffers:

    • {1, M} -> index_B (for Buffer B)
    • {0, N} -> index_A (for Buffer A)

  3. Map Collision in Tint: During Tint’s raise phase in third_party/dawn/src/tint/lang/glsl/writer/raise/raise.cc, Tint resolves the pre-remapped keys of bindpoint_to_size_index to their post-remapped targets to build a final size_indices map:

    std::unordered_map<BindingPoint, uint32_t> size_indices;
for (auto& pair : options.array_length_from_uniform.bindpoint_to_size_index) {
    auto& bp = pair.first;
    auto& index = pair.second;
    where = remapper_data.find(bp);
    if (where != remapper_data.end()) {
        size_indices[where->second] = index;
    } else {
        size_indices[bp] = index;
    }
}
    • When processing Buffer B ({1, M}), Tint resolves it via remapper_data to {0, N} and sets size_indices[{0, N}] = index_B.
    • When processing Buffer A ({0, N}), Tint does not find it in remapper_data and falls back to the else block, setting size_indices[{0, N}] = index_A.
    • This causes a key collision inside the std::unordered_map. Buffer A’s entry overwrites Buffer B’s size index configuration, setting it to index_A instead of index_B.

  4. Robustness Bypass: When Tint substitutes .length() calls or injects robustness-clamping bound checks for Buffer B (binding point {0, N}), it retrieves the size from index_A (the index corresponding to Buffer A). If the attacker binds a very large buffer to Buffer A and a small buffer to Buffer B, the runtime clamping checks for Buffer B will clamp against the size of Buffer A, allowing out-of-bounds reads and writes on the memory backing Buffer B.

Note: These steps are based on static code analysis; our tooling does not currently have the environment or capability to run a live dynamic Proof-of-Concept to verify execution.

Impact

This flaw allows a shader to bypass WebGPU/Tint’s software robustness safety guarantees. On GLES backends (such as older Android devices with PowerVR or Nvidia GPUs), the GPU process is typically unsandboxed, potentially allowing GPU memory corruption or sandbox escape.

Suggested Fix

Modify GenerateArrayLengthFromuniformData in third_party/dawn/src/dawn/native/opengl/ShaderModuleGL.cpp to filter out storage buffers that are not visible to the shader stage currently being compiled, matching the visibility filter used in GenerateBindingRemapping:

for (BindingIndex binding : bgl->GetBufferIndices()) {
    const BindingInfo& bindingInfo = bgl->GetBindingInfo(binding);
    if (!(bindingInfo.visibility & StageBit(stage))) {
        continue;
    }
    // ...

Evaluated with Chrome root at commit: b1520ef4a76878853a31f0943b565e42060edec8

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
Uploaded PoC files
On This Page