PatchGap
Chrome · Dawn

CVE-2026-9907

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

Overview

High
Severity
—
CVSS
No
Exploited ITW
Fixed
Fix Status
ImpactOut of bounds read in Dawn
DescriptionOut of bounds read in Dawn
ComponentDawn
Bug ClassOOB
Tracker499091269
Fix commit6558ab4d77ba (dawn) +68/-10
CISA KEVNot listed
CreditedGoogle
Disclosed2026-05-27

Changed Functions

FunctionChangeNotes
GenerateBindingRemapping
src/dawn/native/d3d12/ShaderModuleD3D12.cpp
 modified
TEST_F
src/tint/lang/hlsl/writer/writer_test.cc
 modified

Files Changed

  • src/dawn/native/d3d12/ShaderModuleD3D12.cpp
  • src/tint/lang/hlsl/writer/writer_test.cc

Patch

view commit
From 6558ab4d77ba037ef69b3949d31831f6ea146fde Mon Sep 17 00:00:00 2001
From: Antonio Maiorano <amaiorano@google.com>
Date: Thu, 23 Apr 2026 14:38:36 -0700
Subject: [PATCH] [tint][hlsl] Fix bindings in ignored_by_robustness_transform not being remapped

Values in the hlsl writer options, ignored_by_robustness_transform were
populated with WGSL space binding values, but were being interpreted as
HLSL space values. Remap them to HLSL space in Dawn.

Bug: 499091269
Change-Id: Ie7d6cdcb9bb945d92d8d26b2df0b778e9ad67ea4
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/304415
Reviewed-by: Corentin Wallez <cwallez@chromium.org>
Reviewed-by: dan sinclair <dsinclair@chromium.org>
Commit-Queue: Antonio Maiorano <amaiorano@google.com>
---

diff --git a/src/dawn/native/d3d12/ShaderModuleD3D12.cpp b/src/dawn/native/d3d12/ShaderModuleD3D12.cpp
index c210f4e..71b6306 100644
--- a/src/dawn/native/d3d12/ShaderModuleD3D12.cpp
+++ b/src/dawn/native/d3d12/ShaderModuleD3D12.cpp
@@ -184,14 +184,14 @@
         };
     }
 
-    tint::Bindings bindings =
-        GenerateBindingRemapping(layout, stage, [&](BindGroupIndex group, BindingIndex index) {
-            const BindGroupLayout* bgl = ToBackend(layout->GetBindGroupLayout(group));
-            return tint::BindingPoint{
-                .group = uint32_t(group),
-                .binding = bgl->GetShaderRegister(index),
-            };
-        });
+    auto ToHLSLBindPoint = [&](BindGroupIndex group, BindingIndex index) {
+        const BindGroupLayout* bgl = ToBackend(layout->GetBindGroupLayout(group));
+        return tint::BindingPoint{
+            .group = uint32_t(group),
+            .binding = bgl->GetShaderRegister(index),
+        };
+    };
+    tint::Bindings bindings = GenerateBindingRemapping(layout, stage, ToHLSLBindPoint);
 
     std::vector<tint::BindingPoint> ignored_by_robustness;
     for (BindGroupIndex group : layout->GetBindGroupLayoutsMask()) {
@@ -233,8 +233,9 @@
             if ((bufferInfo.type == wgpu::BufferBindingType::Storage ||
                  bufferInfo.type == wgpu::BufferBindingType::ReadOnlyStorage) &&
                 !bufferInfo.hasDynamicOffset) {
-                ignored_by_robustness.emplace_back(tint::BindingPoint{
-                    .group = uint32_t(group), .binding = uint32_t(bindingInfo.binding)});
+                BindingIndex bindingIndex =
+                    bgl->AsBindingIndex(bgl->GetAPIBindingIndex(bindingInfo.binding));
+                ignored_by_robustness.emplace_back(ToHLSLBindPoint(group, bindingIndex));
             }
         }
 
