github/wgpu
1
1
Fork 0
mirror of https://github.com/gfx-rs/wgpu.git synced 2026-04-22 03:02:01 -04:00
Code Issues Packages Projects Releases Wiki Activity

Support bitwise Or on booleans

Browse Source
This commit is contained in:
João Capucho
2021-12-15 22:30:33 +00:00
committed by Dzmitry Malyshau
parent 8ffd6ba929
commit c64d5eff50
9 changed files with 73 additions and 14 deletions
Download Patch File Download Diff File Expand all files Collapse all files

13
src/back/glsl/mod.rs
View File

@@ -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,
},
};

5
src/back/spv/block.rs
View File

@@ -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,

12
src/valid/expression.rs
View File

@@ -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,

5
tests/in/operators.wgsl
View File

@@ -68,6 +68,10 @@ fn scalar_times_matrix() {
let assertion: mat4x4<f32> = 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();
}

5
tests/out/glsl/operators.main.Compute.glsl
View File

@@ -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;
}

6
tests/out/hlsl/operators.hlsl
View File

@@ -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;
}

6
tests/out/msl/operators.msl
View File

@@ -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;
}

30
tests/out/spv/operators.spvasm
View File

@@ -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

5
tests/out/wgsl/operators.wgsl
View File

@@ -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;
}