From 29c28f75e0f35fd7e88f899e2be5807370c36fe9 Mon Sep 17 00:00:00 2001
From: Peter McNeeley <[email protected]>
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 <[email protected]>
Commit-Queue: Peter McNeeley <[email protected]>
Reviewed-by: Natalie Chouinard <[email protected]>
---

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

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)