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>

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 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;
                     }

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

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.