PatchGap
Chrome · Dawn

CVE-2026-9889

OOB in Dawn
Critical chrome OOB 📄 Reporter bug report 🔧 Commit mapped

Overview

Critical
Severity
—
CVSS
No
Exploited ITW
Fixed
Fix Status
ImpactOut of bounds read and write in Dawn
DescriptionOut of bounds read and write in Dawn
ComponentDawn
Bug ClassOOB
Tracker511727159
Fix commit0bc97cb2f50a (dawn) +127/-0
CISA KEVNot listed
CreditedGoogle
Disclosed2026-05-27

Changed Functions

FunctionChangeNotes
if
src/dawn/native/opengl/CommandBufferGL.cpp
 modified
if
src/dawn/native/opengl/PipelineLayoutGL.cpp
 modified
GLArrayLengthOverflowTest
src/dawn/tests/end2end/OpArrayLengthTests.cpp
 modified
TEST_P
src/dawn/tests/end2end/OpArrayLengthTests.cpp
 modified
for
src/dawn/tests/end2end/OpArrayLengthTests.cpp
 modified

Files Changed

  • src/dawn/native/opengl/CommandBufferGL.cpp
  • src/dawn/native/opengl/PipelineLayoutGL.cpp
  • src/dawn/tests/end2end/OpArrayLengthTests.cpp

Patch

view commit
From 0bc97cb2f50a092b4f332c8cb2b3285de7ccc39d Mon Sep 17 00:00:00 2001
From: Stephen White <senorblanco@chromium.org>
Date: Thu, 21 May 2026 10:10:38 -0700
Subject: [PATCH] GL: fix issue with ShaderStage::None bindings.

If more than 96 bindings with visibility ShaderStage::None are used,
they will exceed kGLMaxShaderStorageBufferBindingsReported (aka
GL_MAX_SHADER_STORAGE_BUFFER_BINDINGS) in the GL backend.
Dawn validation will not catch this, since it skips such bindings.
The fix is to skip all of the bindings with visibility ShaderStage::None
when computing binding indices in pipeline layout construction and
when applying bind groups in the GL backend.

Bug: 511727159
Change-Id: Ibc873a5afab4e73c12ad25b7d7cedcf48f35ca6e
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/309876
Reviewed-by: Corentin Wallez <cwallez@chromium.org>
Commit-Queue: Stephen White <senorblanco@chromium.org>
---

