Chrome · Tint › Tint
CVE-2026-9918
Logic Error in Tint (Tint)
Overview
High
Severity
9.6
CVSS
No
Exploited ITW
Fixed
Fix Status
Background
- Tint / Dawn SPIR-V writer
- Tint is WebGPU’s shader compiler; for the Vulkan backend it lowers WGSL to SPIR-V. Dawn applies device-specific “workarounds” (polyfills) during this lowering to route around driver bugs.
- AMD Mesa ACO optimizer bug (crbug.com/448294721)
- A known front-end-optimizer defect in the Mesa RADV driver that memory-corrupts when folding floating-point modifiers (
fneg/fabs) into cross-lane subgroup instructions. The bug is in the third-party driver, not in Chrome. - `UnaryPolyfill` / float-negation & -abs toggles
- The mitigation: instead of emitting
OpFNegate/OpFAbs, Tint rewrites them as bit manipulation (XOR/AND of the sign/magnitude bits) so the driver optimizer never sees a foldable modifier. Driven by device toggles set for bad AMD Mesa drivers. - `shader-f16`
- A WebGPU extension exposing 16-bit floats; modern AMD hardware supports f16 and subgroup ops with the same VOP3 modifier logic as f32.
Root Cause Analysis
The mitigation was scoped to 32-bit floats only. The toggles VulkanPolyfillF32Negation/VulkanPolyfillF32Abs and the Tint workarounds polyfill_unary_f32_negation/polyfill_f32_abs/polyfill_length_scalar_f32/polyfill_distance_scalar_f32 all keyed on f32, so a WGSL shader that performed negation/abs/length/distance on f16 values bypassed the polyfill entirely and Tint emitted native OpFNegate/OpFAbs. On a vulnerable AMD Mesa driver the ACO optimizer would then fold the 16-bit modifier into a subgroup primitive and re-trigger the driver memory-corruption that the polyfill existed to prevent.
The patch generalizes the mitigation from “f32” to “float” so it covers f16 as well. In Toggles.cpp/Toggles.h the toggles are renamed VulkanPolyfillFloatNegation/VulkanPolyfillFloatAbs (“Polyfill f32 and f16 negation/abs”); PhysicalDeviceVk.cpp enables these generalized toggles for pre-25.3 AMD Mesa; and ShaderModuleVk.cpp wires them to renamed Tint workarounds polyfill_float_negation/polyfill_float_abs/polyfill_length_scalar_float/polyfill_distance_scalar_float. A new end-to-end test (PolyfillFloatUnary) exercises abs/neg/length/distance for both f32 and, when ShaderF16 is available, f16. The classification here matches the advisory’s generic “Inappropriate implementation” label: this is a LogicError — an incomplete defense-in-depth mitigation. The actual memory-safety sink (the ACO wild-deref) lives in the third-party Mesa driver and is not present in this Dawn/Tint diff (inferred/external).
Key insight
A driver-bug workaround was hard-coded to only match
f32 operands, so the identical unsafe code pattern on f16 operands slipped past the polyfill and reached the buggy AMD Mesa optimizer unchanged.Attack Path
- Malicious page
Serve a page that opens a WebGPU device requesting
shader-f16andsubgroupsfeatures (inferred: standard web reachability for WebGPU-enabled origins). - f16 subgroup shader
Submit a compute shader (
enable f16; enable subgroups;) that applies negation/abs to an f16 value inside a subgroup builtin, e.g.subgroupMin(-f16(x)). - Polyfill skipped
Tint lowers the op; because the result type is
f16, the f32-onlyUnaryPolyfillcheck does not fire and the SPIR-V writer emits rawOpFNegate/OpFAbs. - Driver compiles The un-polyfilled SPIR-V reaches Mesa RADV in the GPU process.
- Driver corruption The ACO optimizer folds the 16-bit modifier into the subgroup op and triggers the wild dereference (memory corruption) during shader compilation.
- Escalation With shaped allocations the compile-time corruption is leveraged toward GPU-process code execution / sandbox escape (inferred; the corruption and any escalation are driver-side and not in this diff).
Impact Assessment
The vulnerable code path executes in the GPU process; the triggering input is an ordinary WebGPU shader, requiring no special privilege beyond a page that can use WebGPU on affected AMD-Mesa (<25.3) hardware (gating precondition). The achieved effect is re-triggering a driver memory-corruption crash; the advisory rates it CVSS 9.6 and posits sandbox escape, but that escalation is not substantiated by this diff — the memory-unsafe sink is in the Mesa driver, and the Dawn/Tint change is purely a mitigation-coverage extension. Honestly stated: the diff fixes an incomplete workaround (LogicError); the memory-safety consequence is external and the RCE/escape is inferred.
Changed Functions
| Function | Change | Notes |
|---|---|---|
ifsrc/dawn/native/vulkan/PhysicalDeviceVk.cpp |
modified | |
PolyfillBuiltinSimpleTestssrc/dawn/tests/end2end/PolyfillBuiltinSimpleTests.cpp |
modified | |
TEST_Psrc/dawn/tests/end2end/PolyfillBuiltinSimpleTests.cpp |
modified | |
ifsrc/dawn/tests/end2end/PolyfillBuiltinSimpleTests.cpp |
modified |
Files Changed
src/dawn/native/Toggles.cppsrc/dawn/native/Toggles.hsrc/dawn/native/vulkan/PhysicalDeviceVk.cppsrc/dawn/native/vulkan/ShaderModuleVk.cppsrc/dawn/native/vulkan/TextureVk.cppsrc/dawn/tests/end2end/PolyfillBuiltinSimpleTests.cpp
Audit Directions
- Width-scoped workaroundsGrep Dawn/Tint toggles and
workaroundsfor any_f32/_f16-specific mitigation that should logically apply to all float widths (and to vector/matrix element types viaDeepestElement). - Subgroup + modifier interactionsReview other builtins that compile to modifier-foldable ops (e.g.
length,distance, dot-product chains) under subgroup/cross-lane contexts on AMD for the same folding hazard. - Driver-version gatingConfirm the pre-25.3 AMD Mesa gate matches the real fixed-driver versions for the f16 case, and that similar Mesa/Intel toggles aren’t similarly type-narrowed (e.g. the CCS-ambiguate and tanh-f16 toggles nearby).
Patch
From 29c28f75e0f35fd7e88f899e2be5807370c36fe9 Mon Sep 17 00:00:00 2001
From: Peter McNeeley <petermcneeley@google.com>
Date: Tue, 21 Apr 2026 08:37:11 -0700
Subject: [PATCH] [tint] Polyfill abs, neg, length, and distance for f16
Bug: 500099471
Change-Id: Ie9aea138776742559f0ab284b5aa79bcc04e07b8
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/302455
Reviewed-by: James Price <jrprice@google.com>
Commit-Queue: Peter McNeeley <petermcneeley@google.com>
Reviewed-by: Natalie Chouinard <chouinard@google.com>
---
diff --git a/src/dawn/native/Toggles.cpp b/src/dawn/native/Toggles.cpp
index 657a8d3..175192f 100644
--- a/src/dawn/native/Toggles.cpp
+++ b/src/dawn/native/Toggles.cpp
@@ -397,13 +397,12 @@
{Toggle::MetalPolyfillTanhF16,
{"metal_polyfill_tanh_f16", "Polyfill tanh with an f16 value for MSL.",
"https://crbug.com/42251267", ToggleStage::Device}},
- {Toggle::VulkanPolyfillF32Negation,
- {"spirv_polyfill_f32_negation",
- "Polyfill f32 negation with bit manipulation in SPIR-V writer.",
+ {Toggle::VulkanPolyfillFloatNegation,
+ {"spirv_polyfill_float_negation", "Polyfill f32 and f16 negation.",
"https://crbug.com/448294721", ToggleStage::Device}},
- {Toggle::VulkanPolyfillF32Abs,
- {"spirv_polyfill_f32_abs", "Polyfill f32 abs with bit manipulation in SPIR-V writer.",
- "https://crbug.com/448294721", ToggleStage::Device}},
+ {Toggle::VulkanPolyfillFloatAbs,
+ {"spirv_polyfill_float_abs", "Polyfill f32 and f16 abs.", "https://crbug.com/448294721",
+ ToggleStage::Device}},
{Toggle::MetalFillEmptyOcclusionQueriesWithZero,
{"metal_fill_empty_occlusion_queries_with_zero",
"Apple GPUs leave stale results in the visibility result buffer instead of writing zero if "
@@ -527,7 +526,7 @@
{Toggle::VulkanClearGen12TextureWithCCSAmbiguateOnCreation,
{"vulkan_clear_gen12_texture_with_ccs_ambiguate_on_creation",
"Clears some R8-like textures to full 0 bits as soon as they are created. This Toggle is "
- "enabled on Intel Gen12 GPUs due to a mesa driver issue.",
+ "enabled on Intel Gen12 GPUs due to a Mesa driver issue.",
"https://crbug.com/chromium/1361662", ToggleStage::Device}},
{Toggle::D3D12UseRootSignatureVersion1_1,
{"d3d12_use_root_signature_version_1_1",
diff --git a/src/dawn/native/Toggles.h b/src/dawn/native/Toggles.h
index d725887..cf32c9d 100644
--- a/src/dawn/native/Toggles.h
+++ b/src/dawn/native/Toggles.h
@@ -107,8 +107,8 @@
MetalPolyfillUnpack2x16snorm,
MetalPolyfillUnpack2x16unorm,
MetalPolyfillTanhF16,
- VulkanPolyfillF32Negation,
- VulkanPolyfillF32Abs,
+ VulkanPolyfillFloatNegation,
+ VulkanPolyfillFloatAbs,
MetalFillEmptyOcclusionQueriesWithZero,
UseBlitForBufferToDepthTextureCopy,
UseBlitForBufferToStencilTextureCopy,
diff --git a/src/dawn/native/vulkan/PhysicalDeviceVk.cpp b/src/dawn/native/vulkan/PhysicalDeviceVk.cpp
index 0df53e8..ca931a7 100644
--- a/src/dawn/native/vulkan/PhysicalDeviceVk.cpp
+++ b/src/dawn/native/vulkan/PhysicalDeviceVk.cpp
@@ -1042,14 +1042,14 @@
}
// AMD Mesa front end optimizer bug for unary negation and abs.
- // Fixed in 25.3 - See crbug.com/448294721
+ // Fixed in 25.3 - See crbug.com/448294721 and crbug.com/500099471
// See crbug.com/93692702 for variations of this bug.
if (IsAmdMesa()) {
const gpu_info::DriverVersion kGoodMesaDriver = {25, 3, 0, 0};
const bool badDriver = GetDriverVersion() < kGoodMesaDriver;
if (badDriver) {
- deviceToggles->Default(Toggle::VulkanPolyfillF32Abs, true);
- deviceToggles->Default(Toggle::VulkanPolyfillF32Negation, true);
+ deviceToggles->Default(Toggle::VulkanPolyfillFloatAbs, true);
+ deviceToggles->Default(Toggle::VulkanPolyfillFloatNegation, true);
}
}
diff --git a/src/dawn/native/vulkan/ShaderModuleVk.cpp b/src/dawn/native/vulkan/ShaderModuleVk.cpp
index 6d1b6ba..534f0d1 100644
--- a/src/dawn/native/vulkan/ShaderModuleVk.cpp
+++ b/src/dawn/native/vulkan/ShaderModuleVk.cpp
@@ -254,15 +254,15 @@
req.tintOptions.disable_workgroup_init =
GetDevice()->IsToggleEnabled(Toggle::DisableWorkgroupInit);
- req.tintOptions.workarounds.polyfill_unary_f32_negation =
- GetDevice()->IsToggleEnabled(Toggle::VulkanPolyfillF32Negation);
+ req.tintOptions.workarounds.polyfill_float_negation =
+ GetDevice()->IsToggleEnabled(Toggle::VulkanPolyfillFloatNegation);
// These polyfills all relate to incorrect backend optimization of fabs.
// See: crbug.com/93692702
- if (GetDevice()->IsToggleEnabled(Toggle::VulkanPolyfillF32Abs)) {
- req.tintOptions.workarounds.polyfill_f32_abs = true;
- req.tintOptions.workarounds.polyfill_length_scalar_f32 = true;
- req.tintOptions.workarounds.polyfill_distance_scalar_f32 = true;
+ if (GetDevice()->IsToggleEnabled(Toggle::VulkanPolyfillFloatAbs)) {
+ req.tintOptions.workarounds.polyfill_float_abs = true;
+ req.tintOptions.workarounds.polyfill_length_scalar_float = true;
+ req.tintOptions.workarounds.polyfill_distance_scalar_float = true;
}
req.tintOptions.disable_polyfill_integer_div_mod =
diff --git a/src/dawn/native/vulkan/TextureVk.cpp b/src/dawn/native/vulkan/TextureVk.cpp
index 7a629ee..a6023d4 100644
--- a/src/dawn/native/vulkan/TextureVk.cpp
+++ b/src/dawn/native/vulkan/TextureVk.cpp
@@ -1562,7 +1562,7 @@
"BindImageMemory"));
// crbug.com/1361662
- // This works around an Intel Gen12 mesa bug due to CCS ambiguates stomping on each other.
+ // This works around an Intel Gen12 Mesa bug due to CCS ambiguates stomping on each other.
// https://gitlab.freedesktop.org/mesa/mesa/-/issues/7301#note_1826367
if (device->IsToggleEnabled(Toggle::VulkanClearGen12TextureWithCCSAmbiguateOnCreation)) {
auto format = GetFormat().format;
diff --git a/src/dawn/tests/end2end/PolyfillBuiltinSimpleTests.cpp b/src/dawn/tests/end2end/PolyfillBuiltinSimpleTests.cpp
index b2e8e23..f481172 100644
--- a/src/dawn/tests/end2end/PolyfillBuiltinSimpleTests.cpp
+++ b/src/dawn/tests/end2end/PolyfillBuiltinSimpleTests.cpp
@@ -38,6 +38,14 @@
class PolyfillBuiltinSimpleTests : public DawnTest {
public:
+ std::vector<wgpu::FeatureName> GetRequiredFeatures() override {
+ std::vector<wgpu::FeatureName> features;
+ if (SupportsFeatures({wgpu::FeatureName::ShaderF16})) {
+ features.push_back(wgpu::FeatureName::ShaderF16);
+ }
+ return features;
+ }
+
wgpu::Buffer CreateBuffer(const std::vector<uint32_t>& data,
wgpu::BufferUsage usage = wgpu::BufferUsage::Storage |
wgpu::BufferUsage::CopySrc) {
@@ -45,6 +53,13 @@
return utils::CreateBufferFromData(device, data.data(), bufferSize, usage);
}
+ wgpu::Buffer CreateBuffer(const std::vector<float>& data,
+ wgpu::BufferUsage usage = wgpu::BufferUsage::Storage |
+ wgpu::BufferUsage::CopySrc) {
+ uint64_t bufferSize = static_cast<uint64_t>(data.size() * sizeof(float));
+ return utils::CreateBufferFromData(device, data.data(), bufferSize, usage);
+ }
+
wgpu::Buffer CreateBuffer(const uint32_t count,
const uint32_t default_val = 0,
wgpu::BufferUsage usage = wgpu::BufferUsage::Storage |
@@ -373,6 +388,74 @@
EXPECT_BUFFER_U32_RANGE_EQ(expected.data(), output, 0, expected.size());
}
+// Some versions of AMD Mesa (prior to 25.3) have a front-end optimizer bug where unary negation
+// and abs operations on floating point values (f32 and f16) are incorrectly optimized or
+// handled, leading to incorrect results.
+// See crbug.com/448294721 and crbug.com/500099471.
+TEST_P(PolyfillBuiltinSimpleTests, PolyfillFloatUnary) {
+ bool hasF16 = device.HasFeature(wgpu::FeatureName::ShaderF16);
+
+ std::string shader = R"(
+ @group(0) @binding(0) var<storage, read> in_f32 : array<f32, 4>;
+ @group(0) @binding(1) var<storage, read_write> out_f32 : array<f32, 8>;
+
+ @compute @workgroup_size(1)
+ fn main() {
+ out_f32[0] = abs(in_f32[0]);
+ out_f32[1] = -in_f32[1];
+ out_f32[2] = length(in_f32[2]);
+ out_f32[3] = distance(in_f32[3], 2.0);
+ )";
+
+ if (hasF16) {
+ shader = "enable f16;\n" + shader;
+ shader += R"(
+ out_f32[4] = f32(abs(f16(in_f32[0])));
+ out_f32[5] = f32(-f16(in_f32[1]));
+ out_f32[6] = f32(length(f16(in_f32[2])));
+ out_f32[7] = f32(distance(f16(in_f32[3]), 2.0h));
+ )";
+ }
+
+ shader += R"(
+ }
+ )";
+
+ wgpu::ComputePipeline pipeline = CreateComputePipeline(shader);
+
+ std::vector<float> input_data = {-1.5f, 2.0f, -2.5f, 1.0f};
+ wgpu::Buffer input = CreateBuffer(input_data, wgpu::BufferUsage::Storage);
+ wgpu::Buffer output = CreateBuffer(8, 0);
+ wgpu::BindGroup bindGroup = utils::MakeBindGroup(device, pipeline.GetBindGroupLayout(0),
+ {
+ {0, input},
+ {1, output},
Loading diff…
Regression Test / PoC
shipped with the fix
diff --git a/src/dawn/tests/end2end/PolyfillBuiltinSimpleTests.cpp b/src/dawn/tests/end2end/PolyfillBuiltinSimpleTests.cpp
index b2e8e23..f481172 100644
--- a/src/dawn/tests/end2end/PolyfillBuiltinSimpleTests.cpp
+++ b/src/dawn/tests/end2end/PolyfillBuiltinSimpleTests.cpp
@@ -38,6 +38,14 @@
class PolyfillBuiltinSimpleTests : public DawnTest {
public:
+ std::vector<wgpu::FeatureName> GetRequiredFeatures() override {
+ std::vector<wgpu::FeatureName> features;
+ if (SupportsFeatures({wgpu::FeatureName::ShaderF16})) {
+ features.push_back(wgpu::FeatureName::ShaderF16);
+ }
+ return features;
+ }
+
wgpu::Buffer CreateBuffer(const std::vector<uint32_t>& data,
wgpu::BufferUsage usage = wgpu::BufferUsage::Storage |
wgpu::BufferUsage::CopySrc) {
@@ -45,6 +53,13 @@
return utils::CreateBufferFromData(device, data.data(), bufferSize, usage);
}
+ wgpu::Buffer CreateBuffer(const std::vector<float>& data,
+ wgpu::BufferUsage usage = wgpu::BufferUsage::Storage |
+ wgpu::BufferUsage::CopySrc) {
+ uint64_t bufferSize = static_cast<uint64_t>(data.size() * sizeof(float));
+ return utils::CreateBufferFromData(device, data.data(), bufferSize, usage);
+ }
+
wgpu::Buffer CreateBuffer(const uint32_t count,
const uint32_t default_val = 0,
wgpu::BufferUsage usage = wgpu::BufferUsage::Storage |
@@ -373,6 +388,74 @@
EXPECT_BUFFER_U32_RANGE_EQ(expected.data(), output, 0, expected.size());
}
+// Some versions of AMD Mesa (prior to 25.3) have a front-end optimizer bug where unary negation
+// and abs operations on floating point values (f32 and f16) are incorrectly optimized or
+// handled, leading to incorrect results.
+// See crbug.com/448294721 and crbug.com/500099471.
+TEST_P(PolyfillBuiltinSimpleTests, PolyfillFloatUnary) {
+ bool hasF16 = device.HasFeature(wgpu::FeatureName::ShaderF16);
+
+ std::string shader = R"(
+ @group(0) @binding(0) var<storage, read> in_f32 : array<f32, 4>;
+ @group(0) @binding(1) var<storage, read_write> out_f32 : array<f32, 8>;
+
+ @compute @workgroup_size(1)
+ fn main() {
+ out_f32[0] = abs(in_f32[0]);
+ out_f32[1] = -in_f32[1];
+ out_f32[2] = length(in_f32[2]);
+ out_f32[3] = distance(in_f32[3], 2.0);
+ )";
+
+ if (hasF16) {
+ shader = "enable f16;\n" + shader;
+ shader += R"(
+ out_f32[4] = f32(abs(f16(in_f32[0])));
+ out_f32[5] = f32(-f16(in_f32[1]));
+ out_f32[6] = f32(length(f16(in_f32[2])));
+ out_f32[7] = f32(distance(f16(in_f32[3]), 2.0h));
+ )";
+ }
+
+ shader += R"(
+ }
+ )";
+
+ wgpu::ComputePipeline pipeline = CreateComputePipeline(shader);
+
+ std::vector<float> input_data = {-1.5f, 2.0f, -2.5f, 1.0f};
+ wgpu::Buffer input = CreateBuffer(input_data, wgpu::BufferUsage::Storage);
+ wgpu::Buffer output = CreateBuffer(8, 0);
+ wgpu::BindGroup bindGroup = utils::MakeBindGroup(device, pipeline.GetBindGroupLayout(0),
+ {
+ {0, input},
+ {1, output},
+ });
+
+ wgpu::CommandBuffer commands;
+ {
+ wgpu::CommandEncoder encoder = device.CreateCommandEncoder();
+ wgpu::ComputePassEncoder pass = encoder.BeginComputePass();
+ pass.SetPipeline(pipeline);
+ pass.SetBindGroup(0, bindGroup);
+ pass.DispatchWorkgroups(1);
+ pass.End();
+ commands = encoder.Finish();
+ }
+
+ queue.Submit(1, &commands);
+
+ std::vector<float> expected = {1.5f, -2.0f, 2.5f, 1.0f, 0.0f, 0.0f, 0.0f, 0.0f};
+ if (hasF16) {
+ expected[4] = 1.5f;
+ expected[5] = -2.0f;
+ expected[6] = 2.5f;
+ expected[7] = 1.0f;
+ }
+
+ EXPECT_BUFFER_FLOAT_RANGE_EQ(expected.data(), output, 0, expected.size());
+}
+
DAWN_INSTANTIATE_TEST(PolyfillBuiltinSimpleTests,
D3D12Backend(),
D3D11Backend(),
@@ -383,6 +466,7 @@
MetalBackend({"scalarize_max_min_clamp"}),
VulkanBackend({"scalarize_max_min_clamp"}),
VulkanBackend({"vulkan_polyfill_switch_with_if"}),
+ VulkanBackend({"spirv_polyfill_float_negation", "spirv_polyfill_float_abs"}),
D3D11Backend({"scalarize_max_min_clamp"}),
OpenGLESBackend());
diff --git a/src/tint/lang/core/ir/transform/builtin_polyfill_test.cc b/src/tint/lang/core/ir/transform/builtin_polyfill_test.cc
index 955cceb..909532e 100644
--- a/src/tint/lang/core/ir/transform/builtin_polyfill_test.cc
+++ b/src/tint/lang/core/ir/transform/builtin_polyfill_test.cc
@@ -1245,6 +1245,80 @@
EXPECT_EQ(expect, str());
}
+TEST_F(IR_BuiltinPolyfillTest, Distance_Scalar_F16) {
+ auto* arg1 = b.FunctionParam("arg1", ty.f16());
+ auto* arg2 = b.FunctionParam("arg2", ty.f16());
+ auto* func = b.Function("foo", ty.f16());
+ func->SetParams({arg1, arg2});
+
+ b.Append(func->Block(), [&] {
+ auto* result = b.Call(ty.f16(), core::BuiltinFn::kDistance, arg1, arg2);
+ b.Return(func, result);
+ });
+
+ auto* src = R"(
+%foo = func(%arg1:f16, %arg2:f16):f16 {
+ $B1: {
+ %4:f16 = distance %arg1, %arg2
+ ret %4
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+%foo = func(%arg1:f16, %arg2:f16):f16 {
+ $B1: {
+ %4:f16 = sub %arg1, %arg2
+ %5:f16 = abs %4
+ ret %5
+ }
+}
+)";
+
+ BuiltinPolyfillConfig config;
+ config.distance_scalar_float = true;
+ Run(BuiltinPolyfill, config);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(IR_BuiltinPolyfillTest, Length_Scalar_F16) {
+ auto* arg = b.FunctionParam("arg", ty.f16());
+ auto* func = b.Function("foo", ty.f16());
+ func->SetParams({arg});
+
+ b.Append(func->Block(), [&] {
+ auto* result = b.Call(ty.f16(), core::BuiltinFn::kLength, arg);
+ b.Return(func, result);
+ });
+
+ auto* src = R"(
+%foo = func(%arg:f16):f16 {
+ $B1: {
+ %3:f16 = length %arg
+ ret %3
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+%foo = func(%arg:f16):f16 {
+ $B1: {
+ %3:f16 = abs %arg
+ ret %3
+ }
+}
+)";
+
+ BuiltinPolyfillConfig config;
+ config.length_scalar_float = true;
+ Run(BuiltinPolyfill, config);
+
+ EXPECT_EQ(expect, str());
+}
+
TEST_F(IR_BuiltinPolyfillTest, Distance_Scalar_F32) {
Build(core::BuiltinFn::kDistance, ty.f32(), Vector{ty.f32(), ty.f32()});
auto* src = R"(
@@ -1268,7 +1342,7 @@
EXPECT_EQ(src, str());
BuiltinPolyfillConfig config;
- config.distance_scalar_f32 = true;
+ config.distance_scalar_float = true;
Run(BuiltinPolyfill, config);
EXPECT_EQ(expect, str());
}
@@ -1288,7 +1362,7 @@
EXPECT_EQ(src, str());
BuiltinPolyfillConfig config;
- config.distance_scalar_f32 = true;
+ config.distance_scalar_float = true;
Run(BuiltinPolyfill, config);
EXPECT_EQ(expect, str());
}
@@ -1315,7 +1389,7 @@
EXPECT_EQ(src, str());
BuiltinPolyfillConfig config;
- config.length_scalar_f32 = true;
+ config.length_scalar_float = true;
Run(BuiltinPolyfill, config);
EXPECT_EQ(expect, str());
}
@@ -1335,7 +1409,7 @@
EXPECT_EQ(src, str());
BuiltinPolyfillConfig config;
- config.length_scalar_f32 = true;
+ config.length_scalar_float = true;
Run(BuiltinPolyfill, config);
EXPECT_EQ(expect, str());
}
diff --git a/src/tint/lang/spirv/writer/raise/unary_polyfill_test.cc b/src/tint/lang/spirv/writer/raise/unary_polyfill_test.cc
index 398146f..7def679 100644
--- a/src/tint/lang/spirv/writer/raise/unary_polyfill_test.cc
+++ b/src/tint/lang/spirv/writer/raise/unary_polyfill_test.cc
@@ -69,7 +69,47 @@
)";
UnaryPolyfillConfig config;
- config.polyfill_f32_negation = true;
+ config.polyfill_float_negation = true;
+ Run(UnaryPolyfill, config);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvWriter_UnaryPolyfillTest, Negation_F16) {
+ auto* arg = b.FunctionParam("arg", ty.f16());
+ auto* func = b.Function("foo", ty.f16());
+ func->SetParams({arg});
+
+ b.Append(func->Block(), [&] {
+ auto* result = b.Negation(arg);
+ b.Return(func, result);
+ });
+
+ auto* src = R"(
+%foo = func(%arg:f16):f16 {
+ $B1: {
+ %3:f16 = negation %arg
+ ret %3
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+%foo = func(%arg:f16):f16 {
+ $B1: {
+ %3:f32 = convert %arg
+ %4:u32 = bitcast<u32> %3
+ %5:u32 = xor %4, 2147483648u
+ %6:f32 = bitcast<f32> %5
+ %7:f16 = convert %6
+ ret %7
+ }
+}
+)";
+
+ UnaryPolyfillConfig config;
+ config.polyfill_float_negation = true;
Run(UnaryPolyfill, config);
EXPECT_EQ(expect, str());
@@ -107,7 +147,7 @@
)";
UnaryPolyfillConfig config;
- config.polyfill_f32_negation = true;
+ config.polyfill_float_negation = true;
Run(UnaryPolyfill, config);
EXPECT_EQ(expect, str());
@@ -145,7 +185,47 @@
)";
... (truncated)
Loading diff…
Original Bug Report
reported by rj...@google.com
Potential GPU Sandbox Escape via Bypassing AMD Mesa Bug Mitigation with f16 in Tint
Flapjack, 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 mitigation for an AMD Mesa driver bug (crbug.com/448294721) in Dawn/Tint incorrectly only covers 32-bit floats. By using 16-bit floats (f16) in WebGPU shaders, an attacker can bypass this polyfill. This allows re-triggering a memory corruption crash in the GPU process, potentially leading to a sandbox escape.
Affected files:
third_party/dawn/src/tint/lang/spirv/writer/raise/unary_polyfill.ccthird_party/dawn/src/tint/lang/spirv/writer/raise/unary_polyfill.h
Estimated timestamp from git blame: 2025-12-01
Background
The AMD Mesa ACO optimizer contains a known vulnerability (tracked as crbug.com/448294721) that results in memory corruption (a ‘wild-deref’ in aco::combine_instruction) when attempting to fold floating-point modifiers like negation (fneg) or absolute value (fabs) into cross-lane subgroup instructions.
Dawn/Tint implemented a mitigation for this issue via UnaryPolyfill, which replaces these operations with bitwise manipulation (XOR 0x80000000 for negation and AND 0x7FFFFFFF for absolute value) before emitting SPIR-V for Vulkan. This prevents the driver’s optimizer from folding the modifiers into subgroup primitives, avoiding the crash.
Vulnerability Details
The implementation of this mitigation in third_party/dawn/src/tint/lang/spirv/writer/raise/unary_polyfill.cc is incomplete because it explicitly checks for 32-bit floats (f32) while ignoring 16-bit floats (f16).
The current implementation explicitly checks for f32 (lines 53-54 and 58-59):
if (config.polyfill_f32_negation && unary->Op() == core::UnaryOp::kNegation &&
unary->Result()->Type()->DeepestElement()->Is<core::type::F32>()) {
// ...
}
// ...
if (config.polyfill_f32_abs && builtin->Func() == core::BuiltinFn::kAbs &&
builtin->Result()->Type()->DeepestElement()->Is<core::type::F32>()) {
// ...
}
WebGPU supports the shader-f16 extension, and modern AMD GPUs support both f16 types and subgroup operations (verified in DeviceVk.cpp and PhysicalDeviceVk.cpp). AMD hardware uses the same VOP3 instruction encodings and modifier logic for both 32-bit and 16-bit operands. Consequently, the same Mesa driver vulnerability can be triggered using 16-bit modifiers.
If an attacker provides a WebGPU shader using f16 types, Tint will bypass the polyfill and emit native OpFNegate or OpFAbs instructions. The Mesa driver will then attempt to fold these 16-bit modifiers into the subgroup primitive, re-triggering the memory corruption in the GPU process during compilation.
Potential Attack Scenario
Note: These are potential steps, as our tooling agent does not have the ability to run code to verify the exploit end-to-end.
- An attacker crafts a malicious webpage that sets up a WebGPU context, requesting a device with both the
shader-f16andsubgroupsfeatures enabled. - The attacker provides a compute shader that explicitly enables these features:
enable f16; enable subgroups;. - The shader contains a subgroup built-in function, passing a negated 16-bit float value, e.g.,
let val = f16(1.0); let result = subgroupMin(-val);. - Tint lowers this to IR. Because the negation is on an
f16type, theUnaryPolyfillpass skips it. - The SPIR-V printer emits a raw
OpFNegateinstruction. - The un-polyfilled SPIR-V is passed to the Vulkan driver (Mesa RADV) in the GPU process.
- The buggy ACO optimizer in Mesa attempts to fold the 16-bit
fneginto the subgroup operation, triggering the wild dereference. - With careful control over shader complexity and memory allocations, an attacker could potentially turn this compilation-time memory corruption into arbitrary code execution, achieving a renderer-to-GPU sandbox escape.
Suggested Fix
The UnaryPolyfill should be extended to handle f16 types when the mitigation is active.
- The checks in
Process()should be updated to allowf16types (e.g., usingIsAnyOf<core::type::F32, core::type::F16>()). - The polyfill logic in
PolyfillF32NegationandPolyfillF32Absshould be renamed and updated to use 16-bit bitwise masks (0x8000_ufor negation and0x7FFF_ufor absolute value) when processingf16operands. Alternatively, separate functions could be added forf16. - Ensure
ty.MatchWidthandb.Bitcastcorrectly handle mappingf16elements tou16before applying the masks.
Evaluated with Chrome root at commit: 09ec9e7cc4d24823d20b6d37cf3d282734f6bf0f
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.
References
On This Page