[spirv] Add workaround for Xclipse bug A certain sequence of bitwise operations followed by two select builtins produces incorrect results on Xclipse GPUs. Guidance from Samsung is that we can avoid the broken optimization by emitting `x < 1u` instead of `x == 0u`. Add a new transform that does this, and add a Dawn E2E test to test it. Guard the workaround on driver version < 25.x, since the issue is fixed in newer drivers. Fixed: 542268656 Change-Id: Ibfe7557a81e9483bcf7be64477e69318b9bc6658 Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/330255 Reviewed-by: dan sinclair <dsinclair@chromium.org> Reviewed-by: Brandon Jones <bajones@chromium.org> Commit-Queue: James Price <jrprice@google.com>
diff --git a/src/dawn/native/Toggles.cpp b/src/dawn/native/Toggles.cpp index f477a1e..46a8107 100644 --- a/src/dawn/native/Toggles.cpp +++ b/src/dawn/native/Toggles.cpp
@@ -828,6 +828,11 @@ "Reject NPOT depth/stencil textures with mipLevelCount > 1. Workaround for mip level " "miscomputation in PowerVR proprietary driver.", "https://crbug.com/540087398", ToggleStage::Device}}, + {Toggle::VulkanReplaceUnsignedCompareZero, + {"vulkan_rewrite_unsigned_compare_zero", + "Workaround for a driver bug where unsigned equality comparisons with zero trigger a buggy " + "peephole optimization on Samsung Xclipse GPUs.", + "https://crbug.com/543420711", ToggleStage::Device}}, {Toggle::WaitIsThreadSafe, {"wait_is_thread_safe", "WaitFor* functions are thread-safe and can be called without the device-lock if implicit "
diff --git a/src/dawn/native/Toggles.h b/src/dawn/native/Toggles.h index b1ef6c1..18e9a3d 100644 --- a/src/dawn/native/Toggles.h +++ b/src/dawn/native/Toggles.h
@@ -197,6 +197,7 @@ UseSpirvReconvergenceMode, VulkanReplaceWorkgroupAtomicStoreWithExchange, VulkanDisallowNPOTDepthStencilMipmaps, + VulkanReplaceUnsignedCompareZero, // Once all backends have been updated to be thread safe for waiting, we can remove this toggle. WaitIsThreadSafe,
diff --git a/src/dawn/native/vulkan/PhysicalDeviceVk.cpp b/src/dawn/native/vulkan/PhysicalDeviceVk.cpp index ab41f17..081b434 100644 --- a/src/dawn/native/vulkan/PhysicalDeviceVk.cpp +++ b/src/dawn/native/vulkan/PhysicalDeviceVk.cpp
@@ -1148,6 +1148,15 @@ // gate this on driver version. // https://crbug.com/487773864 deviceToggles->Default(Toggle::VulkanReplaceWorkgroupAtomicStoreWithExchange, true); + + // Samsung Xclipse GPUs produce incorrect results from certain combinations of bitwise + // operations and select instructions when comparing unsigned values with zero. + // Fixed in driver versions >= 25.x. + // https://crbug.com/543420711 + const gpu_info::DriverVersion kFixedDriverVersion = {25, 0, 0, 0}; + if (GetDriverVersion() < kFixedDriverVersion) { + deviceToggles->Default(Toggle::VulkanReplaceUnsignedCompareZero, true); + } } if (IsSwiftshader()) {
diff --git a/src/dawn/native/vulkan/ShaderModuleVk.cpp b/src/dawn/native/vulkan/ShaderModuleVk.cpp index 6f9c206..6f6bd16 100644 --- a/src/dawn/native/vulkan/ShaderModuleVk.cpp +++ b/src/dawn/native/vulkan/ShaderModuleVk.cpp
@@ -358,6 +358,8 @@ GetDevice()->IsToggleEnabled(Toggle::VulkanCooperativeMatrixStrideIsMatrixElements); req.tintOptions.workarounds.replace_workgroup_atomic_store_with_exchange = GetDevice()->IsToggleEnabled(Toggle::VulkanReplaceWorkgroupAtomicStoreWithExchange); + req.tintOptions.workarounds.replace_unsigned_compare_zero = + GetDevice()->IsToggleEnabled(Toggle::VulkanReplaceUnsignedCompareZero); // Pass matrices to user functions by pointer on Qualcomm devices to workaround a known bug. // See crbug.com/tint/2045.
diff --git a/src/dawn/tests/end2end/ShaderTests.cpp b/src/dawn/tests/end2end/ShaderTests.cpp index d4776f9..cd990d1 100644 --- a/src/dawn/tests/end2end/ShaderTests.cpp +++ b/src/dawn/tests/end2end/ShaderTests.cpp
@@ -3728,6 +3728,65 @@ readbackBuffer.Unmap(); } +// Regression test for a bug on Xclipse GPUs that affects certain combinations of bitwise operations +// and select builtins. +// https://crbug.com/542268656 +TEST_P(ShaderTests, XclipseBitwiseAndSelectBuiltinBug) { + wgpu::ComputePipelineDescriptor pDesc; + pDesc.compute.module = utils::CreateShaderModule(device, R"( +@group(0) @binding(0) var<storage> inputs: array<u32>; +@group(0) @binding(1) var<storage, read_write> outputs: array<u32>; + +@compute @workgroup_size(16) +fn main(@builtin(global_invocation_id) global_invocation_id: vec3u) { + let in_index = global_invocation_id.x; + if (in_index >= 3) { + return; + } + + let value = (inputs[in_index] >> 24u) & 0xff; + let edge = value & 0xfu; + + let out_index = select( + select(0u, 1u, edge == 0u), + 2u, + edge == 8u + ); + + outputs[out_index] = in_index; +} +)"); + wgpu::ComputePipeline pipeline = device.CreateComputePipeline(&pDesc); + + std::vector<uint32_t> inputValues = { + 0u << 24u, // edge=0 -> out_index=1 -> writes 0 at outputs[1] + 3u << 24u, // edge=3 -> out_index=0 -> writes 1 at outputs[0] + 8u << 24u, // edge=8 -> out_index=2 -> writes 2 at outputs[2] + }; + wgpu::Buffer inputBuffer = utils::CreateBufferFromData(device, inputValues.data(), + inputValues.size() * sizeof(uint32_t), + wgpu::BufferUsage::Storage); + + wgpu::Buffer outputBuffer = CreateBuffer(inputValues.size()); + + wgpu::BindGroup bg = utils::MakeBindGroup(device, pipeline.GetBindGroupLayout(0), + {{0, inputBuffer}, {1, outputBuffer}}); + + wgpu::CommandEncoder encoder = device.CreateCommandEncoder(); + wgpu::ComputePassEncoder pass = encoder.BeginComputePass(); + pass.SetBindGroup(0, bg); + pass.SetPipeline(pipeline); + pass.DispatchWorkgroups(1); + pass.End(); + + wgpu::CommandBuffer commands = encoder.Finish(); + queue.Submit(1, &commands); + + EXPECT_BUFFER_U32_EQ(1u, outputBuffer, 0); + EXPECT_BUFFER_U32_EQ(0u, outputBuffer, 4); + EXPECT_BUFFER_U32_EQ(2u, outputBuffer, 8); +} + DAWN_INSTANTIATE_TEST(ShaderTests, D3D11Backend(), D3D12Backend(),
diff --git a/src/tint/lang/spirv/writer/common/options.h b/src/tint/lang/spirv/writer/common/options.h index 4cdf4d7..982533e 100644 --- a/src/tint/lang/spirv/writer/common/options.h +++ b/src/tint/lang/spirv/writer/common/options.h
@@ -146,6 +146,10 @@ /// Set to `true` to replace atomicStore in workgroup memory with atomicExchange. bool replace_workgroup_atomic_store_with_exchange = false; + /// Set to `true` to replace unsigned equality comparisons with zero by relational + /// comparisons against one. + bool replace_unsigned_compare_zero = false; + TINT_REFLECT(Workarounds, polyfill_case_switch, scalarize_max_min_clamp, @@ -162,7 +166,8 @@ polyfill_saturate_as_min_max_f16, cooperative_matrix_stride_is_matrix_elements, collapse_subgroup_min_max, - replace_workgroup_atomic_store_with_exchange); + replace_workgroup_atomic_store_with_exchange, + replace_unsigned_compare_zero); }; /// Any options which are controlled by the presence/absence of a vulkan extension.
diff --git a/src/tint/lang/spirv/writer/raise/BUILD.bazel b/src/tint/lang/spirv/writer/raise/BUILD.bazel index 791396f..0d8fdb1 100644 --- a/src/tint/lang/spirv/writer/raise/BUILD.bazel +++ b/src/tint/lang/spirv/writer/raise/BUILD.bazel
@@ -53,6 +53,7 @@ "pass_matrix_by_pointer.cc", "raise.cc", "remove_unreachable_in_loop_continuing.cc", + "replace_unsigned_compare_zero.cc", "resource_table_helper.cc", "shader_io.cc", "unary_polyfill.cc", @@ -69,6 +70,7 @@ "pass_matrix_by_pointer.h", "raise.h", "remove_unreachable_in_loop_continuing.h", + "replace_unsigned_compare_zero.h", "resource_table_helper.h", "shader_io.h", "unary_polyfill.h", @@ -117,6 +119,7 @@ "merge_return_test.cc", "pass_matrix_by_pointer_test.cc", "remove_unreachable_in_loop_continuing_test.cc", + "replace_unsigned_compare_zero_test.cc", "shader_io_test.cc", "unary_polyfill_test.cc", "var_for_dynamic_index_test.cc",
diff --git a/src/tint/lang/spirv/writer/raise/BUILD.cmake b/src/tint/lang/spirv/writer/raise/BUILD.cmake index fb9fc63..1701445 100644 --- a/src/tint/lang/spirv/writer/raise/BUILD.cmake +++ b/src/tint/lang/spirv/writer/raise/BUILD.cmake
@@ -61,6 +61,8 @@ lang/spirv/writer/raise/raise.h lang/spirv/writer/raise/remove_unreachable_in_loop_continuing.cc lang/spirv/writer/raise/remove_unreachable_in_loop_continuing.h + lang/spirv/writer/raise/replace_unsigned_compare_zero.cc + lang/spirv/writer/raise/replace_unsigned_compare_zero.h lang/spirv/writer/raise/resource_table_helper.cc lang/spirv/writer/raise/resource_table_helper.h lang/spirv/writer/raise/shader_io.cc @@ -119,6 +121,7 @@ lang/spirv/writer/raise/merge_return_test.cc lang/spirv/writer/raise/pass_matrix_by_pointer_test.cc lang/spirv/writer/raise/remove_unreachable_in_loop_continuing_test.cc + lang/spirv/writer/raise/replace_unsigned_compare_zero_test.cc lang/spirv/writer/raise/shader_io_test.cc lang/spirv/writer/raise/unary_polyfill_test.cc lang/spirv/writer/raise/var_for_dynamic_index_test.cc
diff --git a/src/tint/lang/spirv/writer/raise/BUILD.gn b/src/tint/lang/spirv/writer/raise/BUILD.gn index 9dd1b09..547e663 100644 --- a/src/tint/lang/spirv/writer/raise/BUILD.gn +++ b/src/tint/lang/spirv/writer/raise/BUILD.gn
@@ -65,6 +65,8 @@ "raise.h", "remove_unreachable_in_loop_continuing.cc", "remove_unreachable_in_loop_continuing.h", + "replace_unsigned_compare_zero.cc", + "replace_unsigned_compare_zero.h", "resource_table_helper.cc", "resource_table_helper.h", "shader_io.cc", @@ -116,6 +118,7 @@ "merge_return_test.cc", "pass_matrix_by_pointer_test.cc", "remove_unreachable_in_loop_continuing_test.cc", + "replace_unsigned_compare_zero_test.cc", "shader_io_test.cc", "unary_polyfill_test.cc", "var_for_dynamic_index_test.cc",
diff --git a/src/tint/lang/spirv/writer/raise/raise.cc b/src/tint/lang/spirv/writer/raise/raise.cc index 6ba9579..3ae1b29 100644 --- a/src/tint/lang/spirv/writer/raise/raise.cc +++ b/src/tint/lang/spirv/writer/raise/raise.cc
@@ -64,6 +64,7 @@ #include "src/tint/lang/spirv/writer/raise/merge_return.h" #include "src/tint/lang/spirv/writer/raise/pass_matrix_by_pointer.h" #include "src/tint/lang/spirv/writer/raise/remove_unreachable_in_loop_continuing.h" +#include "src/tint/lang/spirv/writer/raise/replace_unsigned_compare_zero.h" #include "src/tint/lang/spirv/writer/raise/resource_table_helper.h" #include "src/tint/lang/spirv/writer/raise/shader_io.h" #include "src/tint/lang/spirv/writer/raise/unary_polyfill.h" @@ -235,6 +236,10 @@ .signed_negation = true, .signed_arithmetic = true, .signed_shiftleft = true}; TINT_CHECK_RESULT(core::ir::transform::SignedIntegerPolyfill(module, signed_integer_cfg)); + if (options.workarounds.replace_unsigned_compare_zero) { + TINT_CHECK_RESULT(raise::ReplaceUnsignedCompareZero(module)); + } + // 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 = {
diff --git a/src/tint/lang/spirv/writer/raise/replace_unsigned_compare_zero.cc b/src/tint/lang/spirv/writer/raise/replace_unsigned_compare_zero.cc new file mode 100644 index 0000000..e00faa9 --- /dev/null +++ b/src/tint/lang/spirv/writer/raise/replace_unsigned_compare_zero.cc
@@ -0,0 +1,93 @@ +// Copyright 2026 The Dawn & Tint Authors +// +// Redistribution and use in source and binary forms, with or without +// modification, are permitted provided that the following conditions are met: +// +// 1. Redistributions of source code must retain the above copyright notice, this +// list of conditions and the following disclaimer. +// +// 2. Redistributions in binary form must reproduce the above copyright notice, +// this list of conditions and the following disclaimer in the documentation +// and/or other materials provided with the distribution. +// +// 3. Neither the name of the copyright holder nor the names of its +// contributors may be used to endorse or promote products derived from +// this software without specific prior written permission. +// +// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" +// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE +// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE +// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE +// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL +// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR +// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER +// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, +// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE +// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. + +#include "src/tint/lang/spirv/writer/raise/replace_unsigned_compare_zero.h" + +#include "src/tint/lang/core/ir/builder.h" +#include "src/tint/lang/core/ir/constant.h" +#include "src/tint/lang/core/ir/core_binary.h" +#include "src/tint/lang/core/ir/module.h" +#include "src/tint/lang/core/ir/validator.h" + +using namespace tint::core::number_suffixes; // NOLINT + +namespace tint::spirv::writer::raise { + +namespace { + +/// PIMPL state for the transform. +struct State { + /// The IR module. + core::ir::Module& ir; + + /// The IR builder. + core::ir::Builder b{ir}; + + /// Process the module. + void Process() { + for (auto* inst : ir.Instructions()) { + auto* binary = inst->As<core::ir::CoreBinary>(); + if (!binary || binary->Op() != core::BinaryOp::kEqual) { + continue; + } + + if (!binary->LHS()->Type()->DeepestElement()->Is<core::type::U32>()) { + continue; + } + + if (IsConstantZero(binary->RHS())) { + binary->SetOp(core::BinaryOp::kLessThan); + binary->SetOperand(core::ir::Binary::kRhsOperandOffset, + b.MatchWidth(1_u, binary->RHS()->Type())); + } else if (IsConstantZero(binary->LHS())) { + binary->SetOp(core::BinaryOp::kGreaterThan); + binary->SetOperand(core::ir::Binary::kLhsOperandOffset, + b.MatchWidth(1_u, binary->LHS()->Type())); + } + } + } + + /// @returns true if @p val is a constant zero value + bool IsConstantZero(const core::ir::Value* val) { + if (auto* c = val->As<core::ir::Constant>()) { + return c->Value()->AllZero(); + } + return false; + } +}; + +} // namespace + +Result<SuccessType> ReplaceUnsignedCompareZero(core::ir::Module& ir) { + AssertValid(ir, "before spirv.ReplaceUnsignedCompareZero"); + + State{ir}.Process(); + + return Success; +} + +} // namespace tint::spirv::writer::raise
diff --git a/src/tint/lang/spirv/writer/raise/replace_unsigned_compare_zero.h b/src/tint/lang/spirv/writer/raise/replace_unsigned_compare_zero.h new file mode 100644 index 0000000..9badf41 --- /dev/null +++ b/src/tint/lang/spirv/writer/raise/replace_unsigned_compare_zero.h
@@ -0,0 +1,48 @@ +// Copyright 2026 The Dawn & Tint Authors +// +// Redistribution and use in source and binary forms, with or without +// modification, are permitted provided that the following conditions are met: +// +// 1. Redistributions of source code must retain the above copyright notice, this +// list of conditions and the following disclaimer. +// +// 2. Redistributions in binary form must reproduce the above copyright notice, +// this list of conditions and the following disclaimer in the documentation +// and/or other materials provided with the distribution. +// +// 3. Neither the name of the copyright holder nor the names of its +// contributors may be used to endorse or promote products derived from +// this software without specific prior written permission. +// +// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" +// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE +// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE +// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE +// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL +// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR +// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER +// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, +// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE +// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. + +#ifndef SRC_TINT_LANG_SPIRV_WRITER_RAISE_REPLACE_UNSIGNED_COMPARE_ZERO_H_ +#define SRC_TINT_LANG_SPIRV_WRITER_RAISE_REPLACE_UNSIGNED_COMPARE_ZERO_H_ + +#include "src/tint/utils/result.h" + +// Forward declarations. +namespace tint::core::ir { +class Module; +} + +namespace tint::spirv::writer::raise { + +/// ReplaceUnsignedCompareZero is a transform that replaces equality comparisons (==) between an +/// unsigned value and zero with relational comparisons (< or >) against one. +/// @param module the module to transform +/// @returns success or failure +Result<SuccessType> ReplaceUnsignedCompareZero(core::ir::Module& module); + +} // namespace tint::spirv::writer::raise + +#endif // SRC_TINT_LANG_SPIRV_WRITER_RAISE_REPLACE_UNSIGNED_COMPARE_ZERO_H_
diff --git a/src/tint/lang/spirv/writer/raise/replace_unsigned_compare_zero_test.cc b/src/tint/lang/spirv/writer/raise/replace_unsigned_compare_zero_test.cc new file mode 100644 index 0000000..34b71ad --- /dev/null +++ b/src/tint/lang/spirv/writer/raise/replace_unsigned_compare_zero_test.cc
@@ -0,0 +1,341 @@ +// Copyright 2026 The Dawn & Tint Authors +// +// Redistribution and use in source and binary forms, with or without +// modification, are permitted provided that the following conditions are met: +// +// 1. Redistributions of source code must retain the above copyright notice, this +// list of conditions and the following disclaimer. +// +// 2. Redistributions in binary form must reproduce the above copyright notice, +// this list of conditions and the following disclaimer in the documentation +// and/or other materials provided with the distribution. +// +// 3. Neither the name of the copyright holder nor the names of its +// contributors may be used to endorse or promote products derived from +// this software without specific prior written permission. +// +// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" +// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE +// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE +// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE +// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL +// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR +// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER +// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, +// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE +// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. + +#include "src/tint/lang/spirv/writer/raise/replace_unsigned_compare_zero.h" + +#include <utility> + +#include "src/tint/lang/core/ir/transform/helper_test.h" + +namespace tint::spirv::writer::raise { +namespace { + +using namespace tint::core::fluent_types; // NOLINT +using namespace tint::core::number_suffixes; // NOLINT + +using SpirvWriter_ReplaceUnsignedCompareZeroTest = core::ir::transform::TransformTest; + +TEST_F(SpirvWriter_ReplaceUnsignedCompareZeroTest, NoModify_SignedEqualZero) { + auto* val = b.FunctionParam("val", ty.i32()); + auto* func = b.Function("foo", ty.bool_()); + func->SetParams({val}); + + b.Append(func->Block(), [&] { + auto* result = b.Equal(val, 0_i); + b.Return(func, result); + }); + + auto* expect = R"( +%foo = func(%val:i32):bool { + $B1: { + %3:bool = eq %val, 0i + ret %3 + } +} +)"; + + Run(ReplaceUnsignedCompareZero); + + EXPECT_EQ(expect, str()); +} + +TEST_F(SpirvWriter_ReplaceUnsignedCompareZeroTest, NoModify_FloatEqualZero) { + auto* val = b.FunctionParam("val", ty.f32()); + auto* func = b.Function("foo", ty.bool_()); + func->SetParams({val}); + + b.Append(func->Block(), [&] { + auto* result = b.Equal(val, 0_f); + b.Return(func, result); + }); + + auto* expect = R"( +%foo = func(%val:f32):bool { + $B1: { + %3:bool = eq %val, 0.0f + ret %3 + } +} +)"; + + Run(ReplaceUnsignedCompareZero); + + EXPECT_EQ(expect, str()); +} + +TEST_F(SpirvWriter_ReplaceUnsignedCompareZeroTest, NoModify_UnsignedEqualNonZero) { + auto* val = b.FunctionParam("val", ty.u32()); + auto* func = b.Function("foo", ty.bool_()); + func->SetParams({val}); + + b.Append(func->Block(), [&] { + auto* result = b.Equal(val, 5_u); + b.Return(func, result); + }); + + auto* expect = R"( +%foo = func(%val:u32):bool { + $B1: { + %3:bool = eq %val, 5u + ret %3 + } +} +)"; + + Run(ReplaceUnsignedCompareZero); + + EXPECT_EQ(expect, str()); +} + +TEST_F(SpirvWriter_ReplaceUnsignedCompareZeroTest, NoModify_UnsignedLessThanZero) { + auto* val = b.FunctionParam("val", ty.u32()); + auto* func = b.Function("foo", ty.bool_()); + func->SetParams({val}); + + b.Append(func->Block(), [&] { + auto* result = b.LessThan(val, 0_u); + b.Return(func, result); + }); + + auto* expect = R"( +%foo = func(%val:u32):bool { + $B1: { + %3:bool = lt %val, 0u + ret %3 + } +} +)"; + + Run(ReplaceUnsignedCompareZero); + + EXPECT_EQ(expect, str()); +} + +TEST_F(SpirvWriter_ReplaceUnsignedCompareZeroTest, Scalar_RhsZero) { + auto* val = b.FunctionParam("val", ty.u32()); + auto* func = b.Function("foo", ty.bool_()); + func->SetParams({val}); + + b.Append(func->Block(), [&] { + auto* result = b.Equal(val, 0_u); + b.Return(func, result); + }); + + auto* src = R"( +%foo = func(%val:u32):bool { + $B1: { + %3:bool = eq %val, 0u + ret %3 + } +} +)"; + EXPECT_EQ(src, str()); + + auto* expect = R"( +%foo = func(%val:u32):bool { + $B1: { + %3:bool = lt %val, 1u + ret %3 + } +} +)"; + + Run(ReplaceUnsignedCompareZero); + + EXPECT_EQ(expect, str()); +} + +TEST_F(SpirvWriter_ReplaceUnsignedCompareZeroTest, Scalar_LhsZero) { + auto* val = b.FunctionParam("val", ty.u32()); + auto* func = b.Function("foo", ty.bool_()); + func->SetParams({val}); + + b.Append(func->Block(), [&] { + auto* result = b.Equal(0_u, val); + b.Return(func, result); + }); + + auto* src = R"( +%foo = func(%val:u32):bool { + $B1: { + %3:bool = eq 0u, %val + ret %3 + } +} +)"; + EXPECT_EQ(src, str()); + + auto* expect = R"( +%foo = func(%val:u32):bool { + $B1: { + %3:bool = gt 1u, %val + ret %3 + } +} +)"; + + Run(ReplaceUnsignedCompareZero); + + EXPECT_EQ(expect, str()); +} + +TEST_F(SpirvWriter_ReplaceUnsignedCompareZeroTest, Vector_RhsZero_Vec2u) { + auto* val = b.FunctionParam("val", ty.vec2u()); + auto* func = b.Function("foo", ty.vec2<bool>()); + func->SetParams({val}); + + b.Append(func->Block(), [&] { + auto* result = b.Equal(val, b.Zero<vec2<u32>>()); + b.Return(func, result); + }); + + auto* src = R"( +%foo = func(%val:vec2<u32>):vec2<bool> { + $B1: { + %3:vec2<bool> = eq %val, vec2<u32>(0u) + ret %3 + } +} +)"; + EXPECT_EQ(src, str()); + + auto* expect = R"( +%foo = func(%val:vec2<u32>):vec2<bool> { + $B1: { + %3:vec2<bool> = lt %val, vec2<u32>(1u) + ret %3 + } +} +)"; + + Run(ReplaceUnsignedCompareZero); + + EXPECT_EQ(expect, str()); +} + +TEST_F(SpirvWriter_ReplaceUnsignedCompareZeroTest, Vector_LhsZero_Vec3u) { + auto* val = b.FunctionParam("val", ty.vec3u()); + auto* func = b.Function("foo", ty.vec3<bool>()); + func->SetParams({val}); + + b.Append(func->Block(), [&] { + auto* result = b.Equal(b.Zero<vec3<u32>>(), val); + b.Return(func, result); + }); + + auto* src = R"( +%foo = func(%val:vec3<u32>):vec3<bool> { + $B1: { + %3:vec3<bool> = eq vec3<u32>(0u), %val + ret %3 + } +} +)"; + EXPECT_EQ(src, str()); + + auto* expect = R"( +%foo = func(%val:vec3<u32>):vec3<bool> { + $B1: { + %3:vec3<bool> = gt vec3<u32>(1u), %val + ret %3 + } +} +)"; + + Run(ReplaceUnsignedCompareZero); + + EXPECT_EQ(expect, str()); +} + +TEST_F(SpirvWriter_ReplaceUnsignedCompareZeroTest, Vector_RhsZero_Vec4u) { + auto* val = b.FunctionParam("val", ty.vec4u()); + auto* func = b.Function("foo", ty.vec4<bool>()); + func->SetParams({val}); + + b.Append(func->Block(), [&] { + auto* result = b.Equal(val, b.Zero<vec4<u32>>()); + b.Return(func, result); + }); + + auto* src = R"( +%foo = func(%val:vec4<u32>):vec4<bool> { + $B1: { + %3:vec4<bool> = eq %val, vec4<u32>(0u) + ret %3 + } +} +)"; + EXPECT_EQ(src, str()); + + auto* expect = R"( +%foo = func(%val:vec4<u32>):vec4<bool> { + $B1: { + %3:vec4<bool> = lt %val, vec4<u32>(1u) + ret %3 + } +} +)"; + + Run(ReplaceUnsignedCompareZero); + + EXPECT_EQ(expect, str()); +} + +TEST_F(SpirvWriter_ReplaceUnsignedCompareZeroTest, BothZero) { + auto* func = b.Function("foo", ty.bool_()); + + b.Append(func->Block(), [&] { + auto* result = b.Equal(0_u, 0_u); + b.Return(func, result); + }); + + auto* src = R"( +%foo = func():bool { + $B1: { + %2:bool = eq 0u, 0u + ret %2 + } +} +)"; + EXPECT_EQ(src, str()); + + auto* expect = R"( +%foo = func():bool { + $B1: { + %2:bool = lt 0u, 1u + ret %2 + } +} +)"; + + Run(ReplaceUnsignedCompareZero); + + EXPECT_EQ(expect, str()); +} + +} // namespace +} // namespace tint::spirv::writer::raise
diff --git a/src/tint/lang/spirv/writer/writer_fuzz.cc b/src/tint/lang/spirv/writer/writer_fuzz.cc index 6412570..ffff431 100644 --- a/src/tint/lang/spirv/writer/writer_fuzz.cc +++ b/src/tint/lang/spirv/writer/writer_fuzz.cc
@@ -88,6 +88,7 @@ bool polyfill_distance_scalar_float; bool collapse_subgroup_min_max; bool replace_workgroup_atomic_store_with_exchange; + bool replace_unsigned_compare_zero; /// Reflect the fields of this class so that it can be used by tint::ForeachField() TINT_REFLECT(FuzzedOptions, @@ -125,7 +126,8 @@ polyfill_length_scalar_float, polyfill_distance_scalar_float, collapse_subgroup_min_max, - replace_workgroup_atomic_store_with_exchange); + replace_workgroup_atomic_store_with_exchange, + replace_unsigned_compare_zero); TINT_REFLECT_HASH_CODE(FuzzedOptions); }; @@ -355,6 +357,8 @@ options.workarounds.collapse_subgroup_min_max = fuzzed_options.collapse_subgroup_min_max; options.workarounds.replace_workgroup_atomic_store_with_exchange = fuzzed_options.replace_workgroup_atomic_store_with_exchange; + options.workarounds.replace_unsigned_compare_zero = + fuzzed_options.replace_unsigned_compare_zero; options.multisampled_framebuffer_fetch = fuzzed_options.multisampled_framebuffer_fetch; TINT_CHECK_RESULT_UNWRAP(output, Generate(module, options));