spirv-reader: specop integer shift
ShiftLeftLogical
ShiftRightLogical
ShiftRightArithmetic
Fixed: 402726662 402726660 402725870
Change-Id: I92a0cf73d51a8d466a2095ea2f03b1c77212faef
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/250514
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 432770c..53c3beb 100644
--- a/src/tint/lang/spirv/reader/parser/parser.cc
+++ b/src/tint/lang/spirv/reader/parser/parser.cc
@@ -439,6 +439,18 @@
case spv::Op::OpSRem:
EmitSpirvExplicitBuiltinCall(inst, spirv::BuiltinFn::kSMod, 3);
break;
+ case spv::Op::OpShiftLeftLogical:
+ EmitSpirvExplicitBuiltinCall(inst, spirv::BuiltinFn::kShiftLeftLogical,
+ 3);
+ break;
+ case spv::Op::OpShiftRightLogical:
+ EmitSpirvExplicitBuiltinCall(inst, spirv::BuiltinFn::kShiftRightLogical,
+ 3);
+ break;
+ case spv::Op::OpShiftRightArithmetic:
+ EmitSpirvExplicitBuiltinCall(
+ inst, spirv::BuiltinFn::kShiftRightArithmetic, 3);
+ break;
default:
TINT_ICE() << "Unknown spec constant operation: " << op;
}
diff --git a/src/tint/lang/spirv/reader/parser/var_test.cc b/src/tint/lang/spirv/reader/parser/var_test.cc
index 59f561e..7678d2f 100644
--- a/src/tint/lang/spirv/reader/parser/var_test.cc
+++ b/src/tint/lang/spirv/reader/parser/var_test.cc
@@ -2660,6 +2660,111 @@
)");
}
+TEST_F(SpirvParserTest, Var_OpSpecConstantOp_ShiftLeftLogical) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %myconst "myconst"
+ %void = OpTypeVoid
+ %u32 = OpTypeInt 32 0
+ %i32 = OpTypeInt 32 1
+ %one = OpConstant %i32 1
+ %two = OpConstant %u32 2
+ %myconst = OpSpecConstantOp %u32 ShiftLeftLogical %one %two
+ %voidfn = OpTypeFunction %void
+ %main = OpFunction %void None %voidfn
+ %main_entry = OpLabel
+ %1 = OpCopyObject %u32 %myconst
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+$B1: { # root
+ %1:u32 = spirv.shift_left_logical<u32> 1i, 2u
+ %myconst:u32 = override %1
+}
+
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B2: {
+ %4:u32 = let %myconst
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, Var_OpSpecConstantOp_ShiftRightLogical) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %myconst "myconst"
+ %void = OpTypeVoid
+ %u32 = OpTypeInt 32 0
+ %i32 = OpTypeInt 32 1
+ %one = OpConstant %i32 1
+ %two = OpConstant %u32 2
+ %myconst = OpSpecConstantOp %u32 ShiftRightLogical %one %two
+ %voidfn = OpTypeFunction %void
+ %main = OpFunction %void None %voidfn
+ %main_entry = OpLabel
+ %1 = OpCopyObject %u32 %myconst
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+$B1: { # root
+ %1:u32 = spirv.shift_right_logical<u32> 1i, 2u
+ %myconst:u32 = override %1
+}
+
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B2: {
+ %4:u32 = let %myconst
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, Var_OpSpecConstantOp_ShiftRightArithmetic) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %myconst "myconst"
+ %void = OpTypeVoid
+ %u32 = OpTypeInt 32 0
+ %i32 = OpTypeInt 32 1
+ %one = OpConstant %i32 1
+ %two = OpConstant %u32 2
+ %myconst = OpSpecConstantOp %u32 ShiftRightArithmetic %one %two
+ %voidfn = OpTypeFunction %void
+ %main = OpFunction %void None %voidfn
+ %main_entry = OpLabel
+ %1 = OpCopyObject %u32 %myconst
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+$B1: { # root
+ %1:u32 = spirv.shift_right_arithmetic<u32> 1i, 2u
+ %myconst:u32 = override %1
+}
+
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B2: {
+ %4:u32 = 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.
