CVE-2026-14420

Critical chrome OOB 📄 Reporter bug report 🔧 Commit mapped

Overview

Critical

Severity

—

CVSS

Exploited ITW

Fixed

Fix Status

ImpactOut of bounds read and write in Dawn

DescriptionOut of bounds read and write in Dawn

ComponentDawn

Bug ClassOOB

Tracker517031505

Fix commit73591422f504 (dawn) +94/-1

CISA KEVNot listed

CreditedGoogle

Disclosed2026-06-30

Changed Functions

Function	Change	Notes
`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

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:

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;
    }

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:

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}.
Array Length Options Stage: GenerateArrayLengthFromuniformData populates options for both buffers:
- {1, M} -> index_B (for Buffer B)
- {0, N} -> index_A (for Buffer A)
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.
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

References

Vendor advisory Fix commit Bug tracker NVD

On This Page