diff --git a/src/tint/lang/hlsl/writer/writer_test.cc b/src/tint/lang/hlsl/writer/writer_test.cc
index f3f3a69..eb872d8 100644
--- a/src/tint/lang/hlsl/writer/writer_test.cc
+++ b/src/tint/lang/hlsl/writer/writer_test.cc
@@ -175,5 +175,62 @@
     EXPECT_EQ(output_.workgroup_info.storage_size, 0x100000000ull);
 }
 
+TEST_F(HlslWriterTest, IgnoredByRobustnessEntriesMustBeRemapped) {
+    // @group(0) @binding(7) var<uniform> u : array<vec4<u32>, 4>;
+    core::ir::Var* uniform_var = b.Var<uniform, array<vec4<u32>, 4>, core::Access::kRead>("u");
+    uniform_var->SetBindingPoint(0, 7);
+    b.ir.root_block->Append(uniform_var);
+
+    // @group(0) @binding(0) var<storage, read_write> idx_src : u32;
+    core::ir::Var* idx_var = b.Var("idx_src", ty.ptr(storage, ty.u32(), core::Access::kReadWrite));
+    idx_var->SetBindingPoint(0, 0);
+    b.ir.root_block->Append(idx_var);
+
+    // @group(0) @binding(1) var<storage, read_write> outp : vec4<u32>;
+    core::ir::Var* out_var =
+        b.Var("outp", ty.ptr(storage, ty.vec4<u32>(), core::Access::kReadWrite));
+    out_var->SetBindingPoint(0, 1);
+    b.ir.root_block->Append(out_var);
+
+    // @compute @workgroup_size(1) fn main() { outp = u[idx_src]; }
+    auto* func = b.ComputeFunction("main");
+    b.Append(func->Block(), [&] {
+        auto* idx = b.Load(idx_var);
+        auto* elem = b.Access(ty.ptr<uniform, vec4<u32>, core::Access::kRead>(), uniform_var, idx);
+        b.Store(out_var, b.Load(elem));
+        b.Return(func);
+    });
+
+    Options options{};
+    options.disable_robustness = false;
+    options.bindings.uniform.emplace(BindingPoint{0, 7}, BindingPoint{0, 0});
+    options.bindings.storage.emplace(BindingPoint{0, 0}, BindingPoint{0, 1});
+    options.bindings.storage.emplace(BindingPoint{0, 1}, BindingPoint{0, 2});
+
+    // We specify to ignore robustness on the storage pointer, idx_src, at hlsl (0,1),
+    // which should be a no-op since we only load from this pointer (no indexing).
+    // Robustness should apply to the uniform buffer at hlsl (0,0), though.
+    // However, when values in 'ignored_by_robustness_transform' contained the wgsl bindpoints,
+    // rather than the hlsl ones, and wgsl (0,0) was passed in for idx_src, the robustness transform
+    // would erroneously ignore the uniform buffer because it got remapped to hlsl (0,0).
+    options.ignored_by_robustness_transform.push_back(BindingPoint{0, 1});
+
+    auto result = Generate(options);
+    ASSERT_EQ(result, Success) << result.Failure().reason << output_.hlsl;
+    EXPECT_EQ(output_.hlsl, R"(
+cbuffer cbuffer_u : register(b0) {
+  uint4 u[4];
+};
+RWByteAddressBuffer idx_src : register(u1);
+RWByteAddressBuffer outp : register(u2);
+[numthreads(1, 1, 1)]
+void main() {
+  uint v = ((min(idx_src.Load(0u), 3u) * 16u) / 16u);
+  outp.Store4(0u, u[v]);
+}
+
+)");
+}
+
 }  // namespace
 }  // namespace tint::hlsl::writer
Loading diff…

Regression Test / PoC

shipped with the fix
diff --git a/src/tint/lang/hlsl/writer/writer_test.cc b/src/tint/lang/hlsl/writer/writer_test.cc
index f3f3a69..eb872d8 100644
--- a/src/tint/lang/hlsl/writer/writer_test.cc
+++ b/src/tint/lang/hlsl/writer/writer_test.cc
@@ -175,5 +175,62 @@
     EXPECT_EQ(output_.workgroup_info.storage_size, 0x100000000ull);
 }
 
