CVE-2026-9907
Overview
Background
- `ignored_by_robustness_transform`
- a Tint HLSL-writer option listing bindings whose software bounds-clamping is skipped (so FXC does not reject OOB access on fixed-size arrays), populated by Dawn in
ShaderModule::Compile(ShaderModuleD3D12.cpp:196,:281). - WGSL vs HLSL binding coordinates
- a binding has an original WGSL
@group/@bindingand a separate HLSL register assigned by Dawn;bgl->GetShaderRegister(index)maps aBindingIndexto the HLSL register. - Tint `BindingRemapper` then `Robustness`
- Tint remaps IR variables to their HLSL registers and then runs the robustness (bounds-clamping) pass, which consults the ignore list using the variable’s current (remapped, HLSL-space) coordinate. (remapper ordering per reporter;
raise.ccnot in shown material) - D3D12 Root CBV
- dynamic uniform buffers are bound as root descriptors with no hardware size, so bounds safety depends entirely on software clamping. (per reporter;
PipelineLayoutD3D12.cppnot shown)
Root Cause Analysis
In the pre-fix ShaderModule::Compile, the HLSL registers for normal bindings are produced by GenerateBindingRemapping via a lambda that returns bgl->GetShaderRegister(index) as the binding (ShaderModuleD3D12.cpp:187-194). But when building the ignored_by_robustness list, the code recorded each non-dynamic storage buffer using its raw WGSL binding: ignored_by_robustness.emplace_back(tint::BindingPoint{.group = uint32_t(group), .binding = uint32_t(bindingInfo.binding)}); (ShaderModuleD3D12.cpp:236-237). This list is passed to Tint as ignored_by_robustness_transform (ShaderModuleD3D12.cpp:281), where it is interpreted in HLSL register space after BindingRemapper has already rewritten variable coordinates. The invariant that should hold — that the ignore list is expressed in the same coordinate space Tint uses when the robustness pass checks it — was violated: a storage buffer’s WGSL {0,0} could match a different buffer (e.g. a dynamic uniform buffer) that got remapped to HLSL register {0,0}, so robustness clamping was skipped on the wrong buffer.
The fix extracts the HLSL-mapping lambda as ToHLSLBindPoint and uses it for the ignore list too: BindingIndex bindingIndex = bgl->AsBindingIndex(bgl->GetAPIBindingIndex(bindingInfo.binding)); ignored_by_robustness.emplace_back(ToHLSLBindPoint(group, bindingIndex)); (added diff at @@ -233,8). This makes the ignore-list coordinates HLSL registers, matching what Tint checks, so clamping is disabled only on the intended storage buffer. The reporter attributes the collision to Dawn’s sorting/BindingRemapper ordering; that ordering code is not in the shown material, so those specifics are (inferred), but the coordinate-space mismatch and its fix are fully shown here.
Attack Path
- Craft a colliding layout
A bind group layout pairs a non-dynamic storage buffer whose WGSL binding and a dynamic uniform buffer whose assigned HLSL register coincide (e.g. both resolve to
{0,0}). (trigger per reporter; sorting code not shown) - Ignore list poisoned with WGSL coordinate
Compilation records the storage buffer’s WGSL
{group, binding}inignored_by_robustness(ShaderModuleD3D12.cpp:236-237). - Robustness skipped on the uniform buffer After Tint remaps the dynamic uniform buffer to the same HLSL coordinate, the robustness pass sees it in the ignore list and skips bounds clamping. (inferred; Tint pass not shown)
- Out-of-bounds read A shader indexes the dynamic uniform buffer out of bounds; with clamping disabled and the Root CBV providing no hardware bound, it reads adjacent GPU memory that is copied to an output buffer. (inferred; hardware/readback path not shown)
Impact Assessment
(inferred).Changed Functions
| Function | Change | Notes |
|---|---|---|
GenerateBindingRemappingsrc/dawn/native/d3d12/ShaderModuleD3D12.cpp |
modified | |
TEST_Fsrc/tint/lang/hlsl/writer/writer_test.cc |
modified |
Files Changed
src/dawn/native/d3d12/ShaderModuleD3D12.cppsrc/tint/lang/hlsl/writer/writer_test.cc
Audit Directions
- Cross-space coordinate listsWhen one component records identifiers (bindings, indices) that another component consumes after a remapping, ensure both sides agree on the coordinate space; reuse the same mapping function rather than hand-building raw coordinates.
- Robustness/clamping opt-out listsTreat any list that disables a safety check (bounds clamping here) as security-sensitive; verify entries cannot alias or be mis-keyed onto a different resource.
- Root descriptors without hardware boundsResources bound as D3D12 Root CBVs rely solely on software clamping, so audit every path that could skip that clamping for such bindings.
Patch
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
Regression Test / PoC
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
Original Bug Report
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.cppthird_party/dawn/src/tint/lang/hlsl/writer/raise/raise.ccthird_party/dawn/src/tint/lang/hlsl/writer/common/option_helpers.ccthird_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.
- Sorting and Register Assignment: In
BindGroupLayoutInternal::SortBindingsCompare, Dawn sorts bindings by type.BindingTypeOrder_DynamicBufferhas 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 assignedBindingIndex 0(which becomes HLSL registerb0). The storage buffer is assignedBindingIndex 1(e.g.,u1). - Poisoning the Ignore List: In
ShaderModuleD3D12.cpp, Dawn iterates over the layout to populate theignored_by_robustnesslist 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. - Tint Transformation Mismatch: In the Tint HLSL writer (
raise.cc), theBindingRemappertransform 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 registerb0). - Robustness Bypass: The
Robustnesstransform runs immediately after. When processing pointer accesses to the dynamic uniform buffer,ShouldClampchecksIsRootVarIgnored. This function reads the buffer’s current (remapped) coordinate{0, 0}and checks theignored_by_robustnesslist. Because the storage buffer previously added{0, 0}based on its WGSL index, the check returnstrue, 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.
- Setup Layout: An attacker creates a WebGPU
BindGroupLayoutwith:- A non-dynamic
Storagebuffer at@group(0) @binding(0). - A
Uniformbuffer withhasDynamicOffset: trueat a higher binding, e.g.,@group(0) @binding(7).
- A non-dynamic
- 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 theignored_by_robustnesslist. Tint remaps the uniform buffer to{0, 0}, matching the ignore list and bypassing software clamping. - OOB Read: The attacker dispatches a shader that accesses an array inside the uniform buffer using a large out-of-bounds index.
- Hardware execution: In
PipelineLayoutD3D12.cpp, dynamic uniform buffers are mapped asD3D12_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). - 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.