diff --git a/src/dawn/native/opengl/CommandBufferGL.cpp b/src/dawn/native/opengl/CommandBufferGL.cpp
index 72448b7..ccd52e3 100644
--- a/src/dawn/native/opengl/CommandBufferGL.cpp
+++ b/src/dawn/native/opengl/CommandBufferGL.cpp
@@ -349,6 +349,9 @@
 
         for (BindingIndex bindingIndex : Range(group->GetLayout()->GetBindingCount())) {
             const BindingInfo& bindingInfo = group->GetLayout()->GetBindingInfo(bindingIndex);
+            if (bindingInfo.visibility == wgpu::ShaderStage::None) {
+                continue;
+            }
             DAWN_TRY(MatchVariant(
                 bindingInfo.bindingLayout,
                 [&](const BufferBindingInfo& layout) -> MaybeError {
diff --git a/src/dawn/native/opengl/PipelineLayoutGL.cpp b/src/dawn/native/opengl/PipelineLayoutGL.cpp
index 98bfff0..8479add 100644
--- a/src/dawn/native/opengl/PipelineLayoutGL.cpp
+++ b/src/dawn/native/opengl/PipelineLayoutGL.cpp
@@ -48,6 +48,9 @@
 
         for (BindingIndex bindingIndex{0}; bindingIndex < bgl->GetBindingCount(); ++bindingIndex) {
             const BindingInfo& bindingInfo = bgl->GetBindingInfo(bindingIndex);
+            if (bindingInfo.visibility == wgpu::ShaderStage::None) {
+                continue;
+            }
             MatchVariant(
                 bindingInfo.bindingLayout,
                 [&](const BufferBindingInfo& layout) {
diff --git a/src/dawn/tests/end2end/OpArrayLengthTests.cpp b/src/dawn/tests/end2end/OpArrayLengthTests.cpp
index e1b4179..10b5ed4 100644
--- a/src/dawn/tests/end2end/OpArrayLengthTests.cpp
+++ b/src/dawn/tests/end2end/OpArrayLengthTests.cpp
@@ -432,5 +432,126 @@
                          VulkanBackend(), WebGPUBackend()},
                         {TieredLimits::No, TieredLimits::Yes});
 
+class GLArrayLengthOverflowTest : public DawnTest {
+  protected:
+    void GetRequiredLimits(const dawn::utils::ComboLimits& supported,
+                           dawn::utils::ComboLimits& required) override {
+        supported.UnlinkedCopyTo(&required);
+    }
+};
+
+// Test that using more than 96 ShaderStage::None bind group entries
+// (which don't count against Dawn's validation limit) don't cause GL
+// errors and failed buffer transfers.
+TEST_P(GLArrayLengthOverflowTest, VisibilityNoneOverflowsArrayLengthBuffer) {
+    DAWN_TEST_UNSUPPORTED_IF(GetSupportedLimits().maxBindGroups < 4);
+
+    constexpr uint32_t kLargeSize = 512u;
+    constexpr uint32_t kSmallSize = 256u;
+    constexpr uint32_t kPadPerGroup = 33;
+
+    wgpu::BufferDescriptor bd;
+    bd.size = kLargeSize;
+    bd.usage = wgpu::BufferUsage::Storage | wgpu::BufferUsage::CopySrc;
+    wgpu::Buffer largeBuf = device.CreateBuffer(&bd);
+
+    // Get the arrayLength() of the passed-in storage buffer, and store it into
+    // the first element of the array.
+    wgpu::ComputePipelineDescriptor primeDesc;
+    primeDesc.compute.module = utils::CreateShaderModule(device, R"(
+        @group(0) @binding(0) var<storage, read_write> a : array<u32>;
+        @compute @workgroup_size(1) fn main() {
+            a[0] = arrayLength(&a);
+        })");
+    wgpu::ComputePipeline primePipeline = device.CreateComputePipeline(&primeDesc);
+    wgpu::BindGroupLayout primeBGL = primePipeline.GetBindGroupLayout(0);
+    wgpu::BindGroup primeBG = utils::MakeBindGroup(device, primeBGL, {{0, largeBuf}});
+
+    {
+        wgpu::CommandEncoder enc = device.CreateCommandEncoder();
+        wgpu::ComputePassEncoder pass = enc.BeginComputePass();
+        pass.SetPipeline(primePipeline);
+        pass.SetBindGroup(0, primeBG);
+        pass.DispatchWorkgroups(1);
+        pass.End();
+        wgpu::CommandBuffer cb = enc.Finish();
+        queue.Submit(1, &cb);
+    }
+
+    bd.size = kSmallSize;
+    bd.usage = wgpu::BufferUsage::Storage | wgpu::BufferUsage::CopySrc;
+    wgpu::Buffer smallBuf = device.CreateBuffer(&bd);
+
+    bd.size = 4;
+    bd.usage = wgpu::BufferUsage::Storage;
+    wgpu::Buffer tinyBuffer = device.CreateBuffer(&bd);
+
+    wgpu::BindGroupLayout bgl0 = utils::MakeBindGroupLayout(
+        device, {{0, wgpu::ShaderStage::Compute, wgpu::BufferBindingType::Storage}});
+
+    std::vector<wgpu::BindGroupLayoutEntry> padEntries(kPadPerGroup);
+    for (uint32_t i = 0; i < kPadPerGroup; i++) {
+        padEntries[i].binding = i;
+        padEntries[i].visibility = wgpu::ShaderStage::None;
+        padEntries[i].buffer.type = wgpu::BufferBindingType::ReadOnlyStorage;
+    }
+    wgpu::BindGroupLayoutDescriptor padDesc;
+    padDesc.entryCount = padEntries.size();
+    padDesc.entries = padEntries.data();
+    wgpu::BindGroupLayout bglPad = device.CreateBindGroupLayout(&padDesc);
+
+    wgpu::BindGroupLayout bgls[] = {bgl0, bglPad, bglPad, bglPad};
+    wgpu::PipelineLayoutDescriptor plDesc;
+    plDesc.bindGroupLayoutCount = 4;
+    plDesc.bindGroupLayouts = bgls;
+    wgpu::PipelineLayout manyBindingsPL = device.CreatePipelineLayout(&plDesc);
+
+    wgpu::ComputePipelineDescriptor manyBindingsDesc;
+    manyBindingsDesc.layout = manyBindingsPL;
+    manyBindingsDesc.compute.module = primeDesc.compute.module;
+    wgpu::ComputePipeline manyBindingsPipeline = device.CreateComputePipeline(&manyBindingsDesc);
+
+    wgpu::BindGroup bg0 = utils::MakeBindGroup(device, bgl0, {{0, smallBuf}});
+
+    std::vector<wgpu::BindGroupEntry> padBinds(kPadPerGroup);
+    for (uint32_t i = 0; i < kPadPerGroup; i++) {
+        padBinds[i].binding = i;
+        padBinds[i].buffer = tinyBuffer;
+    }
+    wgpu::BindGroupDescriptor bgPadDesc;
+    bgPadDesc.layout = bglPad;
+    bgPadDesc.entryCount = padBinds.size();
+    bgPadDesc.entries = padBinds.data();
+    wgpu::BindGroup bgPad = device.CreateBindGroup(&bgPadDesc);
+
+    {
+        wgpu::CommandEncoder enc = device.CreateCommandEncoder();
+        wgpu::ComputePassEncoder pass = enc.BeginComputePass();
+        pass.SetPipeline(manyBindingsPipeline);
+        pass.SetBindGroup(0, bg0);
+        pass.SetBindGroup(1, bgPad);
+        pass.SetBindGroup(2, bgPad);
+        pass.SetBindGroup(3, bgPad);
+        pass.DispatchWorkgroups(1);
+        pass.End();
+        wgpu::CommandBuffer cb = enc.Finish();
+        queue.Submit(1, &cb);
+    }
+
+    // Check that the stored arrayLength is the (new) small buffer length
+    // and not the (stale) large buffer length.
+    EXPECT_BUFFER_U32_EQ(kSmallSize / 4u, smallBuf, 0);
+}
+
+DAWN_INSTANTIATE_TEST(GLArrayLengthOverflowTest,
+                      D3D11Backend(),
+                      D3D12Backend(),
+                      MetalBackend(),
+                      OpenGLBackend(),
+                      OpenGLESBackend(),
+                      OpenGLESBackend({"gl_use_array_length_from_uniform"}),
+                      VulkanBackend(),
+                      WebGPUBackend());
+
 }  // anonymous namespace
 }  // namespace dawn
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 e1b4179..10b5ed4 100644
--- a/src/dawn/tests/end2end/OpArrayLengthTests.cpp
+++ b/src/dawn/tests/end2end/OpArrayLengthTests.cpp
@@ -432,5 +432,126 @@
                          VulkanBackend(), WebGPUBackend()},
                         {TieredLimits::No, TieredLimits::Yes});
 
+class GLArrayLengthOverflowTest : public DawnTest {
+  protected:
+    void GetRequiredLimits(const dawn::utils::ComboLimits& supported,
+                           dawn::utils::ComboLimits& required) override {
+        supported.UnlinkedCopyTo(&required);
+    }
+};
+
+// Test that using more than 96 ShaderStage::None bind group entries
+// (which don't count against Dawn's validation limit) don't cause GL
+// errors and failed buffer transfers.
+TEST_P(GLArrayLengthOverflowTest, VisibilityNoneOverflowsArrayLengthBuffer) {
+    DAWN_TEST_UNSUPPORTED_IF(GetSupportedLimits().maxBindGroups < 4);
+
+    constexpr uint32_t kLargeSize = 512u;
+    constexpr uint32_t kSmallSize = 256u;
+    constexpr uint32_t kPadPerGroup = 33;
+
+    wgpu::BufferDescriptor bd;
+    bd.size = kLargeSize;
+    bd.usage = wgpu::BufferUsage::Storage | wgpu::BufferUsage::CopySrc;
+    wgpu::Buffer largeBuf = device.CreateBuffer(&bd);
+
+    // Get the arrayLength() of the passed-in storage buffer, and store it into
+    // the first element of the array.
+    wgpu::ComputePipelineDescriptor primeDesc;
+    primeDesc.compute.module = utils::CreateShaderModule(device, R"(
+        @group(0) @binding(0) var<storage, read_write> a : array<u32>;
+        @compute @workgroup_size(1) fn main() {
+            a[0] = arrayLength(&a);
+        })");
+    wgpu::ComputePipeline primePipeline = device.CreateComputePipeline(&primeDesc);
+    wgpu::BindGroupLayout primeBGL = primePipeline.GetBindGroupLayout(0);
+    wgpu::BindGroup primeBG = utils::MakeBindGroup(device, primeBGL, {{0, largeBuf}});
+
+    {
+        wgpu::CommandEncoder enc = device.CreateCommandEncoder();
+        wgpu::ComputePassEncoder pass = enc.BeginComputePass();
+        pass.SetPipeline(primePipeline);
+        pass.SetBindGroup(0, primeBG);
+        pass.DispatchWorkgroups(1);
+        pass.End();
+        wgpu::CommandBuffer cb = enc.Finish();
+        queue.Submit(1, &cb);
+    }
+
+    bd.size = kSmallSize;
+    bd.usage = wgpu::BufferUsage::Storage | wgpu::BufferUsage::CopySrc;
+    wgpu::Buffer smallBuf = device.CreateBuffer(&bd);
+
+    bd.size = 4;
+    bd.usage = wgpu::BufferUsage::Storage;
+    wgpu::Buffer tinyBuffer = device.CreateBuffer(&bd);
+
+    wgpu::BindGroupLayout bgl0 = utils::MakeBindGroupLayout(
+        device, {{0, wgpu::ShaderStage::Compute, wgpu::BufferBindingType::Storage}});
+
+    std::vector<wgpu::BindGroupLayoutEntry> padEntries(kPadPerGroup);
+    for (uint32_t i = 0; i < kPadPerGroup; i++) {
+        padEntries[i].binding = i;
+        padEntries[i].visibility = wgpu::ShaderStage::None;
+        padEntries[i].buffer.type = wgpu::BufferBindingType::ReadOnlyStorage;
+    }
+    wgpu::BindGroupLayoutDescriptor padDesc;
+    padDesc.entryCount = padEntries.size();
+    padDesc.entries = padEntries.data();
+    wgpu::BindGroupLayout bglPad = device.CreateBindGroupLayout(&padDesc);
+
+    wgpu::BindGroupLayout bgls[] = {bgl0, bglPad, bglPad, bglPad};
+    wgpu::PipelineLayoutDescriptor plDesc;
+    plDesc.bindGroupLayoutCount = 4;
+    plDesc.bindGroupLayouts = bgls;
+    wgpu::PipelineLayout manyBindingsPL = device.CreatePipelineLayout(&plDesc);
+
+    wgpu::ComputePipelineDescriptor manyBindingsDesc;
+    manyBindingsDesc.layout = manyBindingsPL;
+    manyBindingsDesc.compute.module = primeDesc.compute.module;
+    wgpu::ComputePipeline manyBindingsPipeline = device.CreateComputePipeline(&manyBindingsDesc);
+
+    wgpu::BindGroup bg0 = utils::MakeBindGroup(device, bgl0, {{0, smallBuf}});
+
+    std::vector<wgpu::BindGroupEntry> padBinds(kPadPerGroup);
+    for (uint32_t i = 0; i < kPadPerGroup; i++) {
+        padBinds[i].binding = i;
+        padBinds[i].buffer = tinyBuffer;
+    }
+    wgpu::BindGroupDescriptor bgPadDesc;
+    bgPadDesc.layout = bglPad;
+    bgPadDesc.entryCount = padBinds.size();
+    bgPadDesc.entries = padBinds.data();
+    wgpu::BindGroup bgPad = device.CreateBindGroup(&bgPadDesc);
+
+    {
+        wgpu::CommandEncoder enc = device.CreateCommandEncoder();
+        wgpu::ComputePassEncoder pass = enc.BeginComputePass();
+        pass.SetPipeline(manyBindingsPipeline);
+        pass.SetBindGroup(0, bg0);
+        pass.SetBindGroup(1, bgPad);
+        pass.SetBindGroup(2, bgPad);
+        pass.SetBindGroup(3, bgPad);
+        pass.DispatchWorkgroups(1);
+        pass.End();
+        wgpu::CommandBuffer cb = enc.Finish();
+        queue.Submit(1, &cb);
+    }
+
+    // Check that the stored arrayLength is the (new) small buffer length
+    // and not the (stale) large buffer length.
+    EXPECT_BUFFER_U32_EQ(kSmallSize / 4u, smallBuf, 0);
+}
+
+DAWN_INSTANTIATE_TEST(GLArrayLengthOverflowTest,
+                      D3D11Backend(),
+                      D3D12Backend(),
+                      MetalBackend(),
+                      OpenGLBackend(),
+                      OpenGLESBackend(),
+                      OpenGLESBackend({"gl_use_array_length_from_uniform"}),
+                      VulkanBackend(),
+                      WebGPUBackend());
+
 }  // anonymous namespace
 }  // namespace dawn
