From c64d5eff5006bc8cac9991efa4240536d5787c6e Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Jo=C3=A3o=20Capucho?= Date: Wed, 15 Dec 2021 22:30:33 +0000 Subject: [PATCH] Support bitwise Or on booleans --- src/back/glsl/mod.rs | 13 +++++++++- src/back/spv/block.rs | 5 +++- src/valid/expression.rs | 12 ++++++++- tests/in/operators.wgsl | 5 ++++ tests/out/glsl/operators.main.Compute.glsl | 5 ++++ tests/out/hlsl/operators.hlsl | 6 +++++ tests/out/msl/operators.msl | 6 +++++ tests/out/spv/operators.spvasm | 30 ++++++++++++++-------- tests/out/wgsl/operators.wgsl | 5 ++++ 9 files changed, 73 insertions(+), 14 deletions(-) diff --git a/src/back/glsl/mod.rs b/src/back/glsl/mod.rs index 082fa685c5..bb0f4c4386 100644 --- a/src/back/glsl/mod.rs +++ b/src/back/glsl/mod.rs @@ -2239,7 +2239,11 @@ impl<'a, W: Write> Writer<'a, W> { // comparison operations on vectors as they are implemented with // builtin functions. // Once again we wrap everything in parentheses to avoid precedence issues - Expression::Binary { op, left, right } => { + Expression::Binary { + mut op, + left, + right, + } => { // Holds `Some(function_name)` if the binary operation is // implemented as a function call use crate::{BinaryOperator as Bo, ScalarKind as Sk, TypeInner as Ti}; @@ -2276,6 +2280,13 @@ impl<'a, W: Write> Writer<'a, W> { Bo::Modulo => BinaryOperation::Modulo, _ => BinaryOperation::Other, }, + (Some(Sk::Bool), Some(Sk::Bool)) => match op { + Bo::InclusiveOr => { + op = crate::BinaryOperator::LogicalOr; + BinaryOperation::Other + } + _ => BinaryOperation::Other, + }, _ => BinaryOperation::Other, }, }; diff --git a/src/back/spv/block.rs b/src/back/spv/block.rs index f0676f0d40..ca05da4d36 100644 --- a/src/back/spv/block.rs +++ b/src/back/spv/block.rs @@ -435,7 +435,10 @@ impl<'w> BlockContext<'w> { }, crate::BinaryOperator::And => spirv::Op::BitwiseAnd, crate::BinaryOperator::ExclusiveOr => spirv::Op::BitwiseXor, - crate::BinaryOperator::InclusiveOr => spirv::Op::BitwiseOr, + crate::BinaryOperator::InclusiveOr => match left_ty_inner.scalar_kind() { + Some(crate::ScalarKind::Bool) => spirv::Op::LogicalOr, + _ => spirv::Op::BitwiseOr, + }, crate::BinaryOperator::LogicalAnd => spirv::Op::LogicalAnd, crate::BinaryOperator::LogicalOr => spirv::Op::LogicalOr, crate::BinaryOperator::ShiftLeft => spirv::Op::ShiftLeftLogical, diff --git a/src/valid/expression.rs b/src/valid/expression.rs index a9dab24cd6..023bf018c8 100644 --- a/src/valid/expression.rs +++ b/src/valid/expression.rs @@ -739,7 +739,17 @@ impl super::Validator { false } }, - Bo::And | Bo::ExclusiveOr | Bo::InclusiveOr => match *left_inner { + Bo::And | Bo::InclusiveOr => match *left_inner { + Ti::Scalar { kind, .. } | Ti::Vector { kind, .. } => match kind { + Sk::Bool | Sk::Sint | Sk::Uint => left_inner == right_inner, + Sk::Float => false, + }, + ref other => { + log::error!("Op {:?} left type {:?}", op, other); + false + } + }, + Bo::ExclusiveOr => match *left_inner { Ti::Scalar { kind, .. } | Ti::Vector { kind, .. } => match kind { Sk::Sint | Sk::Uint => left_inner == right_inner, Sk::Bool | Sk::Float => false, diff --git a/tests/in/operators.wgsl b/tests/in/operators.wgsl index 673bd55cde..0c727eae2e 100644 --- a/tests/in/operators.wgsl +++ b/tests/in/operators.wgsl @@ -68,6 +68,10 @@ fn scalar_times_matrix() { let assertion: mat4x4 = 2.0 * model; } +fn binary() { + let a = true | false; +} + [[stage(compute), workgroup_size(1)]] fn main() { let a = builtins(); @@ -77,4 +81,5 @@ fn main() { let e = constructors(); modulo(); scalar_times_matrix(); + binary(); } diff --git a/tests/out/glsl/operators.main.Compute.glsl b/tests/out/glsl/operators.main.Compute.glsl index 9a9fa5c899..eb9dd8490e 100644 --- a/tests/out/glsl/operators.main.Compute.glsl +++ b/tests/out/glsl/operators.main.Compute.glsl @@ -60,6 +60,10 @@ void scalar_times_matrix() { mat4x4 assertion = (2.0 * model); } +void binary() { + bool a_2 = (true || false); +} + void main() { vec4 _e4 = builtins(); vec4 _e5 = splat(); @@ -68,6 +72,7 @@ void main() { float _e9 = constructors(); modulo(); scalar_times_matrix(); + binary(); return; } diff --git a/tests/out/hlsl/operators.hlsl b/tests/out/hlsl/operators.hlsl index a359689f68..a3f7f8a08c 100644 --- a/tests/out/hlsl/operators.hlsl +++ b/tests/out/hlsl/operators.hlsl @@ -73,6 +73,11 @@ void scalar_times_matrix() float4x4 assertion = mul(model, 2.0); } +void binary() +{ + bool a_2 = (true | false); +} + [numthreads(1, 1, 1)] void main() { @@ -83,5 +88,6 @@ void main() const float _e9 = constructors(); modulo(); scalar_times_matrix(); + binary(); return; } diff --git a/tests/out/msl/operators.msl b/tests/out/msl/operators.msl index f801979895..97b7490feb 100644 --- a/tests/out/msl/operators.msl +++ b/tests/out/msl/operators.msl @@ -69,6 +69,11 @@ void scalar_times_matrix( metal::float4x4 assertion = 2.0 * model; } +void binary( +) { + bool a_2 = true | false; +} + kernel void main_( ) { metal::float4 _e4 = builtins(); @@ -78,5 +83,6 @@ kernel void main_( float _e9 = constructors(); modulo(); scalar_times_matrix(); + binary(); return; } diff --git a/tests/out/spv/operators.spvasm b/tests/out/spv/operators.spvasm index bc931da334..081bfbb23d 100644 --- a/tests/out/spv/operators.spvasm +++ b/tests/out/spv/operators.spvasm @@ -1,12 +1,12 @@ ; SPIR-V ; Version: 1.1 ; Generator: rspirv -; Bound: 139 +; Bound: 144 OpCapability Shader %1 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 -OpEntryPoint GLCompute %129 "main" -OpExecutionMode %129 LocalSize 1 1 1 +OpEntryPoint GLCompute %133 "main" +OpExecutionMode %133 LocalSize 1 1 1 OpMemberDecorate %23 0 Offset 0 OpMemberDecorate %23 1 Offset 16 %2 = OpTypeVoid @@ -165,13 +165,21 @@ OpFunctionEnd %128 = OpLabel OpBranch %130 %130 = OpLabel -%131 = OpFunctionCall %19 %30 -%132 = OpFunctionCall %19 %55 -%133 = OpFunctionCall %8 %72 -%134 = OpVectorShuffle %22 %25 %25 0 1 2 -%135 = OpFunctionCall %22 %83 %134 -%136 = OpFunctionCall %4 %95 -%137 = OpFunctionCall %2 %107 -%138 = OpFunctionCall %2 %120 +%131 = OpLogicalOr %10 %9 %12 +OpReturn +OpFunctionEnd +%133 = OpFunction %2 None %108 +%132 = OpLabel +OpBranch %134 +%134 = OpLabel +%135 = OpFunctionCall %19 %30 +%136 = OpFunctionCall %19 %55 +%137 = OpFunctionCall %8 %72 +%138 = OpVectorShuffle %22 %25 %25 0 1 2 +%139 = OpFunctionCall %22 %83 %138 +%140 = OpFunctionCall %4 %95 +%141 = OpFunctionCall %2 %107 +%142 = OpFunctionCall %2 %120 +%143 = OpFunctionCall %2 %129 OpReturn OpFunctionEnd \ No newline at end of file diff --git a/tests/out/wgsl/operators.wgsl b/tests/out/wgsl/operators.wgsl index de5ec99f8e..37dd0bed93 100644 --- a/tests/out/wgsl/operators.wgsl +++ b/tests/out/wgsl/operators.wgsl @@ -58,6 +58,10 @@ fn scalar_times_matrix() { let assertion = (2.0 * model); } +fn binary() { + let a_2 = (true | false); +} + [[stage(compute), workgroup_size(1, 1, 1)]] fn main() { let _e4 = builtins(); @@ -67,5 +71,6 @@ fn main() { let _e9 = constructors(); modulo(); scalar_times_matrix(); + binary(); return; }