+TEST_F(HlslWriterTest, IgnoredByRobustnessEntriesMustBeRemapped) {
+    // @group(0) @binding(7) var<uniform> u : array<vec4<u32>, 4>;
+    core::ir::Var* uniform_var = b.Var<uniform, array<vec4<u32>, 4>, core::Access::kRead>("u");
+    uniform_var->SetBindingPoint(0, 7);
+    b.ir.root_block->Append(uniform_var);
+
+    // @group(0) @binding(0) var<storage, read_write> idx_src : u32;
+    core::ir::Var* idx_var = b.Var("idx_src", ty.ptr(storage, ty.u32(), core::Access::kReadWrite));
+    idx_var->SetBindingPoint(0, 0);
+    b.ir.root_block->Append(idx_var);
+
+    // @group(0) @binding(1) var<storage, read_write> outp : vec4<u32>;
+    core::ir::Var* out_var =
+        b.Var("outp", ty.ptr(storage, ty.vec4<u32>(), core::Access::kReadWrite));
+    out_var->SetBindingPoint(0, 1);
+    b.ir.root_block->Append(out_var);
+
+    // @compute @workgroup_size(1) fn main() { outp = u[idx_src]; }
+    auto* func = b.ComputeFunction("main");
+    b.Append(func->Block(), [&] {
+        auto* idx = b.Load(idx_var);
+        auto* elem = b.Access(ty.ptr<uniform, vec4<u32>, core::Access::kRead>(), uniform_var, idx);
+        b.Store(out_var, b.Load(elem));
+        b.Return(func);
+    });
+
+    Options options{};
+    options.disable_robustness = false;
+    options.bindings.uniform.emplace(BindingPoint{0, 7}, BindingPoint{0, 0});
+    options.bindings.storage.emplace(BindingPoint{0, 0}, BindingPoint{0, 1});
+    options.bindings.storage.emplace(BindingPoint{0, 1}, BindingPoint{0, 2});
+
+    // We specify to ignore robustness on the storage pointer, idx_src, at hlsl (0,1),
+    // which should be a no-op since we only load from this pointer (no indexing).
+    // Robustness should apply to the uniform buffer at hlsl (0,0), though.
+    // However, when values in 'ignored_by_robustness_transform' contained the wgsl bindpoints,
+    // rather than the hlsl ones, and wgsl (0,0) was passed in for idx_src, the robustness transform
+    // would erroneously ignore the uniform buffer because it got remapped to hlsl (0,0).
+    options.ignored_by_robustness_transform.push_back(BindingPoint{0, 1});
+
+    auto result = Generate(options);
+    ASSERT_EQ(result, Success) << result.Failure().reason << output_.hlsl;
+    EXPECT_EQ(output_.hlsl, R"(
+cbuffer cbuffer_u : register(b0) {
+  uint4 u[4];
+};
+RWByteAddressBuffer idx_src : register(u1);
+RWByteAddressBuffer outp : register(u2);
+[numthreads(1, 1, 1)]
+void main() {
+  uint v = ((min(idx_src.Load(0u), 3u) * 16u) / 16u);
+  outp.Store4(0u, u[v]);
+}
+
+)");
+}
+
 }  // namespace
 }  // namespace tint::hlsl::writer
Loading diff…

Original Bug Report

reported by vm...@google.com

Potential Dawn D3D12 OOB Read via Coordinate Collision in Robustness Transform

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 security team.

Overview: A coordinate system mismatch between Dawn’s D3D12 backend and Tint’s robustness transform allows an attacker to bypass bounds checking for dynamic uniform buffers. By forcing a coordinate collision in the bind group layout, a shader can perform out-of-bounds reads via Root CBVs. This can result in a cross-origin information leak of adjacent GPU heap memory.

Affected files:

  • third_party/dawn/src/dawn/native/d3d12/ShaderModuleD3D12.cpp
  • third_party/dawn/src/tint/lang/hlsl/writer/raise/raise.cc
  • third_party/dawn/src/tint/lang/hlsl/writer/common/option_helpers.cc
  • third_party/dawn/src/tint/lang/core/ir/transform/robustness.cc