Loading diff…

Original Bug Report

reported by vm...@google.com

Dawn GLES: Bypassed SSBO limit via ShaderStage::None allows OOB GPU memory access

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 without the Chrome Security team. Please see https://chromium.googlesource.com/chromium/src/+/main/docs/security/ai-generated-security-bugs-faq.md for more information.

Overview: Storage buffer bindings with wgpu::ShaderStage::None visibility bypass per-stage limit validation, allowing a pipeline to exceed internal SSBO capacity. This inflates internal SSBO indices, causing a silent failure when uploading length data to an internal uniform buffer. The stale or zeroed buffer values cause Tint’s robustness transforms to underflow, resulting in an ineffective clamp and allowing arbitrary OOB GPU memory reads/writes.

Affected files:

  • third_party/dawn/src/dawn/native/BindingInfo.cpp
  • third_party/dawn/src/dawn/native/opengl/PipelineLayoutGL.cpp
  • third_party/dawn/src/dawn/native/opengl/CommandBufferGL.cpp
  • third_party/dawn/src/dawn/native/opengl/DeviceGL.cpp
  • third_party/dawn/src/dawn/native/opengl/PipelineLayoutGL.h
  • third_party/dawn/src/dawn/native/opengl/ShaderModuleGL.cpp
  • third_party/dawn/src/dawn/native/opengl/BufferGL.cpp
  • third_party/dawn/src/dawn/native/opengl/UtilsGL.h
  • third_party/dawn/src/tint/lang/core/ir/transform/array_length_from_uniform.cc

