tint: Fix MSL generation of '&' and '|' with booleans

The bitwise-and and bitwise-or binary operators on booleans result in an integer.
Explicitly cast this back to a boolean.

Fixed: tint:1540
Fixed: tint:1541
Change-Id: I395176f291e6080c88b8cff18e14ed6cd1234074
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/90501
Commit-Queue: Ben Clayton <bclayton@google.com>
Reviewed-by: James Price <jrprice@google.com>

src/tint/writer/msl/generator_impl.cc

@@ -513,6 +513,22 @@ bool GeneratorImpl::EmitBinary(std::ostream& out, const ast::BinaryExpression* e
         return true;
     }
 
+    // Handle '&' and '|' of booleans.
+    if ((expr->IsAnd() || expr->IsOr()) && lhs_type->Is<sem::Bool>()) {
+        out << "bool";
+        ScopedParen sp(out);
+        if (!EmitExpression(out, expr->lhs)) {
+            return false;
+        }
+        if (!emit_op()) {
+            return false;
+        }
+        if (!EmitExpression(out, expr->rhs)) {
+            return false;
+        }
+        return true;
+    }
+
     // Emit as usual
     ScopedParen sp(out);
     if (!EmitExpression(out, expr->lhs)) {

src/tint/writer/msl/generator_impl_binary_test.cc

@@ -171,5 +171,31 @@ TEST_F(MslBinaryTest, ModVec3F32) {
     EXPECT_EQ(out.str(), "fmod(left, right)");
 }
 
+TEST_F(MslBinaryTest, BoolAnd) {
+    auto* left = Var("left", nullptr, Expr(true));
+    auto* right = Var("right", nullptr, Expr(false));
+    auto* expr = create<ast::BinaryExpression>(ast::BinaryOp::kAnd, Expr(left), Expr(right));
+    WrapInFunction(left, right, expr);
+
+    GeneratorImpl& gen = Build();
+
+    std::stringstream out;
+    ASSERT_TRUE(gen.EmitExpression(out, expr)) << gen.error();
+    EXPECT_EQ(out.str(), "bool(left & right)");
+}
+
+TEST_F(MslBinaryTest, BoolOr) {
+    auto* left = Var("left", nullptr, Expr(true));
+    auto* right = Var("right", nullptr, Expr(false));
+    auto* expr = create<ast::BinaryExpression>(ast::BinaryOp::kOr, Expr(left), Expr(right));
+    WrapInFunction(left, right, expr);
+
+    GeneratorImpl& gen = Build();
+
+    std::stringstream out;
+    ASSERT_TRUE(gen.EmitExpression(out, expr)) << gen.error();
+    EXPECT_EQ(out.str(), "bool(left | right)");
+}
+
 }  // namespace
 }  // namespace tint::writer::msl

test/tint/bug/tint/1540.wgsl

@@ -0,0 +1,10 @@
+struct S {
+    e: bool,
+}
+
+@stage(compute)
+@workgroup_size(1)
+fn main() {
+    var b : bool;
+    var v = S(true & b);
+}

test/tint/bug/tint/1540.wgsl.expected.glsl

@@ -0,0 +1,16 @@
+#version 310 es
+
+struct S {
+  bool e;
+};
+
+void tint_symbol() {
+  bool b = false;
+  S v = S(bool(uint(true) & uint(b)));
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  tint_symbol();
+  return;
+}

test/tint/bug/tint/1540.wgsl.expected.hlsl

@@ -0,0 +1,10 @@
+struct S {
+  bool e;
+};
+
+[numthreads(1, 1, 1)]
+void main() {
+  bool b = false;
+  S v = {(true & b)};
+  return;
+}

test/tint/bug/tint/1540.wgsl.expected.msl

@@ -0,0 +1,13 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct S {
+  bool e;
+};
+
+kernel void tint_symbol() {
+  bool b = false;
+  S v = {.e=bool(true & b)};
+  return;
+}
+

