spirv-reader: specop support bitwise And, Or, Xor
Fixed: 402726781 402726328 402726435
Change-Id: Iab5d1d44ef5e7aa766c7f45ce8ebaf68e42ca0d5
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/250475
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 cc904ca..b5fdb87 100644
--- a/src/tint/lang/spirv/reader/parser/parser.cc
+++ b/src/tint/lang/spirv/reader/parser/parser.cc
@@ -346,6 +346,15 @@
}
switch (static_cast<spv::Op>(op)) {
+ case spv::Op::OpBitwiseAnd:
+ EmitSpirvExplicitBuiltinCall(inst, spirv::BuiltinFn::kBitwiseAnd, 3);
+ break;
+ case spv::Op::OpBitwiseOr:
+ EmitSpirvExplicitBuiltinCall(inst, spirv::BuiltinFn::kBitwiseOr, 3);
+ break;
+ case spv::Op::OpBitwiseXor:
+ EmitSpirvExplicitBuiltinCall(inst, spirv::BuiltinFn::kBitwiseXor, 3);
+ break;
case spv::Op::OpLogicalAnd:
EmitBinary(inst, core::BinaryOp::kAnd, 3);
break;
diff --git a/src/tint/lang/spirv/reader/parser/var_test.cc b/src/tint/lang/spirv/reader/parser/var_test.cc
index 0b23840..11f570e 100644
--- a/src/tint/lang/spirv/reader/parser/var_test.cc
+++ b/src/tint/lang/spirv/reader/parser/var_test.cc
@@ -2195,6 +2195,111 @@
)");
}
+TEST_F(SpirvParserTest, Var_OpSpecConstantOp_BitwiseAnd) {
+ 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 BitwiseAnd %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.bitwise_and<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_BitwiseOr) {
+ 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 BitwiseOr %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.bitwise_or<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_BitwiseXor) {
+ 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 BitwiseXor %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.bitwise_xor<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.
