[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}, + }); + + 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.cc b/src/tint/lang/core/ir/transform/builtin_polyfill.cc index 6284528..ce5a3e4 100644 --- a/src/tint/lang/core/ir/transform/builtin_polyfill.cc +++ b/src/tint/lang/core/ir/transform/builtin_polyfill.cc
@@ -101,8 +101,8 @@ } break; case core::BuiltinFn::kDistance: - if (config.distance_scalar_f32 && - builtin->Args()[0]->Type()->Is<core::type::F32>()) { + if (config.distance_scalar_float && + builtin->Args()[0]->Type()->IsFloatScalar()) { worklist.Push(builtin); } break; @@ -135,8 +135,8 @@ } break; case core::BuiltinFn::kLength: - if (config.length_scalar_f32 && - builtin->Args()[0]->Type()->Is<core::type::F32>()) { + if (config.length_scalar_float && + builtin->Args()[0]->Type()->IsFloatScalar()) { worklist.Push(builtin); } break; @@ -237,7 +237,7 @@ Degrees(builtin); break; case core::BuiltinFn::kDistance: - DistanceScalarF32(builtin); + DistanceScalarFloat(builtin); break; case core::BuiltinFn::kSmoothstep: SmoothStep(builtin); @@ -258,7 +258,7 @@ InsertBits(builtin); break; case core::BuiltinFn::kLength: - LengthScalarF32(builtin); + LengthScalarFloat(builtin); break; case core::BuiltinFn::kRadians: Radians(builtin); @@ -577,9 +577,9 @@ call->Destroy(); } - /// Polyfill a `distance()` builtin call for scalar f32. + /// Polyfill a `distance()` builtin call for scalar f32 and f16. /// @param call the builtin call instruction - void DistanceScalarF32(ir::CoreBuiltinCall* call) { + void DistanceScalarFloat(ir::CoreBuiltinCall* call) { // distance(x, y) -> abs(x - y) b.InsertBefore(call, [&] { auto* sub = b.Subtract(call->Args()[0], call->Args()[1]); @@ -892,9 +892,9 @@ } } - /// Polyfill a `length()` builtin call for scalar f32. + /// Polyfill a `length()` builtin call for scalar f32 and f16. /// @param call the builtin call instruction - void LengthScalarF32(ir::CoreBuiltinCall* call) { + void LengthScalarFloat(ir::CoreBuiltinCall* call) { // length(x) -> abs(x) b.InsertBefore(call, [&] { b.CallWithResult(call->DetachResult(), core::BuiltinFn::kAbs, call->Args()[0]); @@ -966,9 +966,9 @@ one = b.MatchWidth(1_h, type); } - // Intel mesa incorrectly performs saturate on vec f16 loads from uniforms. + // Intel Mesa incorrectly performs saturate on vec f16 loads from uniforms. // Note: to avoid compiler pattern matching, we do the min then the max which is - // functionally different than doing the max then the min for high/low swapped (this doesnt + // functionally different than doing the max then the min for high/low swapped (this doesn't // matter in the case with saturate). See crbug.com/448873316 if (config.saturate_as_min_max && is_vec_f16) { b.InsertBefore(call, [&] {
diff --git a/src/tint/lang/core/ir/transform/builtin_polyfill.h b/src/tint/lang/core/ir/transform/builtin_polyfill.h index a04f5f1..6d4c81c 100644 --- a/src/tint/lang/core/ir/transform/builtin_polyfill.h +++ b/src/tint/lang/core/ir/transform/builtin_polyfill.h
@@ -71,8 +71,8 @@ bool count_trailing_zeros = false; /// Should `degrees()` be polyfilled? bool degrees = false; - /// Should `distance()` be polyfilled for scalar f32? - bool distance_scalar_f32 = false; + /// Should `distance()` be polyfilled for scalar f32 and f16? + bool distance_scalar_float = false; /// How should `extractBits()` be polyfilled? BuiltinPolyfillLevel extract_bits = BuiltinPolyfillLevel::kNone; /// Should `firstLeadingBit()` be polyfilled? @@ -83,8 +83,8 @@ bool fwidth_fine = false; /// How should `insertBits()` be polyfilled? BuiltinPolyfillLevel insert_bits = BuiltinPolyfillLevel::kNone; - /// Should `length()` be polyfilled for scalar f32? - bool length_scalar_f32 = false; + /// Should `length()` be polyfilled for scalar f32 and f16? + bool length_scalar_float = false; /// Should `radians()` be polyfilled? bool radians = false; /// Should `reflect()` be polyfilled for vec2<f32>? @@ -105,22 +105,24 @@ bool pack_unpack_4x8_norm = false; /// Should `subgroupBroadcast(f16)` be polyfilled? bool subgroup_broadcast_f16 = false; - // Should 'saturate(f16)' be polyfilled with min and max. + /// Should 'saturate(f16)' be polyfilled with min and max. bool saturate_as_min_max = false; /// Reflection for this class TINT_REFLECT(BuiltinPolyfillConfig, clamp_int, + clamp_float, + abs_signed_int, count_leading_zeros, count_trailing_zeros, degrees, - distance_scalar_f32, + distance_scalar_float, extract_bits, first_leading_bit, first_trailing_bit, fwidth_fine, insert_bits, - length_scalar_f32, + length_scalar_float, radians, reflect_vec2_f32, saturate,
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/common/options.h b/src/tint/lang/spirv/writer/common/options.h index 9c52337..077559b 100644 --- a/src/tint/lang/spirv/writer/common/options.h +++ b/src/tint/lang/spirv/writer/common/options.h
@@ -105,17 +105,17 @@ /// Set to `true` to always pass matrices to user functions by pointer instead of by value. bool pass_matrix_by_pointer = false; - /// Set to `true` to generate polyfill for f32 negation. - bool polyfill_unary_f32_negation = false; + /// Set to `true` to generate polyfill for f32 and f16 negation. + bool polyfill_float_negation = false; - /// Set to `true` to generate polyfill for f32 abs. - bool polyfill_f32_abs = false; + /// Set to `true` to generate polyfill for f32 and f16 abs. + bool polyfill_float_abs = false; - /// Set to `true` to generate polyfill for length(scalar f32). - bool polyfill_length_scalar_f32 = false; + /// Set to `true` to generate polyfill for length(scalar f32 and f16). + bool polyfill_length_scalar_float = false; - /// Set to `true` to generate polyfill for distance(scalar f32). - bool polyfill_distance_scalar_f32 = false; + /// Set to `true` to generate polyfill for distance(scalar f32 and f16). + bool polyfill_distance_scalar_float = false; /// Set to `true` to generate polyfill for f16 saturate. bool polyfill_saturate_as_min_max_f16 = false; @@ -134,10 +134,10 @@ texture_sample_compare_2d_polyfill, polyfill_subgroup_broadcast_f16, pass_matrix_by_pointer, - polyfill_unary_f32_negation, - polyfill_f32_abs, - polyfill_length_scalar_f32, - polyfill_distance_scalar_f32, + polyfill_float_negation, + polyfill_float_abs, + polyfill_length_scalar_float, + polyfill_distance_scalar_float, polyfill_saturate_as_min_max_f16, cooperative_matrix_stride_is_matrix_elements); };
diff --git a/src/tint/lang/spirv/writer/raise/raise.cc b/src/tint/lang/spirv/writer/raise/raise.cc index 2eaf56e..a70ac65 100644 --- a/src/tint/lang/spirv/writer/raise/raise.cc +++ b/src/tint/lang/spirv/writer/raise/raise.cc
@@ -130,8 +130,8 @@ core_polyfills.pack_4xu8_clamp = true; core_polyfills.pack_unpack_4x8_norm = options.workarounds.polyfill_pack_unpack_4x8_norm; core_polyfills.abs_signed_int = true; - core_polyfills.length_scalar_f32 = options.workarounds.polyfill_length_scalar_f32; - core_polyfills.distance_scalar_f32 = options.workarounds.polyfill_distance_scalar_f32; + core_polyfills.length_scalar_float = options.workarounds.polyfill_length_scalar_float; + core_polyfills.distance_scalar_float = options.workarounds.polyfill_distance_scalar_float; core_polyfills.subgroup_broadcast_f16 = options.workarounds.polyfill_subgroup_broadcast_f16; core_polyfills.saturate_as_min_max = options.workarounds.polyfill_saturate_as_min_max_f16; TINT_CHECK_RESULT(core::ir::transform::BuiltinPolyfill(module, core_polyfills)); @@ -224,11 +224,11 @@ .signed_negation = true, .signed_arithmetic = true, .signed_shiftleft = true}; TINT_CHECK_RESULT(core::ir::transform::SignedIntegerPolyfill(module, signed_integer_cfg)); - // AMD mesa front end optimizer bug for unary negation and abs. - // Fixed in 25.3 - See crbug.com/448294721 + // AMD Mesa front end optimizer bug for unary f32 and f16 negation and abs. + // Fixed in 25.3 - See crbug.com/448294721 and crbug.com/500099471 raise::UnaryPolyfillConfig unary_polyfill_cfg = { - .polyfill_f32_negation = options.workarounds.polyfill_unary_f32_negation, - .polyfill_f32_abs = options.workarounds.polyfill_f32_abs}; + .polyfill_float_negation = options.workarounds.polyfill_float_negation, + .polyfill_float_abs = options.workarounds.polyfill_float_abs}; TINT_CHECK_RESULT(raise::UnaryPolyfill(module, unary_polyfill_cfg));
diff --git a/src/tint/lang/spirv/writer/raise/unary_polyfill.cc b/src/tint/lang/spirv/writer/raise/unary_polyfill.cc index d9750d0..f4c623a 100644 --- a/src/tint/lang/spirv/writer/raise/unary_polyfill.cc +++ b/src/tint/lang/spirv/writer/raise/unary_polyfill.cc
@@ -30,6 +30,7 @@ #include "src/tint/lang/core/ir/builder.h" #include "src/tint/lang/core/ir/module.h" #include "src/tint/lang/core/ir/validator.h" +#include "src/tint/utils/ice/ice.h" #include "src/tint/utils/result.h" using namespace tint::core::fluent_types; // NOLINT @@ -46,36 +47,89 @@ core::type::Manager& ty{ir.Types()}; void Process() { - Vector<core::ir::Unary*, 4> unary_worklist; - Vector<core::ir::CoreBuiltinCall*, 4> builtin_worklist; + // Find all f16 instructions and promote them to f32. + Vector<core::ir::Unary*, 4> f16_unary_worklist; + Vector<core::ir::CoreBuiltinCall*, 4> f16_builtin_worklist; for (auto* inst : ir.Instructions()) { if (auto* unary = inst->As<core::ir::Unary>()) { - if (config.polyfill_f32_negation && unary->Op() == core::UnaryOp::kNegation && - unary->Result()->Type()->DeepestElement()->Is<core::type::F32>()) { - unary_worklist.Push(unary); + if (config.polyfill_float_negation && unary->Op() == core::UnaryOp::kNegation && + unary->Result()->Type()->DeepestElement()->Is<core::type::F16>()) { + f16_unary_worklist.Push(unary); } } else if (auto* builtin = inst->As<core::ir::CoreBuiltinCall>()) { - if (config.polyfill_f32_abs && builtin->Func() == core::BuiltinFn::kAbs && - builtin->Result()->Type()->DeepestElement()->Is<core::type::F32>()) { - builtin_worklist.Push(builtin); + if (config.polyfill_float_abs && builtin->Func() == core::BuiltinFn::kAbs && + builtin->Result()->Type()->DeepestElement()->Is<core::type::F16>()) { + f16_builtin_worklist.Push(builtin); } } } - for (auto* unary : unary_worklist) { + for (auto* unary : f16_unary_worklist) { + TINT_ASSERT(unary->Op() == core::UnaryOp::kNegation); + PromoteToF32(unary); + } + for (auto* builtin : f16_builtin_worklist) { + TINT_ASSERT(builtin->Func() == core::BuiltinFn::kAbs); + PromoteToF32(builtin); + } + + // Find all f32 instructions (including newly promoted ones) and polyfill them. + Vector<core::ir::Unary*, 4> f32_unary_worklist; + Vector<core::ir::CoreBuiltinCall*, 4> f32_builtin_worklist; + for (auto* inst : ir.Instructions()) { + if (auto* unary = inst->As<core::ir::Unary>()) { + if (config.polyfill_float_negation && unary->Op() == core::UnaryOp::kNegation && + unary->Result()->Type()->DeepestElement()->Is<core::type::F32>()) { + f32_unary_worklist.Push(unary); + } + } else if (auto* builtin = inst->As<core::ir::CoreBuiltinCall>()) { + if (config.polyfill_float_abs && builtin->Func() == core::BuiltinFn::kAbs && + builtin->Result()->Type()->DeepestElement()->Is<core::type::F32>()) { + f32_builtin_worklist.Push(builtin); + } + } + } + + for (auto* unary : f32_unary_worklist) { + TINT_ASSERT(unary->Op() == core::UnaryOp::kNegation); PolyfillF32Negation(unary); } - for (auto* builtin : builtin_worklist) { + for (auto* builtin : f32_builtin_worklist) { + TINT_ASSERT(builtin->Func() == core::BuiltinFn::kAbs); PolyfillF32Abs(builtin); } } + void PromoteToF32(core::ir::Unary* unary) { + auto* val = unary->Val(); + auto* type = val->Type(); + b.InsertBefore(unary, [&] { + auto* f32_ty = ty.MatchWidth(ty.f32(), type); + auto* f32_val = b.Convert(f32_ty, val); + auto* f32_neg = b.Negation(f32_val); + b.ConvertWithResult(unary->DetachResult(), f32_neg->Result()); + }); + unary->Destroy(); + } + + void PromoteToF32(core::ir::CoreBuiltinCall* builtin) { + auto* val = builtin->Args()[0]; + auto* type = val->Type(); + b.InsertBefore(builtin, [&] { + auto* f32_ty = ty.MatchWidth(ty.f32(), type); + auto* f32_val = b.Convert(f32_ty, val); + auto* f32_abs = b.Call(f32_ty, core::BuiltinFn::kAbs, f32_val); + b.ConvertWithResult(builtin->DetachResult(), f32_abs->Result()); + }); + builtin->Destroy(); + } + void PolyfillF32Negation(core::ir::Unary* unary) { auto* val = unary->Val(); auto* type = val->Type(); - // AMD mesa front end optimizer bug for unary negation and abs. - // Fixed in 25.3 - See crbug.com/448294721 + // AMD Mesa front end optimizer bug for unary f32 (and indirectly f16) negation and abs. + // Fixed in 25.3 - See crbug.com/448294721 and crbug.com/500099471 // Note we use bitcast as a hammer to avoid the optimizer seeing through other possible // workarounds. b.InsertBefore(unary, [&] { @@ -92,8 +146,8 @@ auto* val = builtin->Args()[0]; auto* type = val->Type(); - // AMD mesa front end optimizer bug for unary negation and abs. - // Fixed in 25.3 - See crbug.com/448294721 + // AMD Mesa front end optimizer bug for unary f32 (and indirectly f16) negation and abs. + // Fixed in 25.3 - See crbug.com/500099471 // Note we use bitcast as a hammer to avoid the optimizer seeing through other possible // workarounds. b.InsertBefore(builtin, [&] {
diff --git a/src/tint/lang/spirv/writer/raise/unary_polyfill.h b/src/tint/lang/spirv/writer/raise/unary_polyfill.h index 36bf2f6..9380c51 100644 --- a/src/tint/lang/spirv/writer/raise/unary_polyfill.h +++ b/src/tint/lang/spirv/writer/raise/unary_polyfill.h
@@ -48,10 +48,10 @@ /// Configuration for the UnaryPolyfill transform. struct UnaryPolyfillConfig { - /// If true, polyfill f32 negation with bit manipulation. - bool polyfill_f32_negation = false; - /// If true, polyfill f32 abs with sign() * x. - bool polyfill_f32_abs = false; + /// If true, polyfill f32 and f16 negation with bit manipulation. + bool polyfill_float_negation = false; + /// If true, polyfill f32 and f16 abs. + bool polyfill_float_abs = false; }; /// UnaryPolyfill is a transform that replaces unary instructions with polyfills.
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 @@ )"; UnaryPolyfillConfig config; - config.polyfill_f32_abs = true; + config.polyfill_float_abs = true; + Run(UnaryPolyfill, config); + + EXPECT_EQ(expect, str()); +} + +TEST_F(SpirvWriter_UnaryPolyfillTest, Abs_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::kAbs, arg); + b.Return(func, result); + }); + + auto* src = R"( +%foo = func(%arg:f16):f16 { + $B1: { + %3:f16 = abs %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 = and %4, 2147483647u + %6:f32 = bitcast<f32> %5 + %7:f16 = convert %6 + ret %7 + } +} +)"; + + UnaryPolyfillConfig config; + config.polyfill_float_abs = true; Run(UnaryPolyfill, config); EXPECT_EQ(expect, str()); @@ -183,7 +263,7 @@ )"; UnaryPolyfillConfig config; - config.polyfill_f32_abs = true; + config.polyfill_float_abs = true; Run(UnaryPolyfill, config); EXPECT_EQ(expect, str()); @@ -210,7 +290,7 @@ EXPECT_EQ(src, str()); UnaryPolyfillConfig config; - config.polyfill_f32_negation = false; + config.polyfill_float_negation = false; Run(UnaryPolyfill, config); EXPECT_EQ(src, str()); @@ -237,7 +317,7 @@ EXPECT_EQ(src, str()); UnaryPolyfillConfig config; - config.polyfill_f32_abs = false; + config.polyfill_float_abs = false; Run(UnaryPolyfill, config); EXPECT_EQ(src, str());
diff --git a/src/tint/lang/spirv/writer/writer_fuzz.cc b/src/tint/lang/spirv/writer/writer_fuzz.cc index 64ce0f7..e8ef81a 100644 --- a/src/tint/lang/spirv/writer/writer_fuzz.cc +++ b/src/tint/lang/spirv/writer/writer_fuzz.cc
@@ -67,8 +67,8 @@ bool subgroup_shuffle_clamped; bool polyfill_subgroup_broadcast_f16; bool pass_matrix_by_pointer; - bool polyfill_unary_f32_negation; - bool polyfill_f32_abs; + bool polyfill_float_negation; + bool polyfill_float_abs; bool use_demote_to_helper_invocation; bool use_storage_input_output_16; bool use_zero_initialize_workgroup_memory; @@ -84,8 +84,8 @@ bool polyfill_saturate_as_min_max_f16; bool multisampled_framebuffer_fetch; bool cooperative_matrix_stride_is_matrix_elements; - bool polyfill_length_scalar_f32; - bool polyfill_distance_scalar_f32; + bool polyfill_length_scalar_float; + bool polyfill_distance_scalar_float; /// Reflect the fields of this class so that it can be used by tint::ForeachField() TINT_REFLECT(FuzzedOptions, @@ -103,8 +103,8 @@ subgroup_shuffle_clamped, polyfill_subgroup_broadcast_f16, pass_matrix_by_pointer, - polyfill_unary_f32_negation, - polyfill_f32_abs, + polyfill_float_negation, + polyfill_float_abs, use_demote_to_helper_invocation, use_storage_input_output_16, use_zero_initialize_workgroup_memory, @@ -120,8 +120,8 @@ polyfill_saturate_as_min_max_f16, multisampled_framebuffer_fetch, cooperative_matrix_stride_is_matrix_elements, - polyfill_length_scalar_f32, - polyfill_distance_scalar_f32); + polyfill_length_scalar_float, + polyfill_distance_scalar_float); TINT_REFLECT_HASH_CODE(FuzzedOptions); }; @@ -314,8 +314,8 @@ options.workarounds.polyfill_subgroup_broadcast_f16 = fuzzed_options.polyfill_subgroup_broadcast_f16; options.workarounds.pass_matrix_by_pointer = fuzzed_options.pass_matrix_by_pointer; - options.workarounds.polyfill_unary_f32_negation = fuzzed_options.polyfill_unary_f32_negation; - options.workarounds.polyfill_f32_abs = fuzzed_options.polyfill_f32_abs; + options.workarounds.polyfill_float_negation = fuzzed_options.polyfill_float_negation; + options.workarounds.polyfill_float_abs = fuzzed_options.polyfill_float_abs; options.extensions.use_demote_to_helper_invocation = fuzzed_options.use_demote_to_helper_invocation; options.extensions.use_storage_input_output_16 = fuzzed_options.use_storage_input_output_16; @@ -335,8 +335,9 @@ fuzzed_options.texture_sample_compare_2d_polyfill; options.workarounds.polyfill_saturate_as_min_max_f16 = fuzzed_options.polyfill_saturate_as_min_max_f16; - options.workarounds.polyfill_length_scalar_f32 = fuzzed_options.polyfill_length_scalar_f32; - options.workarounds.polyfill_distance_scalar_f32 = fuzzed_options.polyfill_distance_scalar_f32; + options.workarounds.polyfill_length_scalar_float = fuzzed_options.polyfill_length_scalar_float; + options.workarounds.polyfill_distance_scalar_float = + fuzzed_options.polyfill_distance_scalar_float; options.workarounds.cooperative_matrix_stride_is_matrix_elements = fuzzed_options.cooperative_matrix_stride_is_matrix_elements; options.multisampled_framebuffer_fetch = fuzzed_options.multisampled_framebuffer_fetch;