test/tint/bug/tint/1540.wgsl.expected.spvasm

@@ -0,0 +1,34 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 17
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %main "main"
+               OpName %b "b"
+               OpName %S "S"
+               OpMemberName %S 0 "e"
+               OpName %v "v"
+               OpMemberDecorate %S 0 Offset 0
+       %void = OpTypeVoid
+          %1 = OpTypeFunction %void
+       %bool = OpTypeBool
+%_ptr_Function_bool = OpTypePointer Function %bool
+          %8 = OpConstantNull %bool
+          %S = OpTypeStruct %bool
+       %true = OpConstantTrue %bool
+%_ptr_Function_S = OpTypePointer Function %S
+         %16 = OpConstantNull %S
+       %main = OpFunction %void None %1
+          %4 = OpLabel
+          %b = OpVariable %_ptr_Function_bool Function %8
+          %v = OpVariable %_ptr_Function_S Function %16
+         %11 = OpLoad %bool %b
+         %12 = OpLogicalAnd %bool %true %11
+         %13 = OpCompositeConstruct %S %12
+               OpStore %v %13
+               OpReturn
+               OpFunctionEnd

test/tint/bug/tint/1540.wgsl.expected.wgsl

@@ -0,0 +1,9 @@
+struct S {
+  e : bool,
+}
+
+@stage(compute) @workgroup_size(1)
+fn main() {
+  var b : bool;
+  var v = S((true & b));
+}

test/tint/bug/tint/1541.wgsl

@@ -0,0 +1,5 @@
+@stage(compute)
+@workgroup_size(1)
+fn main() {
+    var v = select(true & true, true, false);
+}

test/tint/bug/tint/1541.wgsl.expected.glsl

@@ -0,0 +1,11 @@
+#version 310 es
+
+void tint_symbol() {
+  bool v = (false ? true : bool(uint(true) & uint(true)));
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  tint_symbol();
+  return;
+}

test/tint/bug/tint/1541.wgsl.expected.hlsl

@@ -0,0 +1,5 @@
+[numthreads(1, 1, 1)]
+void main() {
+  bool v = (false ? true : (true & true));
+  return;
+}

test/tint/bug/tint/1541.wgsl.expected.msl

@@ -0,0 +1,8 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void tint_symbol() {
+  bool v = select(bool(true & true), true, false);
+  return;
+}
+

test/tint/bug/tint/1541.wgsl.expected.spvasm

@@ -0,0 +1,26 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 13
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %main "main"
+               OpName %v "v"
+       %void = OpTypeVoid
+          %1 = OpTypeFunction %void
+       %bool = OpTypeBool
+      %false = OpConstantFalse %bool
+       %true = OpConstantTrue %bool
+%_ptr_Function_bool = OpTypePointer Function %bool
+         %12 = OpConstantNull %bool
+       %main = OpFunction %void None %1
+          %4 = OpLabel
+          %v = OpVariable %_ptr_Function_bool Function %12
+          %9 = OpLogicalAnd %bool %true %true
+          %5 = OpSelect %bool %false %true %9
+               OpStore %v %5
+               OpReturn
+               OpFunctionEnd

test/tint/bug/tint/1541.wgsl.expected.wgsl

@@ -0,0 +1,4 @@
+@stage(compute) @workgroup_size(1)
+fn main() {
+  var v = select((true & true), true, false);
+}

test/tint/expressions/binary/bit-and/scalar-scalar/bool.wgsl

@@ -0,0 +1,6 @@
+@stage(compute) @workgroup_size(1)
+fn f() {
+    let a = true;
+    let b = false;
+    let r : bool = a & b;
+}

test/tint/expressions/binary/bit-and/scalar-scalar/bool.wgsl.expected.glsl

@@ -0,0 +1,11 @@
+#version 310 es
+
+void f() {
+  bool r = bool(uint(true) & uint(false));
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  f();
+  return;
+}

test/tint/expressions/binary/bit-and/scalar-scalar/bool.wgsl.expected.hlsl

