diff --git a/src/tint/writer/msl/generator_impl.cc b/src/tint/writer/msl/generator_impl.cc index 2f87cb257a..743e03b58a 100644 --- a/src/tint/writer/msl/generator_impl.cc +++ b/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()) { + 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)) { diff --git a/src/tint/writer/msl/generator_impl_binary_test.cc b/src/tint/writer/msl/generator_impl_binary_test.cc index 5262b97eec..02daae2550 100644 --- a/src/tint/writer/msl/generator_impl_binary_test.cc +++ b/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::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::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 diff --git a/test/tint/bug/tint/1540.wgsl b/test/tint/bug/tint/1540.wgsl new file mode 100644 index 0000000000..95c51fb1d2 --- /dev/null +++ b/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); +} diff --git a/test/tint/bug/tint/1540.wgsl.expected.glsl b/test/tint/bug/tint/1540.wgsl.expected.glsl new file mode 100644 index 0000000000..5e3a4100f0 --- /dev/null +++ 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; +} diff --git a/test/tint/bug/tint/1540.wgsl.expected.hlsl b/test/tint/bug/tint/1540.wgsl.expected.hlsl new file mode 100644 index 0000000000..396172fd90 --- /dev/null +++ b/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; +} diff --git a/test/tint/bug/tint/1540.wgsl.expected.msl b/test/tint/bug/tint/1540.wgsl.expected.msl new file mode 100644 index 0000000000..387b86be6e --- /dev/null +++ b/test/tint/bug/tint/1540.wgsl.expected.msl @@ -0,0 +1,13 @@ +#include + +using namespace metal; +struct S { + bool e; +}; + +kernel void tint_symbol() { + bool b = false; + S v = {.e=bool(true & b)}; + return; +} + diff --git a/test/tint/bug/tint/1540.wgsl.expected.spvasm b/test/tint/bug/tint/1540.wgsl.expected.spvasm new file mode 100644 index 0000000000..2a2d626882 --- /dev/null +++ b/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 diff --git a/test/tint/bug/tint/1540.wgsl.expected.wgsl b/test/tint/bug/tint/1540.wgsl.expected.wgsl new file mode 100644 index 0000000000..50f6bf3ddf --- /dev/null +++ b/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)); +} diff --git a/test/tint/bug/tint/1541.wgsl b/test/tint/bug/tint/1541.wgsl new file mode 100644 index 0000000000..44d04f04ec --- /dev/null +++ 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); +} diff --git a/test/tint/bug/tint/1541.wgsl.expected.glsl b/test/tint/bug/tint/1541.wgsl.expected.glsl new file mode 100644 index 0000000000..260a8e522e --- /dev/null +++ b/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; +} diff --git a/test/tint/bug/tint/1541.wgsl.expected.hlsl b/test/tint/bug/tint/1541.wgsl.expected.hlsl new file mode 100644 index 0000000000..3eb2b64c3b --- /dev/null +++ b/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; +} diff --git a/test/tint/bug/tint/1541.wgsl.expected.msl b/test/tint/bug/tint/1541.wgsl.expected.msl new file mode 100644 index 0000000000..d339bd2621 --- /dev/null +++ b/test/tint/bug/tint/1541.wgsl.expected.msl @@ -0,0 +1,8 @@ +#include + +using namespace metal; +kernel void tint_symbol() { + bool v = select(bool(true & true), true, false); + return; +} + diff --git a/test/tint/bug/tint/1541.wgsl.expected.spvasm b/test/tint/bug/tint/1541.wgsl.expected.spvasm new file mode 100644 index 0000000000..415f251f31 --- /dev/null +++ b/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 diff --git a/test/tint/bug/tint/1541.wgsl.expected.wgsl b/test/tint/bug/tint/1541.wgsl.expected.wgsl new file mode 100644 index 0000000000..ed2d91b708 --- /dev/null +++ b/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); +} diff --git a/test/tint/expressions/binary/bit-and/scalar-scalar/bool.wgsl b/test/tint/expressions/binary/bit-and/scalar-scalar/bool.wgsl new file mode 100644 index 0000000000..ffe06b24de --- /dev/null +++ b/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; +} diff --git a/test/tint/expressions/binary/bit-and/scalar-scalar/bool.wgsl.expected.glsl b/test/tint/expressions/binary/bit-and/scalar-scalar/bool.wgsl.expected.glsl new file mode 100644 index 0000000000..f82b0dc0ae --- /dev/null +++ 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; +} diff --git a/test/tint/expressions/binary/bit-and/scalar-scalar/bool.wgsl.expected.hlsl b/test/tint/expressions/binary/bit-and/scalar-scalar/bool.wgsl.expected.hlsl new file mode 100644 index 0000000000..17423e395b --- /dev/null +++ b/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; +} diff --git a/test/tint/expressions/binary/bit-and/scalar-scalar/bool.wgsl.expected.msl b/test/tint/expressions/binary/bit-and/scalar-scalar/bool.wgsl.expected.msl new file mode 100644 index 0000000000..48bc530e57 --- /dev/null +++ b/test/tint/expressions/binary/bit-and/scalar-scalar/bool.wgsl.expected.msl @@ -0,0 +1,10 @@ +#include + +using namespace metal; +kernel void f() { + bool const a = true; + bool const b = false; + bool const r = bool(a & b); + return; +} + diff --git a/test/tint/expressions/binary/bit-and/scalar-scalar/bool.wgsl.expected.spvasm b/test/tint/expressions/binary/bit-and/scalar-scalar/bool.wgsl.expected.spvasm new file mode 100644 index 0000000000..a3279258fb --- /dev/null +++ b/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 diff --git a/test/tint/expressions/binary/bit-and/scalar-scalar/bool.wgsl.expected.wgsl b/test/tint/expressions/binary/bit-and/scalar-scalar/bool.wgsl.expected.wgsl new file mode 100644 index 0000000000..0c1fb6bf5c --- /dev/null +++ b/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); +} diff --git a/test/tint/expressions/binary/bit-and/vec3-vec3/bool.wgsl b/test/tint/expressions/binary/bit-and/vec3-vec3/bool.wgsl new file mode 100644 index 0000000000..a463c0ed0d --- /dev/null +++ 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(true, true, false); + let b = vec3(true, false, true); + let r : vec3 = a & b; +} diff --git a/test/tint/expressions/binary/bit-and/vec3-vec3/bool.wgsl.expected.glsl b/test/tint/expressions/binary/bit-and/vec3-vec3/bool.wgsl.expected.glsl new file mode 100644 index 0000000000..2c08551b7b --- /dev/null +++ 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; +} diff --git a/test/tint/expressions/binary/bit-and/vec3-vec3/bool.wgsl.expected.hlsl b/test/tint/expressions/binary/bit-and/vec3-vec3/bool.wgsl.expected.hlsl new file mode 100644 index 0000000000..0975dc9df6 --- /dev/null +++ b/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; +} diff --git a/test/tint/expressions/binary/bit-and/vec3-vec3/bool.wgsl.expected.msl b/test/tint/expressions/binary/bit-and/vec3-vec3/bool.wgsl.expected.msl new file mode 100644 index 0000000000..f8eed1d343 --- /dev/null +++ b/test/tint/expressions/binary/bit-and/vec3-vec3/bool.wgsl.expected.msl @@ -0,0 +1,10 @@ +#include + +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; +} + diff --git a/test/tint/expressions/binary/bit-and/vec3-vec3/bool.wgsl.expected.spvasm b/test/tint/expressions/binary/bit-and/vec3-vec3/bool.wgsl.expected.spvasm new file mode 100644 index 0000000000..6c965db038 --- /dev/null +++ b/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 diff --git a/test/tint/expressions/binary/bit-and/vec3-vec3/bool.wgsl.expected.wgsl b/test/tint/expressions/binary/bit-and/vec3-vec3/bool.wgsl.expected.wgsl new file mode 100644 index 0000000000..6610d13dae --- /dev/null +++ b/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(true, true, false); + let b = vec3(true, false, true); + let r : vec3 = (a & b); +}