Estimated timestamp from git blame: 2025-12-04

Summary

A potential vulnerability in Dawn’s OpenGL/GLES backend allows for out-of-bounds (OOB) reads and writes in the GPU process.

The issue stems from a combination of factors:

  1. A validation gap allows the creation of a BindGroupLayout with an excessive number of storage buffers by setting their visibility to wgpu::ShaderStage::None.
  2. The OpenGL backend blindly assigns an internal ssboIndex to all storage buffers, causing the index to exceed the internal tracker’s maximum capacity (96).
  3. Updates to the internal uniform buffer holding SSBO array lengths fail silently due to an ignored GL_INVALID_VALUE error in Release builds.
  4. When the internal uniform buffer retains stale or out-of-bounds (zero) length values, integer underflows in Tint’s robustness and array length transforms result in massive bounds clamps, bypassing security checks.

Technical Details and Potential Exploitation Steps

Our codebase investigator agent analyzed the flow and suggests the following steps could trigger the vulnerability:

  1. Prime the Uniform Buffer (Optional but possible): An attacker dispatches a valid compute shader using a normal layout, binding a maximum-sized storage buffer. This populates Dawn’s internal mArrayLengthBuffer (a 384-byte buffer designed to hold up to 96 SSBO lengths) with very large length values on the GPU.

  2. Bypass Validation with ShaderStage::None: The attacker constructs a malicious BindGroupLayoutDescriptor containing over 96 dummy storage buffer bindings. By setting the visibility of these dummy bindings to wgpu::ShaderStage::None, they bypass the maxStorageBuffersPerShaderStage limit (16). In BindingInfo.cpp, IncrementBindingCounts skips incrementing the stage counters for None visibility (IterateStages(0) is empty). The total bindings easily fit within the maxBindingsPerBindGroup limit (1000).

  3. Inflate ssboIndex: The attacker creates a PipelineLayout with this layout. The PipelineLayoutGL constructor sequentially assigns an ssboIndex to every storage buffer binding, inflating the index of the target storage buffer far past 96.

  4. Silent Buffer Upload Failure: The attacker dispatches a malicious shader that performs OOB dynamic array indexing on the target storage buffer.

    • CommandBufferGL::ApplyBindGroup calculates an update range for mArrayLengthBuffer based on the inflated ssboIndex.
    • ApplyInternalArrayLengthUniforms calls glBufferSubData with an update size that exceeds the physically allocated 384 bytes of mArrayLengthBuffer.
    • The OpenGL driver returns GL_INVALID_VALUE, aborting the upload. However, because the call is wrapped in DAWN_GL_TRY, the error is silently ignored in Release builds (where DAWN_ENABLE_ASSERTS is disabled). The uniform buffer on the GPU is not updated.

  5. Robustness Bypass via Underflow: The malicious shader executes. Because the uniform buffer wasn’t updated, a read from the out-of-bounds index returns 0 (due to standard GL robustness rules).

    • In array_length_from_uniform.cc, the array size is calculated as total_buffer_size - member->Offset(). If total_buffer_size is 0, this unsigned subtraction underflows to a massive value.
    • Even if the offset is 0, robustness.cc calculates the clamp limit as arrayLength() - 1. 0 - 1 underflows to 0xFFFFFFFF.
    • The clamp operation becomes min(index, 0xFFFFFFFF), which effectively disables the clamp, allowing the attacker-controlled index to be used directly for OOB GPU memory reads and writes.

Impact

This vulnerability provides a potential path to arbitrary shader-driven out-of-bounds read and write capabilities in the GPU virtual address space. On Android platforms where the GPU process is frequently unsandboxed, this could lead to a sandbox escape and Remote Code Execution (RCE) from unprivileged Web content.

Suggested Fixes

  1. Validation: In BindingInfo.cpp, enforce a hard, cross-stage limit on the total number of storage buffer bindings allowed in a PipelineLayout, regardless of shader stage visibility. Do not rely solely on per-stage accumulated counts.
  2. Error Handling: Ensure that failures during internal buffer synchronization (e.g., glBufferSubData in ApplyInternalArrayLengthUniforms) cause the dispatch to fail safely, rather than silently ignoring the error via DAWN_GL_TRY in non-assert builds.
  3. Tint Robustness: Guard against integer underflows in array_length_from_uniform.cc and robustness.cc. Ensure that if a buffer size is smaller than the array offset (or exactly 0), the resulting clamped array length does not wrap around to a massive maximum value.

Evaluated with Chrome root at commit: eca8648a4e1cdfdda68c495a6003059fed641955

Results so far have been promising, but there can be wrong deductions. If this proves to be a false positive, please close as WAI; 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
On This Page