From 2b475ecc9626107ef9fb4f2acb6d64188407e3fa Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Jo=C3=A3o=20Capucho?= Date: Tue, 6 Jul 2021 10:43:22 +0100 Subject: [PATCH] [spv-out] Support object-wise select --- src/back/spv/block.rs | 35 ++++++- tests/in/operators.wgsl | 9 ++ tests/out/glsl/operators.main.Compute.glsl | 8 ++ tests/out/glsl/operators.main.Fragment.glsl | 35 +++++++ tests/out/hlsl/operators.hlsl | 9 ++ tests/out/msl/operators.msl | 9 ++ tests/out/spv/operators.spvasm | 102 +++++++++++--------- tests/out/wgsl/operators.wgsl | 8 ++ 8 files changed, 170 insertions(+), 45 deletions(-) create mode 100644 tests/out/glsl/operators.main.Fragment.glsl diff --git a/src/back/spv/block.rs b/src/back/spv/block.rs index 7de2359360..f45203076d 100644 --- a/src/back/spv/block.rs +++ b/src/back/spv/block.rs @@ -933,10 +933,43 @@ impl<'w> BlockContext<'w> { reject, } => { let id = self.gen_id(); - let condition_id = self.cached(condition); + let mut condition_id = self.cached(condition); let accept_id = self.cached(accept); let reject_id = self.cached(reject); + let condition_ty = self.fun_info[condition] + .ty + .inner_with(&self.ir_module.types); + let object_ty = self.fun_info[accept].ty.inner_with(&self.ir_module.types); + + if let ( + &crate::TypeInner::Scalar { + kind: crate::ScalarKind::Bool, + width, + }, + &crate::TypeInner::Vector { size, .. }, + ) = (condition_ty, object_ty) + { + self.temp_list.clear(); + self.temp_list.resize(size as usize, condition_id); + + let bool_vector_type_id = + self.get_type_id(LookupType::Local(LocalType::Value { + vector_size: Some(size), + kind: crate::ScalarKind::Bool, + width, + pointer_class: None, + }))?; + + let id = self.gen_id(); + block.body.push(Instruction::composite_construct( + bool_vector_type_id, + id, + &self.temp_list, + )); + condition_id = id + } + let instruction = Instruction::select(result_type_id, id, condition_id, accept_id, reject_id); block.body.push(instruction); diff --git a/tests/in/operators.wgsl b/tests/in/operators.wgsl index cb6da5df09..9bbf9ba675 100644 --- a/tests/in/operators.wgsl +++ b/tests/in/operators.wgsl @@ -9,8 +9,17 @@ fn unary() -> i32 { if (!true) { return a; } else { return ~a; }; } +fn selection() -> vec4 { + let vector1 = vec4(1.0); + let vector2 = vec4(1.0); + let condition = true; + let a = select(0, 1, condition); + return select(vector1, vector2, condition); +} + [[stage(compute), workgroup_size(1)]] fn main() { let a = splat(); let b = unary(); + let c = selection(); } diff --git a/tests/out/glsl/operators.main.Compute.glsl b/tests/out/glsl/operators.main.Compute.glsl index caad84e6aa..232cd4c067 100644 --- a/tests/out/glsl/operators.main.Compute.glsl +++ b/tests/out/glsl/operators.main.Compute.glsl @@ -19,9 +19,17 @@ int unary() { } } +vec4 selection() { + vec4 vector1_ = vec4(1.0); + vec4 vector2_ = vec4(1.0); + int a = (true ? 0 : 1); + return (true ? vector1_ : vector2_); +} + void main() { vec4 _expr0 = splat(); int _expr1 = unary(); + vec4 _expr2 = selection(); return; } diff --git a/tests/out/glsl/operators.main.Fragment.glsl b/tests/out/glsl/operators.main.Fragment.glsl new file mode 100644 index 0000000000..0c1eaa482f --- /dev/null +++ b/tests/out/glsl/operators.main.Fragment.glsl @@ -0,0 +1,35 @@ +#version 310 es + +precision highp float; + +layout(location = 0) out vec4 _fs2p_location0; + +vec4 splat() { + vec2 a = (((vec2(1.0) + vec2(2.0)) - vec2(3.0)) / vec2(4.0)); + ivec4 b = (ivec4(5) % ivec4(2)); + return (a.xyxy + vec4(b)); +} + +int unary() { + if ((! true)) { + return 1; + } else { + return (~ 1); + } +} + +int unary1() { + if ((! true)) { + return 1; + } else { + return (~ 1); + } +} + +void main() { + vec4 vector1_ = vec4(1.0); + vec4 vector2_ = vec4(1.0); + _fs2p_location0 = (true ? vector1_ : vector2_); + return; +} + diff --git a/tests/out/hlsl/operators.hlsl b/tests/out/hlsl/operators.hlsl index 2448e11c46..b582a44520 100644 --- a/tests/out/hlsl/operators.hlsl +++ b/tests/out/hlsl/operators.hlsl @@ -14,10 +14,19 @@ int unary() } } +float4 selection() +{ + float4 vector1_ = float4(1.0.xxxx); + float4 vector2_ = float4(1.0.xxxx); + int a = (true ? 0 : 1); + return (true ? vector1_ : vector2_); +} + [numthreads(1, 1, 1)] void main() { const float4 _e0 = splat(); const int _e1 = unary(); + const float4 _e2 = selection(); return; } diff --git a/tests/out/msl/operators.msl b/tests/out/msl/operators.msl index 3749cb47c4..9fd7f3b4a7 100644 --- a/tests/out/msl/operators.msl +++ b/tests/out/msl/operators.msl @@ -18,9 +18,18 @@ int unary( } } +metal::float4 selection( +) { + metal::float4 vector1_ = metal::float4(1.0); + metal::float4 vector2_ = metal::float4(1.0); + int a = true ? 0 : 1; + return true ? vector1_ : vector2_; +} + kernel void main1( ) { metal::float4 _e0 = splat(); int _e1 = unary(); + metal::float4 _e2 = selection(); return; } diff --git a/tests/out/spv/operators.spvasm b/tests/out/spv/operators.spvasm index aefe1c99fd..f845036094 100644 --- a/tests/out/spv/operators.spvasm +++ b/tests/out/spv/operators.spvasm @@ -1,12 +1,12 @@ ; SPIR-V ; Version: 1.0 ; Generator: rspirv -; Bound: 50 +; Bound: 61 OpCapability Shader %1 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 -OpEntryPoint GLCompute %45 "main" -OpExecutionMode %45 LocalSize 1 1 1 +OpEntryPoint GLCompute %55 "main" +OpExecutionMode %55 LocalSize 1 1 1 %2 = OpTypeVoid %4 = OpTypeFloat 32 %3 = OpConstant %4 1.0 @@ -19,52 +19,66 @@ OpExecutionMode %45 LocalSize 1 1 1 %11 = OpConstant %9 1 %13 = OpTypeBool %12 = OpConstantTrue %13 -%14 = OpTypeVector %4 4 -%17 = OpTypeFunction %14 -%19 = OpTypeVector %4 2 -%27 = OpTypeVector %9 4 -%36 = OpTypeFunction %9 -%43 = OpConstantNull %9 -%46 = OpTypeFunction %2 -%16 = OpFunction %14 None %17 -%15 = OpLabel -OpBranch %18 -%18 = OpLabel -%20 = OpCompositeConstruct %19 %5 %5 -%21 = OpCompositeConstruct %19 %3 %3 -%22 = OpFAdd %19 %21 %20 -%23 = OpCompositeConstruct %19 %6 %6 -%24 = OpFSub %19 %22 %23 -%25 = OpCompositeConstruct %19 %7 %7 -%26 = OpFDiv %19 %24 %25 -%28 = OpCompositeConstruct %27 %8 %8 %8 %8 -%29 = OpCompositeConstruct %27 %10 %10 %10 %10 -%30 = OpSMod %27 %28 %29 -%31 = OpVectorShuffle %14 %26 %26 0 1 0 1 -%32 = OpConvertSToF %14 %30 -%33 = OpFAdd %14 %31 %32 -OpReturnValue %33 +%14 = OpConstant %9 0 +%15 = OpTypeVector %4 4 +%18 = OpTypeFunction %15 +%20 = OpTypeVector %4 2 +%28 = OpTypeVector %9 4 +%37 = OpTypeFunction %9 +%44 = OpConstantNull %9 +%52 = OpTypeVector %13 4 +%56 = OpTypeFunction %2 +%17 = OpFunction %15 None %18 +%16 = OpLabel +OpBranch %19 +%19 = OpLabel +%21 = OpCompositeConstruct %20 %5 %5 +%22 = OpCompositeConstruct %20 %3 %3 +%23 = OpFAdd %20 %22 %21 +%24 = OpCompositeConstruct %20 %6 %6 +%25 = OpFSub %20 %23 %24 +%26 = OpCompositeConstruct %20 %7 %7 +%27 = OpFDiv %20 %25 %26 +%29 = OpCompositeConstruct %28 %8 %8 %8 %8 +%30 = OpCompositeConstruct %28 %10 %10 %10 %10 +%31 = OpSMod %28 %29 %30 +%32 = OpVectorShuffle %15 %27 %27 0 1 0 1 +%33 = OpConvertSToF %15 %31 +%34 = OpFAdd %15 %32 %33 +OpReturnValue %34 OpFunctionEnd -%35 = OpFunction %9 None %36 -%34 = OpLabel -OpBranch %37 -%37 = OpLabel -%38 = OpLogicalNot %13 %12 -OpSelectionMerge %39 None -OpBranchConditional %38 %40 %41 -%40 = OpLabel -OpReturnValue %11 +%36 = OpFunction %9 None %37 +%35 = OpLabel +OpBranch %38 +%38 = OpLabel +%39 = OpLogicalNot %13 %12 +OpSelectionMerge %40 None +OpBranchConditional %39 %41 %42 %41 = OpLabel -%42 = OpNot %9 %11 -OpReturnValue %42 -%39 = OpLabel +OpReturnValue %11 +%42 = OpLabel +%43 = OpNot %9 %11 OpReturnValue %43 +%40 = OpLabel +OpReturnValue %44 OpFunctionEnd -%45 = OpFunction %2 None %46 -%44 = OpLabel +%46 = OpFunction %15 None %18 +%45 = OpLabel OpBranch %47 %47 = OpLabel -%48 = OpFunctionCall %14 %16 -%49 = OpFunctionCall %9 %35 +%48 = OpCompositeConstruct %15 %3 %3 %3 %3 +%49 = OpCompositeConstruct %15 %3 %3 %3 %3 +%50 = OpSelect %9 %12 %14 %11 +%53 = OpCompositeConstruct %52 %12 %12 %12 %12 +%51 = OpSelect %15 %53 %48 %49 +OpReturnValue %51 +OpFunctionEnd +%55 = OpFunction %2 None %56 +%54 = OpLabel +OpBranch %57 +%57 = OpLabel +%58 = OpFunctionCall %15 %17 +%59 = OpFunctionCall %9 %36 +%60 = OpFunctionCall %15 %46 OpReturn OpFunctionEnd \ No newline at end of file diff --git a/tests/out/wgsl/operators.wgsl b/tests/out/wgsl/operators.wgsl index 06e3ca1dc2..bbd3c9086f 100644 --- a/tests/out/wgsl/operators.wgsl +++ b/tests/out/wgsl/operators.wgsl @@ -12,9 +12,17 @@ fn unary() -> i32 { } } +fn selection() -> vec4 { + let vector1_: vec4 = vec4(1.0); + let vector2_: vec4 = vec4(1.0); + let a: i32 = select(0, 1, true); + return select(vector1_, vector2_, true); +} + [[stage(compute), workgroup_size(1, 1, 1)]] fn main() { let _e0: vec4 = splat(); let _e1: i32 = unary(); + let _e2: vec4 = selection(); return; }