@@ -0,0 +1,5 @@
+[numthreads(1, 1, 1)]
+void f() {
+  const bool r = (true & false);
+  return;
+}

test/tint/expressions/binary/bit-and/scalar-scalar/bool.wgsl.expected.msl

@@ -0,0 +1,10 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void f() {
+  bool const a = true;
+  bool const b = false;
+  bool const r = bool(a & b);
+  return;
+}
+

test/tint/expressions/binary/bit-and/scalar-scalar/bool.wgsl.expected.spvasm

@@ -0,0 +1,20 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 9
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %f "f"
+               OpExecutionMode %f LocalSize 1 1 1
+               OpName %f "f"
+       %void = OpTypeVoid
+          %1 = OpTypeFunction %void
+       %bool = OpTypeBool
+       %true = OpConstantTrue %bool
+      %false = OpConstantFalse %bool
+          %f = OpFunction %void None %1
+          %4 = OpLabel
+          %8 = OpLogicalAnd %bool %true %false
+               OpReturn
+               OpFunctionEnd

test/tint/expressions/binary/bit-and/scalar-scalar/bool.wgsl.expected.wgsl

@@ -0,0 +1,6 @@
+@stage(compute) @workgroup_size(1)
+fn f() {
+  let a = true;
+  let b = false;
+  let r : bool = (a & b);
+}

test/tint/expressions/binary/bit-and/vec3-vec3/bool.wgsl

@@ -0,0 +1,6 @@
+@stage(compute) @workgroup_size(1)
+fn f() {
+    let a = vec3<bool>(true, true, false);
+    let b = vec3<bool>(true, false, true);
+    let r : vec3<bool> = a & b;
+}

test/tint/expressions/binary/bit-and/vec3-vec3/bool.wgsl.expected.glsl

@@ -0,0 +1,13 @@
+#version 310 es
+
+void f() {
+  bvec3 a = bvec3(true, true, false);
+  bvec3 b = bvec3(true, false, true);
+  bvec3 r = bvec3(uvec3(a) & uvec3(b));
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  f();
+  return;
+}

test/tint/expressions/binary/bit-and/vec3-vec3/bool.wgsl.expected.hlsl

@@ -0,0 +1,7 @@
+[numthreads(1, 1, 1)]
+void f() {
+  const bool3 a = bool3(true, true, false);
+  const bool3 b = bool3(true, false, true);
+  const bool3 r = (a & b);
+  return;
+}

test/tint/expressions/binary/bit-and/vec3-vec3/bool.wgsl.expected.msl

@@ -0,0 +1,10 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void f() {
+  bool3 const a = bool3(true, true, false);
+  bool3 const b = bool3(true, false, true);
+  bool3 const r = (a & b);
+  return;
+}
+

test/tint/expressions/binary/bit-and/vec3-vec3/bool.wgsl.expected.spvasm

@@ -0,0 +1,23 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 12
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %f "f"
+               OpExecutionMode %f LocalSize 1 1 1
+               OpName %f "f"
+       %void = OpTypeVoid
+          %1 = OpTypeFunction %void
+       %bool = OpTypeBool
+     %v3bool = OpTypeVector %bool 3
+       %true = OpConstantTrue %bool
+      %false = OpConstantFalse %bool
+          %9 = OpConstantComposite %v3bool %true %true %false
+         %10 = OpConstantComposite %v3bool %true %false %true
+          %f = OpFunction %void None %1
+          %4 = OpLabel
+         %11 = OpLogicalAnd %v3bool %9 %10
+               OpReturn
+               OpFunctionEnd

test/tint/expressions/binary/bit-and/vec3-vec3/bool.wgsl.expected.wgsl

@@ -0,0 +1,6 @@
+@stage(compute) @workgroup_size(1)
+fn f() {
+  let a = vec3<bool>(true, true, false);
+  let b = vec3<bool>(true, false, true);
+  let r : vec3<bool> = (a & b);
+}