Estimated timestamp from git blame: 2025-09-24

Summary

A potential vulnerability exists in Dawn’s D3D12 backend where a coordinate mismatch between WGSL bindings and HLSL registers can cause Tint to erroneously disable software robustness clamping for dynamic uniform buffers. Because D3D12 uses Root Descriptor CBVs for dynamic uniform buffers—which lack hardware bounds checking—this bypass allows an attacker to read out-of-bounds GPU memory, leading to a cross-origin information leak.

Root Cause Analysis

The issue stems from how Dawn specifies which buffers should bypass Tint’s software bounds clamping (to prevent HLSL compiler errors on fixed-size arrays) and the order in which Tint applies its internal IR transformations.

  1. Sorting and Register Assignment: In BindGroupLayoutInternal::SortBindingsCompare, Dawn sorts bindings by type. BindingTypeOrder_DynamicBuffer has the highest priority (0). If a layout contains a dynamic uniform buffer and a non-dynamic storage buffer, the dynamic buffer is sorted first and assigned BindingIndex 0 (which becomes HLSL register b0). The storage buffer is assigned BindingIndex 1 (e.g., u1).
  2. Poisoning the Ignore List: In ShaderModuleD3D12.cpp, Dawn iterates over the layout to populate the ignored_by_robustness list with non-dynamic storage buffers. However, it records the buffer using its original WGSL coordinate, not its assigned HLSL register. If the storage buffer was at WGSL @binding(0), the entry {group: 0, binding: 0} is added to the ignore list.
  3. Tint Transformation Mismatch: In the Tint HLSL writer (raise.cc), the BindingRemapper transform runs first, mutating the IR variables from their WGSL coordinates to their new HLSL register assignments. The dynamic uniform buffer’s coordinate is changed to {0, 0} (because it was assigned register b0).
  4. Robustness Bypass: The Robustness transform runs immediately after. When processing pointer accesses to the dynamic uniform buffer, ShouldClamp checks IsRootVarIgnored. This function reads the buffer’s current (remapped) coordinate {0, 0} and checks the ignored_by_robustness list. Because the storage buffer previously added {0, 0} based on its WGSL index, the check returns true, and bounds clamping is erroneously disabled.

Potential Exploitation Steps

Note: These are theoretical steps based on code analysis; our tooling agent cannot execute code to verify.

  1. Setup Layout: An attacker creates a WebGPU BindGroupLayout with:
    • A non-dynamic Storage buffer at @group(0) @binding(0).
    • A Uniform buffer with hasDynamicOffset: true at a higher binding, e.g., @group(0) @binding(7).
  2. Trigger Collision: Due to Dawn’s sorting, the dynamic uniform buffer is assigned HLSL register 0 (b0). The storage buffer causes the WGSL coordinate {0, 0} to be added to the ignored_by_robustness list. Tint remaps the uniform buffer to {0, 0}, matching the ignore list and bypassing software clamping.
  3. OOB Read: The attacker dispatches a shader that accesses an array inside the uniform buffer using a large out-of-bounds index.
  4. Hardware execution: In PipelineLayoutD3D12.cpp, dynamic uniform buffers are mapped as D3D12_ROOT_DESCRIPTOR1 (Root CBVs). Because Root CBVs do not contain size information, the D3D12 hardware relies entirely on the (now bypassed) software clamping. The shader successfully reads arbitrary adjacent memory within the D3D12 GPU heap (up to the ~64KB constant buffer limit).
  5. Exfiltration: The leaked cross-origin data is written to an output buffer and mapped back to the CPU.

Suggested Fix

In the Tint HLSL writer (third_party/dawn/src/tint/lang/hlsl/writer/common/option_helpers.cc), the ignored_by_robustness_transform list should be remapped to HLSL registers using the RemapperData before the Robustness pass is executed. This logic should be integrated into PopulateBindingRelatedOptions, mirroring how the array_length_from_uniform configuration is already handled.

Evaluated with Chrome root at commit: ff3d2b74fa39431785bd60e51463b08fcc71ee33

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