spirv-reader: specop integer comparisons
IEqual
INotEqual
SGreaterThan
SGreaterThanEqual
SLessThan
SLessThanEqual
UGreaterThan
UGreaterThanEqual
ULessThan
ULessThanEqual
Fixed: 402727480 402727143 402727302 402725914 402726714 402727617 402726907 402727620 402726892 402727337
Change-Id: I33fe259205d7e39227cb7a3f556765f5643de76a
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/250494
Commit-Queue: David Neto <dneto@google.com>
Reviewed-by: James Price <jrprice@google.com>
diff --git a/src/tint/lang/spirv/reader/parser/parser.cc b/src/tint/lang/spirv/reader/parser/parser.cc
index b5fdb87..432770c 100644
--- a/src/tint/lang/spirv/reader/parser/parser.cc
+++ b/src/tint/lang/spirv/reader/parser/parser.cc
@@ -355,6 +355,36 @@
case spv::Op::OpBitwiseXor:
EmitSpirvExplicitBuiltinCall(inst, spirv::BuiltinFn::kBitwiseXor, 3);
break;
+ case spv::Op::OpIEqual:
+ EmitSpirvBuiltinCall(inst, spirv::BuiltinFn::kEqual, 3);
+ break;
+ case spv::Op::OpINotEqual:
+ EmitSpirvBuiltinCall(inst, spirv::BuiltinFn::kNotEqual, 3);
+ break;
+ case spv::Op::OpSGreaterThan:
+ EmitSpirvBuiltinCall(inst, spirv::BuiltinFn::kSGreaterThan, 3);
+ break;
+ case spv::Op::OpSGreaterThanEqual:
+ EmitSpirvBuiltinCall(inst, spirv::BuiltinFn::kSGreaterThanEqual, 3);
+ break;
+ case spv::Op::OpSLessThan:
+ EmitSpirvBuiltinCall(inst, spirv::BuiltinFn::kSLessThan, 3);
+ break;
+ case spv::Op::OpSLessThanEqual:
+ EmitSpirvBuiltinCall(inst, spirv::BuiltinFn::kSLessThanEqual, 3);
+ break;
+ case spv::Op::OpUGreaterThan:
+ EmitSpirvBuiltinCall(inst, spirv::BuiltinFn::kUGreaterThan, 3);
+ break;
+ case spv::Op::OpUGreaterThanEqual:
+ EmitSpirvBuiltinCall(inst, spirv::BuiltinFn::kUGreaterThanEqual, 3);
+ break;
+ case spv::Op::OpULessThan:
+ EmitSpirvBuiltinCall(inst, spirv::BuiltinFn::kULessThan, 3);
+ break;
+ case spv::Op::OpULessThanEqual:
+ EmitSpirvBuiltinCall(inst, spirv::BuiltinFn::kULessThanEqual, 3);
+ break;
case spv::Op::OpLogicalAnd:
EmitBinary(inst, core::BinaryOp::kAnd, 3);
break;
@@ -2958,8 +2988,11 @@
inst.result_id());
}
- void EmitSpirvBuiltinCall(const spvtools::opt::Instruction& inst, spirv::BuiltinFn fn) {
- Emit(b_.Call<spirv::ir::BuiltinCall>(Type(inst.type_id()), fn, Args(inst, 2)),
+ void EmitSpirvBuiltinCall(const spvtools::opt::Instruction& inst,
+ spirv::BuiltinFn fn,
+ uint32_t first_operand_idx = 2) {
+ Emit(b_.Call<spirv::ir::BuiltinCall>(Type(inst.type_id()), fn,
+ Args(inst, first_operand_idx)),
inst.result_id());
}
diff --git a/src/tint/lang/spirv/reader/parser/var_test.cc b/src/tint/lang/spirv/reader/parser/var_test.cc
index 11f570e..59f561e 100644
--- a/src/tint/lang/spirv/reader/parser/var_test.cc
+++ b/src/tint/lang/spirv/reader/parser/var_test.cc
@@ -2300,6 +2300,366 @@
)");
}
+TEST_F(SpirvParserTest, Var_OpSpecConstantOp_IEqual) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %myconst "myconst"
+ %void = OpTypeVoid
+ %bool = OpTypeBool
+ %u32 = OpTypeInt 32 0
+ %i32 = OpTypeInt 32 1
+ %one = OpConstant %i32 1
+ %two = OpConstant %u32 2
+ %myconst = OpSpecConstantOp %bool IEqual %one %two
+ %voidfn = OpTypeFunction %void
+ %main = OpFunction %void None %voidfn
+ %main_entry = OpLabel
+ %1 = OpCopyObject %bool %myconst
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+$B1: { # root
+ %1:bool = spirv.equal 1i, 2u
+ %myconst:bool = override %1
+}
+
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B2: {
+ %4:bool = let %myconst
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, Var_OpSpecConstantOp_INotEqual) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %myconst "myconst"
+ %void = OpTypeVoid
+ %bool = OpTypeBool
+ %u32 = OpTypeInt 32 0
+ %i32 = OpTypeInt 32 1
+ %one = OpConstant %i32 1
+ %two = OpConstant %u32 2
+ %myconst = OpSpecConstantOp %bool INotEqual %one %two
+ %voidfn = OpTypeFunction %void
+ %main = OpFunction %void None %voidfn
+ %main_entry = OpLabel
+ %1 = OpCopyObject %bool %myconst
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+$B1: { # root
+ %1:bool = spirv.not_equal 1i, 2u
+ %myconst:bool = override %1
+}
+
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B2: {
+ %4:bool = let %myconst
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, Var_OpSpecConstantOp_SGreaterThan) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %myconst "myconst"
+ %void = OpTypeVoid
+ %bool = OpTypeBool
+ %u32 = OpTypeInt 32 0
+ %i32 = OpTypeInt 32 1
+ %one = OpConstant %i32 1
+ %two = OpConstant %u32 2
+ %myconst = OpSpecConstantOp %bool SGreaterThan %one %two
+ %voidfn = OpTypeFunction %void
+ %main = OpFunction %void None %voidfn
+ %main_entry = OpLabel
+ %1 = OpCopyObject %bool %myconst
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+$B1: { # root
+ %1:bool = spirv.s_greater_than 1i, 2u
+ %myconst:bool = override %1
+}
+
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B2: {
+ %4:bool = let %myconst
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, Var_OpSpecConstantOp_SGreaterThanEqual) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %myconst "myconst"
+ %void = OpTypeVoid
+ %bool = OpTypeBool
+ %u32 = OpTypeInt 32 0
+ %i32 = OpTypeInt 32 1
+ %one = OpConstant %i32 1
+ %two = OpConstant %u32 2
+ %myconst = OpSpecConstantOp %bool SGreaterThanEqual %one %two
+ %voidfn = OpTypeFunction %void
+ %main = OpFunction %void None %voidfn
+ %main_entry = OpLabel
+ %1 = OpCopyObject %bool %myconst
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+$B1: { # root
+ %1:bool = spirv.s_greater_than_equal 1i, 2u
+ %myconst:bool = override %1
+}
+
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B2: {
+ %4:bool = let %myconst
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, Var_OpSpecConstantOp_SLessThan) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %myconst "myconst"
+ %void = OpTypeVoid
+ %bool = OpTypeBool
+ %u32 = OpTypeInt 32 0
+ %i32 = OpTypeInt 32 1
+ %one = OpConstant %i32 1
+ %two = OpConstant %u32 2
+ %myconst = OpSpecConstantOp %bool SLessThan %one %two
+ %voidfn = OpTypeFunction %void
+ %main = OpFunction %void None %voidfn
+ %main_entry = OpLabel
+ %1 = OpCopyObject %bool %myconst
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+$B1: { # root
+ %1:bool = spirv.s_less_than 1i, 2u
+ %myconst:bool = override %1
+}
+
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B2: {
+ %4:bool = let %myconst
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, Var_OpSpecConstantOp_SLessThanEqual) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %myconst "myconst"
+ %void = OpTypeVoid
+ %bool = OpTypeBool
+ %u32 = OpTypeInt 32 0
+ %i32 = OpTypeInt 32 1
+ %one = OpConstant %i32 1
+ %two = OpConstant %u32 2
+ %myconst = OpSpecConstantOp %bool SLessThanEqual %one %two
+ %voidfn = OpTypeFunction %void
+ %main = OpFunction %void None %voidfn
+ %main_entry = OpLabel
+ %1 = OpCopyObject %bool %myconst
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+$B1: { # root
+ %1:bool = spirv.s_less_than_equal 1i, 2u
+ %myconst:bool = override %1
+}
+
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B2: {
+ %4:bool = let %myconst
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, Var_OpSpecConstantOp_UGreaterThan) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %myconst "myconst"
+ %void = OpTypeVoid
+ %bool = OpTypeBool
+ %u32 = OpTypeInt 32 0
+ %i32 = OpTypeInt 32 1
+ %one = OpConstant %i32 1
+ %two = OpConstant %u32 2
+ %myconst = OpSpecConstantOp %bool UGreaterThan %one %two
+ %voidfn = OpTypeFunction %void
+ %main = OpFunction %void None %voidfn
+ %main_entry = OpLabel
+ %1 = OpCopyObject %bool %myconst
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+$B1: { # root
+ %1:bool = spirv.u_greater_than 1i, 2u
+ %myconst:bool = override %1
+}
+
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B2: {
+ %4:bool = let %myconst
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, Var_OpSpecConstantOp_UGreaterThanEqual) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %myconst "myconst"
+ %void = OpTypeVoid
+ %bool = OpTypeBool
+ %u32 = OpTypeInt 32 0
+ %i32 = OpTypeInt 32 1
+ %one = OpConstant %i32 1
+ %two = OpConstant %u32 2
+ %myconst = OpSpecConstantOp %bool UGreaterThanEqual %one %two
+ %voidfn = OpTypeFunction %void
+ %main = OpFunction %void None %voidfn
+ %main_entry = OpLabel
+ %1 = OpCopyObject %bool %myconst
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+$B1: { # root
+ %1:bool = spirv.u_greater_than_equal 1i, 2u
+ %myconst:bool = override %1
+}
+
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B2: {
+ %4:bool = let %myconst
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, Var_OpSpecConstantOp_ULessThan) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %myconst "myconst"
+ %void = OpTypeVoid
+ %bool = OpTypeBool
+ %u32 = OpTypeInt 32 0
+ %i32 = OpTypeInt 32 1
+ %one = OpConstant %i32 1
+ %two = OpConstant %u32 2
+ %myconst = OpSpecConstantOp %bool ULessThan %one %two
+ %voidfn = OpTypeFunction %void
+ %main = OpFunction %void None %voidfn
+ %main_entry = OpLabel
+ %1 = OpCopyObject %bool %myconst
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+$B1: { # root
+ %1:bool = spirv.u_less_than 1i, 2u
+ %myconst:bool = override %1
+}
+
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B2: {
+ %4:bool = let %myconst
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, Var_OpSpecConstantOp_ULessThanEqual) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %myconst "myconst"
+ %void = OpTypeVoid
+ %bool = OpTypeBool
+ %u32 = OpTypeInt 32 0
+ %i32 = OpTypeInt 32 1
+ %one = OpConstant %i32 1
+ %two = OpConstant %u32 2
+ %myconst = OpSpecConstantOp %bool ULessThanEqual %one %two
+ %voidfn = OpTypeFunction %void
+ %main = OpFunction %void None %voidfn
+ %main_entry = OpLabel
+ %1 = OpCopyObject %bool %myconst
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+$B1: { # root
+ %1:bool = spirv.u_less_than_equal 1i, 2u
+ %myconst:bool = override %1
+}
+
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B2: {
+ %4:bool = let %myconst
+ ret
+ }
+}
+)");
+}
+
// In the case of all literals, SPIR-V opt treats the `OpSpecConstantComposite` as an
// `OpConstantComposite` so it appears in the constant manager already. This then needs no handling
// on our side.
