[spirv-reader][ir] Convert OpLogicalNot

Convert an `OpLogicalNot` into a unary `not` in WGSL. Bug: 391486707 Change-Id: I5bcc4ff8111b6a7e5ffc55b976a2b1a4f91f175a Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/227076 Reviewed-by: James Price <[email protected]> Commit-Queue: dan sinclair <[email protected]>
google · Feb 20, 2025 · 143208b · 143208b
1 parent 3296fa5
commit 143208b
Show file tree

Hide file tree

Showing 2 changed files with 69 additions and 0 deletions.
diff --git a/src/tint/lang/spirv/reader/parser/logical_test.cc b/src/tint/lang/spirv/reader/parser/logical_test.cc
@@ -622,5 +622,63 @@ INSTANTIATE_TEST_SUITE_P(SpirvParserTest,
                                          SpirvLogicalParam{"Equal", "eq"},
                                          SpirvLogicalParam{"NotEqual", "neq"}));
 
+TEST_F(SpirvParserTest, LogicalNot_Scalar) {
+    EXPECT_IR(R"(
+               OpCapability Shader
+               OpCapability Float16
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+       %void = OpTypeVoid
+       %bool = OpTypeBool
+     %v2bool = OpTypeVector %bool 2
+       %true = OpConstantTrue %bool
+     %v2true = OpConstantComposite %v2bool %true %true
+    %ep_type = OpTypeFunction %void
+       %main = OpFunction %void None %ep_type
+ %main_start = OpLabel
+          %1 = OpLogicalNot %bool %true
+               OpReturn
+               OpFunctionEnd
+)",
+              R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+  $B1: {
+    %2:bool = not true
+    ret
+  }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, LogicalNot_Vector) {
+    EXPECT_IR(R"(
+               OpCapability Shader
+               OpCapability Float16
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+       %void = OpTypeVoid
+       %bool = OpTypeBool
+     %v2bool = OpTypeVector %bool 2
+       %true = OpConstantTrue %bool
+     %v2true = OpConstantComposite %v2bool %true %true
+    %ep_type = OpTypeFunction %void
+       %main = OpFunction %void None %ep_type
+ %main_start = OpLabel
+          %1 = OpLogicalNot %v2bool %v2true
+               OpReturn
+               OpFunctionEnd
+)",
+              R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+  $B1: {
+    %2:vec2<bool> = not vec2<bool>(true)
+    ret
+  }
+}
+)");
+}
+
 }  // namespace
 }  // namespace tint::spirv::reader
diff --git a/src/tint/lang/spirv/reader/parser/parser.cc b/src/tint/lang/spirv/reader/parser/parser.cc
@@ -856,6 +856,9 @@ class Parser {
                 case spv::Op::OpLogicalNotEqual:
                     EmitBinary(inst, core::BinaryOp::kNotEqual);
                     break;
+                case spv::Op::OpLogicalNot:
+                    EmitUnary(inst, core::UnaryOp::kNot);
+                    break;
                 default:
                     TINT_UNIMPLEMENTED()
                         << "unhandled SPIR-V instruction: " << static_cast<uint32_t>(inst.opcode());
@@ -1283,6 +1286,14 @@ class Parser {
         Emit(access, inst.result_id());
     }
 
+    /// @param inst the SPIR-V instruction
+    /// @param op the unary operator to use
+    void EmitUnary(const spvtools::opt::Instruction& inst, core::UnaryOp op) {
+        auto* val = Value(inst.GetSingleWordOperand(2));
+        auto* unary = b_.Unary(op, Type(inst.type_id()), val);
+        Emit(unary, inst.result_id());
+    }
+
     /// @param inst the SPIR-V instruction
     /// @param op the binary operator to use
     void EmitBinary(const spvtools::opt::Instruction& inst, core::BinaryOp op) {