spirv-reader: specop support Select

The translation only supports OpSpecConstantOp with a scalar result type.
ICE in the case of a vector result type.

Fixed: 402727458
Change-Id: I0e7db48831c9a63a2bf4575a78cf3cbb2bcfa784
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/250674
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 72d475b..d0a49af 100644
--- a/src/tint/lang/spirv/reader/parser/parser.cc
+++ b/src/tint/lang/spirv/reader/parser/parser.cc
@@ -462,6 +462,16 @@
                             TINT_ICE() << "can't translate OpSpecConstantOp with VectorShuffle: "
                                           "OpSpecConstantOp maps to a WGSL override declaration, "
                                           "but WGSL overrides must have scalar type";
+                        case spv::Op::OpSelect:
+                            if (!Type(inst.type_id())->IsScalar()) {
+                                TINT_ICE()
+                                    << "can't translate OpSpecConstantOp with Select that returns "
+                                       "a vector: "
+                                       "OpSpecConstantOp maps to a WGSL override declaration, "
+                                       "but WGSL overrides must have scalar type";
+                            }
+                            EmitSpirvBuiltinCall(inst, spirv::BuiltinFn::kSelect, 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 e47a1f7..63ab072 100644
--- a/src/tint/lang/spirv/reader/parser/var_test.cc
+++ b/src/tint/lang/spirv/reader/parser/var_test.cc
@@ -2802,6 +2802,47 @@
 )");
 }
 
+TEST_F(SpirvParserTest, Var_OpSpecConstantOp_Select_CondScalar_ArgsScalar) {
+    // In SPIR-V the arg types must be the same as the result type.
+    // WGSL only supports overrides with scalar type, so only test the
+    // case where the result type is scalar.
+    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
+       %true = OpConstantTrue %bool
+      %false = OpConstantFalse %bool
+        %one = OpConstant %i32 1
+        %two = OpConstant %i32 2
+    %myconst = OpSpecConstantOp %i32 Select %true %one %two
+     %voidfn = OpTypeFunction %void
+       %main = OpFunction %void None %voidfn
+ %main_entry = OpLabel
+          %1 = OpCopyObject %i32 %myconst
+               OpReturn
+               OpFunctionEnd
+)",
+              R"(
+$B1: {  # root
+  %1:i32 = spirv.select true, 1i, 2i
+  %myconst:i32 = override %1
+}
+
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+  $B2: {
+    %4:i32 = 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.