[spirv-reader][ir] Convert the `OpFUnord` instructions.
This CL converts the `OpFUnordEqual`, `OpFUnordNotEqual`,
`OpFUnordGreaterThan`, `OpFUnordGreaterThanEqual`, `OpFUnordLessThan`,
and `OpFUnordLessThanEqual` into the equivalent IR instructions.
Bug: 391485740, 391487150, 391487131, 391486668, 391486344, 391486006
Change-Id: I4140a05f9d09177e788e97c48e6df9bcdd4f40cc
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/226897
Commit-Queue: dan sinclair <dsinclair@chromium.org>
Reviewed-by: James Price <jrprice@google.com>
diff --git a/src/tint/lang/spirv/reader/parser/logical_test.cc b/src/tint/lang/spirv/reader/parser/logical_test.cc
index fd83e1c..630dadc 100644
--- a/src/tint/lang/spirv/reader/parser/logical_test.cc
+++ b/src/tint/lang/spirv/reader/parser/logical_test.cc
@@ -39,8 +39,8 @@
return out;
}
-using SpirvParser_LogicalTest = SpirvParserTestWithParam<SpirvLogicalParam>;
-TEST_P(SpirvParser_LogicalTest, FOrd_Scalar) {
+using SpirvParser_FOrdLogicalTest = SpirvParserTestWithParam<SpirvLogicalParam>;
+TEST_P(SpirvParser_FOrdLogicalTest, Scalar) {
auto params = GetParam();
EXPECT_IR(R"(
OpCapability Shader
@@ -76,7 +76,7 @@
)");
}
-TEST_P(SpirvParser_LogicalTest, FOrd_Vector) {
+TEST_P(SpirvParser_FOrdLogicalTest, Vector) {
auto params = GetParam();
EXPECT_IR(R"(
OpCapability Shader
@@ -113,7 +113,7 @@
}
INSTANTIATE_TEST_SUITE_P(SpirvParser,
- SpirvParser_LogicalTest,
+ SpirvParser_FOrdLogicalTest,
testing::Values(SpirvLogicalParam{"Equal", "eq"},
SpirvLogicalParam{"NotEqual", "neq"},
SpirvLogicalParam{"GreaterThan", "gt"},
@@ -121,5 +121,89 @@
SpirvLogicalParam{"LessThan", "lt"},
SpirvLogicalParam{"LessThanEqual", "lte"}));
+using SpirvParser_FUnordLogicalTest = SpirvParserTestWithParam<SpirvLogicalParam>;
+TEST_P(SpirvParser_FUnordLogicalTest, FOrd_Scalar) {
+ auto params = GetParam();
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpCapability Float16
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ %void = OpTypeVoid
+ %bool = OpTypeBool
+ %f32 = OpTypeFloat 32
+ %one = OpConstant %f32 1
+ %two = OpConstant %f32 2
+ %v2bool = OpTypeVector %bool 2
+ %v2float = OpTypeVector %f32 2
+ %v2one = OpConstantComposite %v2float %one %one
+ %v2two = OpConstantComposite %v2float %two %two
+ %ep_type = OpTypeFunction %void
+ %main = OpFunction %void None %ep_type
+ %main_start = OpLabel
+ %1 = OpFUnord)" +
+ params.spv_name + R"( %bool %one %two
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:bool = )" + params.wgsl_name +
+ R"( 1.0f, 2.0f
+ %3:bool = not %2
+ ret
+ }
+}
+)");
+}
+
+TEST_P(SpirvParser_FUnordLogicalTest, FOrd_Vector) {
+ auto params = GetParam();
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpCapability Float16
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ %void = OpTypeVoid
+ %bool = OpTypeBool
+ %f32 = OpTypeFloat 32
+ %one = OpConstant %f32 1
+ %two = OpConstant %f32 2
+ %v2bool = OpTypeVector %bool 2
+ %v2float = OpTypeVector %f32 2
+ %v2one = OpConstantComposite %v2float %one %one
+ %v2two = OpConstantComposite %v2float %two %two
+ %ep_type = OpTypeFunction %void
+ %main = OpFunction %void None %ep_type
+ %main_start = OpLabel
+ %1 = OpFUnord)" +
+ params.spv_name + R"( %v2bool %v2one %v2two
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:vec2<bool> = )" +
+ params.wgsl_name + R"( vec2<f32>(1.0f), vec2<f32>(2.0f)
+ %3:vec2<bool> = not %2
+ ret
+ }
+}
+)");
+}
+
+INSTANTIATE_TEST_SUITE_P(SpirvParser,
+ SpirvParser_FUnordLogicalTest,
+ testing::Values(SpirvLogicalParam{"Equal", "neq"},
+ SpirvLogicalParam{"NotEqual", "eq"},
+ SpirvLogicalParam{"GreaterThan", "lte"},
+ SpirvLogicalParam{"GreaterThanEqual", "lt"},
+ SpirvLogicalParam{"LessThan", "gte"},
+ SpirvLogicalParam{"LessThanEqual", "gt"}));
+
} // namespace
} // namespace tint::spirv::reader
diff --git a/src/tint/lang/spirv/reader/parser/parser.cc b/src/tint/lang/spirv/reader/parser/parser.cc
index bc98553..045cbcc 100644
--- a/src/tint/lang/spirv/reader/parser/parser.cc
+++ b/src/tint/lang/spirv/reader/parser/parser.cc
@@ -740,6 +740,24 @@
case spv::Op::OpFOrdLessThanEqual:
EmitBinary(inst, core::BinaryOp::kLessThanEqual);
break;
+ case spv::Op::OpFUnordEqual:
+ EmitInvertedBinary(inst, core::BinaryOp::kNotEqual);
+ break;
+ case spv::Op::OpFUnordNotEqual:
+ EmitInvertedBinary(inst, core::BinaryOp::kEqual);
+ break;
+ case spv::Op::OpFUnordGreaterThan:
+ EmitInvertedBinary(inst, core::BinaryOp::kLessThanEqual);
+ break;
+ case spv::Op::OpFUnordGreaterThanEqual:
+ EmitInvertedBinary(inst, core::BinaryOp::kLessThan);
+ break;
+ case spv::Op::OpFUnordLessThan:
+ EmitInvertedBinary(inst, core::BinaryOp::kGreaterThanEqual);
+ break;
+ case spv::Op::OpFUnordLessThanEqual:
+ EmitInvertedBinary(inst, core::BinaryOp::kGreaterThan);
+ break;
case spv::Op::OpISub:
EmitSpirvExplicitBuiltinCall(inst, spirv::BuiltinFn::kSub);
break;
@@ -1256,6 +1274,18 @@
Emit(binary, inst.result_id());
}
+ /// @param inst the SPIR-V instruction
+ /// @param op the binary operator to use
+ void EmitInvertedBinary(const spvtools::opt::Instruction& inst, core::BinaryOp op) {
+ auto* lhs = Value(inst.GetSingleWordOperand(2));
+ auto* rhs = Value(inst.GetSingleWordOperand(3));
+ auto* binary = b_.Binary(op, Type(inst.type_id()), lhs, rhs);
+ EmitWithoutSpvResult(binary);
+
+ auto* res = b_.Not(Type(inst.type_id()), binary);
+ Emit(res, inst.result_id());
+ }
+
/// @param inst the SPIR-V instruction for OpCompositeExtract
void EmitCompositeExtract(const spvtools::opt::Instruction& inst) {
Vector<core::ir::Value*, 4> indices;
