commit c178ddf9a6082ba4276228f34ae970ccb84bdb36
author dan sinclair <dsinclair@chromium.org> Thu Feb 20 10:06:25 2025 -0800
committer Dawn LUCI CQ <dawn-scoped@luci-project-accounts.iam.gserviceaccount.com> Thu Feb 20 10:06:25 2025 -0800
tree 40c7505d26d69044682bab359f28fff8898e4486
parent 458fc9935c8116484b6970764e65fdb9d725e05b

[spirv-reader][ir] Support remaining OpFOrd instructions.

Add conversions for `OpFOrdNotEqual`, `OpFOrdGreaterThan`,
`OpFOrdGreaterThanEqual`, `OpFOrdLessThan`, and `OpFOrdLessThanEqual`.

Bug: 391486260, 391486721, 391487147, 391486343, 397725475
Change-Id: Ic82b9d26290ae0d36d1bf3cab4c5b4074426f047
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/226915
Reviewed-by: James Price <jrprice@google.com>
Commit-Queue: dan sinclair <dsinclair@chromium.org>

src/tint/lang/spirv/reader/parser/logical_test.cc

diff --git a/src/tint/lang/spirv/reader/parser/logical_test.cc b/src/tint/lang/spirv/reader/parser/logical_test.cc
index 7358a7f..fd83e1c 100644
--- a/src/tint/lang/spirv/reader/parser/logical_test.cc
+++ b/src/tint/lang/spirv/reader/parser/logical_test.cc
@@ -30,7 +30,18 @@
 namespace tint::spirv::reader {
 namespace {
 
-TEST_F(SpirvParserTest, FOrdEqual_Scalar) {
+struct SpirvLogicalParam {
+    std::string spv_name;
+    std::string wgsl_name;
+};
+[[maybe_unused]] inline std::ostream& operator<<(std::ostream& out, SpirvLogicalParam c) {
+    out << c.spv_name;
+    return out;
+}
+
+using SpirvParser_LogicalTest = SpirvParserTestWithParam<SpirvLogicalParam>;
+TEST_P(SpirvParser_LogicalTest, FOrd_Scalar) {
+    auto params = GetParam();
     EXPECT_IR(R"(
                OpCapability Shader
                OpCapability Float16
@@ -49,21 +60,24 @@
     %ep_type = OpTypeFunction %void
        %main = OpFunction %void None %ep_type
  %main_start = OpLabel
-          %1 = OpFOrdEqual %bool %one %two
+          %1 = OpFOrd)" +
+                  params.spv_name + R"( %bool %one %two
                OpReturn
                OpFunctionEnd
 )",
               R"(
 %main = @compute @workgroup_size(1u, 1u, 1u) func():void {
   $B1: {
-    %2:bool = eq 1.0f, 2.0f
+    %2:bool = )" + params.wgsl_name +
+                  R"( 1.0f, 2.0f
     ret
   }
 }
 )");
 }
 
-TEST_F(SpirvParserTest, FOrdEqual_Vector) {
+TEST_P(SpirvParser_LogicalTest, FOrd_Vector) {
+    auto params = GetParam();
     EXPECT_IR(R"(
                OpCapability Shader
                OpCapability Float16
@@ -82,19 +96,30 @@
     %ep_type = OpTypeFunction %void
        %main = OpFunction %void None %ep_type
  %main_start = OpLabel
-          %1 = OpFOrdEqual %v2bool %v2one %v2two
+          %1 = OpFOrd)" +
+                  params.spv_name + R"( %v2bool %v2one %v2two
                OpReturn
                OpFunctionEnd
 )",
               R"(
 %main = @compute @workgroup_size(1u, 1u, 1u) func():void {
   $B1: {
-    %2:vec2<bool> = eq vec2<f32>(1.0f), vec2<f32>(2.0f)
+    %2:vec2<bool> = )" +
+                  params.wgsl_name + R"( vec2<f32>(1.0f), vec2<f32>(2.0f)
     ret
   }
 }
 )");
 }
 
+INSTANTIATE_TEST_SUITE_P(SpirvParser,
+                         SpirvParser_LogicalTest,
+                         testing::Values(SpirvLogicalParam{"Equal", "eq"},
+                                         SpirvLogicalParam{"NotEqual", "neq"},
+                                         SpirvLogicalParam{"GreaterThan", "gt"},
+                                         SpirvLogicalParam{"GreaterThanEqual", "gte"},
+                                         SpirvLogicalParam{"LessThan", "lt"},
+                                         SpirvLogicalParam{"LessThanEqual", "lte"}));
+
 }  // namespace
 }  // namespace tint::spirv::reader

src/tint/lang/spirv/reader/parser/parser.cc

diff --git a/src/tint/lang/spirv/reader/parser/parser.cc b/src/tint/lang/spirv/reader/parser/parser.cc
index 7c944d8..bc98553 100644
--- a/src/tint/lang/spirv/reader/parser/parser.cc
+++ b/src/tint/lang/spirv/reader/parser/parser.cc
@@ -725,6 +725,21 @@
                 case spv::Op::OpFOrdEqual:
                     EmitBinary(inst, core::BinaryOp::kEqual);
                     break;
+                case spv::Op::OpFOrdNotEqual:
+                    EmitBinary(inst, core::BinaryOp::kNotEqual);
+                    break;
+                case spv::Op::OpFOrdGreaterThan:
+                    EmitBinary(inst, core::BinaryOp::kGreaterThan);
+                    break;
+                case spv::Op::OpFOrdGreaterThanEqual:
+                    EmitBinary(inst, core::BinaryOp::kGreaterThanEqual);
+                    break;
+                case spv::Op::OpFOrdLessThan:
+                    EmitBinary(inst, core::BinaryOp::kLessThan);
+                    break;
+                case spv::Op::OpFOrdLessThanEqual:
+                    EmitBinary(inst, core::BinaryOp::kLessThanEqual);
+                    break;
                 case spv::Op::OpISub:
                     EmitSpirvExplicitBuiltinCall(inst, spirv::BuiltinFn::kSub);
                     break;