From 13665466453b457fa4a46eacff0655cd5bd515a2 Mon Sep 17 00:00:00 2001 From: Dzmitry Malyshau Date: Fri, 28 May 2021 22:11:39 -0400 Subject: [PATCH] [glsl-out] improve array types --- src/back/glsl/mod.rs | 69 ++++++++++++++++++--------------- tests/in/globals.wgsl | 3 +- tests/out/boids.Compute.glsl | 3 -- tests/out/globals.Compute.glsl | 2 + tests/out/globals.msl | 4 +- tests/out/globals.spvasm | 32 ++++++++------- tests/out/globals.wgsl | 3 +- tests/out/quad-vert.Vertex.glsl | 4 -- tests/out/shadow.Fragment.glsl | 2 - tests/out/skybox.Vertex.glsl | 1 - 10 files changed, 66 insertions(+), 57 deletions(-) diff --git a/src/back/glsl/mod.rs b/src/back/glsl/mod.rs index c74011a68d..985c385c02 100644 --- a/src/back/glsl/mod.rs +++ b/src/back/glsl/mod.rs @@ -564,6 +564,27 @@ impl<'a, W: Write> Writer<'a, W> { self.collect_reflection_info() } + fn write_array_size(&mut self, size: ArraySize) -> BackendResult { + write!(self.out, "[")?; + + // Write the array size + // Writes nothing if `ArraySize::Dynamic` + // Panics if `ArraySize::Constant` has a constant that isn't an uint + match size { + ArraySize::Constant(const_handle) => match self.module.constants[const_handle].inner { + ConstantInner::Scalar { + width: _, + value: ScalarValue::Uint(size), + } => write!(self.out, "{}", size)?, + _ => unreachable!(), + }, + ArraySize::Dynamic => (), + } + + write!(self.out, "]")?; + Ok(()) + } + /// Helper method used to write value types /// /// # Notes @@ -615,27 +636,7 @@ impl<'a, W: Write> Writer<'a, W> { // GLSL arrays are written as `type name[size]` // Current code is written arrays only as `[size]` // Base `type` and `name` should be written outside - TypeInner::Array { base: _, size, .. } => { - write!(self.out, "[")?; - - // Write the array size - // Writes nothing if `ArraySize::Dynamic` - // Panics if `ArraySize::Constant` has a constant that isn't an uint - match size { - ArraySize::Constant(const_handle) => { - match self.module.constants[const_handle].inner { - ConstantInner::Scalar { - width: _, - value: ScalarValue::Uint(size), - } => write!(self.out, "{}", size)?, - _ => unreachable!(), - } - } - ArraySize::Dynamic => (), - } - - write!(self.out, "]")? - } + TypeInner::Array { size, .. } => self.write_array_size(size)?, // Panic if either Image, Sampler, Pointer, or a Struct is being written // // Write all variants instead of `_` so that if new variants are added a @@ -674,6 +675,8 @@ impl<'a, W: Write> Writer<'a, W> { write!(self.out, "{}", name)?; Ok(()) } + // glsl array has the size separated from the base type + TypeInner::Array { base, .. } => self.write_type(base), ref other => self.write_value_type(other), } } @@ -756,14 +759,14 @@ impl<'a, W: Write> Writer<'a, W> { // Finally write the global name and end the global with a `;` and a newline // Leading space is important let global_name = self.get_global_name(handle, global); - let global_str = - if let Some(default_value) = zero_init_value_str(&self.module.types[global.ty].inner) { - format!("{} = {}", global_name, default_value) - } else { - global_name - }; - writeln!(self.out, " {};", global_str)?; - writeln!(self.out)?; + write!(self.out, " {}", global_name)?; + if let TypeInner::Array { size, .. } = self.module.types[global.ty].inner { + self.write_array_size(size)?; + } + if let Some(default_value) = zero_init_value_str(&self.module.types[global.ty].inner) { + write!(self.out, " = {}", default_value)?; + }; + writeln!(self.out, ";")?; Ok(()) } @@ -1117,7 +1120,11 @@ impl<'a, W: Write> Writer<'a, W> { write!(self.out, "{}", INDENT)?; match self.module.types[member.ty].inner { - TypeInner::Array { base, .. } => { + TypeInner::Array { + base, + size, + stride: _, + } => { // GLSL arrays are written as `type name[size]` let ty_name = match self.module.types[base].inner { // Write scalar type by backend so as not to depend on the front-end implementation @@ -1134,7 +1141,7 @@ impl<'a, W: Write> Writer<'a, W> { &self.names[&NameKey::StructMember(handle, idx as u32)] )?; // Write [size] - self.write_type(member.ty)?; + self.write_array_size(size)?; // Newline is important writeln!(self.out, ";")?; } diff --git a/tests/in/globals.wgsl b/tests/in/globals.wgsl index 3f20e53efc..bab28b3fe8 100644 --- a/tests/in/globals.wgsl +++ b/tests/in/globals.wgsl @@ -2,8 +2,9 @@ let Foo: bool = true; -var wg : array; +var wg : array; [[stage(compute), workgroup_size(1)]] fn main() { + wg[3] = 1.0; } diff --git a/tests/out/boids.Compute.glsl b/tests/out/boids.Compute.glsl index e518b6cc76..e54432d5cc 100644 --- a/tests/out/boids.Compute.glsl +++ b/tests/out/boids.Compute.glsl @@ -18,16 +18,13 @@ uniform SimParams_block_0 { float rule2Scale; float rule3Scale; } _group_0_binding_0; - readonly buffer Particles_block_1 { Particle particles[]; } _group_0_binding_1; - buffer Particles_block_2 { Particle particles[]; } _group_0_binding_2; - void main() { uvec3 global_invocation_id = gl_GlobalInvocationID; vec2 vPos; diff --git a/tests/out/globals.Compute.glsl b/tests/out/globals.Compute.glsl index 63328eab18..2fba504b70 100644 --- a/tests/out/globals.Compute.glsl +++ b/tests/out/globals.Compute.glsl @@ -4,8 +4,10 @@ precision highp float; layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; +shared float wg[10]; void main() { + wg[3] = 1.0; return; } diff --git a/tests/out/globals.msl b/tests/out/globals.msl index 638fa22ede..435c090f27 100644 --- a/tests/out/globals.msl +++ b/tests/out/globals.msl @@ -3,10 +3,12 @@ constexpr constant bool Foo = true; struct type2 { - float inner[10]; + float inner[10u]; }; kernel void main1( + threadgroup type2& wg ) { + wg.inner[3] = 1.0; return; } diff --git a/tests/out/globals.spvasm b/tests/out/globals.spvasm index 0078cf2e8f..2668d178d0 100644 --- a/tests/out/globals.spvasm +++ b/tests/out/globals.spvasm @@ -1,26 +1,32 @@ ; SPIR-V ; Version: 1.0 ; Generator: rspirv -; Bound: 15 +; Bound: 20 OpCapability Shader %1 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 -OpEntryPoint GLCompute %12 "main" -OpExecutionMode %12 LocalSize 1 1 1 -OpDecorate %8 ArrayStride 4 +OpEntryPoint GLCompute %15 "main" +OpExecutionMode %15 LocalSize 1 1 1 +OpDecorate %11 ArrayStride 4 %2 = OpTypeVoid %4 = OpTypeBool %3 = OpConstantTrue %4 -%6 = OpTypeInt 32 1 +%6 = OpTypeInt 32 0 %5 = OpConstant %6 10 -%7 = OpTypeFloat 32 -%8 = OpTypeArray %7 %5 -%10 = OpTypePointer Workgroup %8 -%9 = OpVariable %10 Workgroup -%13 = OpTypeFunction %2 -%12 = OpFunction %2 None %13 -%11 = OpLabel -OpBranch %14 +%8 = OpTypeInt 32 1 +%7 = OpConstant %8 3 +%10 = OpTypeFloat 32 +%9 = OpConstant %10 1.0 +%11 = OpTypeArray %10 %5 +%13 = OpTypePointer Workgroup %11 +%12 = OpVariable %13 Workgroup +%16 = OpTypeFunction %2 +%18 = OpTypePointer Workgroup %10 +%15 = OpFunction %2 None %16 %14 = OpLabel +OpBranch %17 +%17 = OpLabel +%19 = OpAccessChain %18 %12 %7 +OpStore %19 %9 OpReturn OpFunctionEnd \ No newline at end of file diff --git a/tests/out/globals.wgsl b/tests/out/globals.wgsl index 87ce46054a..36856539fa 100644 --- a/tests/out/globals.wgsl +++ b/tests/out/globals.wgsl @@ -1,8 +1,9 @@ let Foo: bool = true; -var wg: array; +var wg: array; [[stage(compute), workgroup_size(1, 1, 1)]] fn main() { + wg[3] = 1.0; return; } diff --git a/tests/out/quad-vert.Vertex.glsl b/tests/out/quad-vert.Vertex.glsl index e1b9853b05..c2f59e712f 100644 --- a/tests/out/quad-vert.Vertex.glsl +++ b/tests/out/quad-vert.Vertex.glsl @@ -11,18 +11,14 @@ struct type10 { }; vec2 v_uv = vec2(0, 0); - vec2 a_uv = vec2(0, 0); - struct gen_gl_PerVertex_block_0 { vec4 gen_gl_Position; float gen_gl_PointSize; float gen_gl_ClipDistance[1]; float gen_gl_CullDistance[1]; } perVertexStruct; - vec2 a_pos = vec2(0, 0); - layout(location = 1) in vec2 _p2vs_location1; layout(location = 0) in vec2 _p2vs_location0; smooth layout(location = 0) out vec2 _vs2fs_location0; diff --git a/tests/out/shadow.Fragment.glsl b/tests/out/shadow.Fragment.glsl index 2042b5f22d..5b9c58de0e 100644 --- a/tests/out/shadow.Fragment.glsl +++ b/tests/out/shadow.Fragment.glsl @@ -11,11 +11,9 @@ struct Light { uniform Globals_block_0 { uvec4 num_lights; } _group_0_binding_0; - readonly buffer Lights_block_1 { Light data[]; } _group_0_binding_1; - uniform highp sampler2DArrayShadow _group_0_binding_2; smooth layout(location = 0) in vec3 _vs2fs_location0; diff --git a/tests/out/skybox.Vertex.glsl b/tests/out/skybox.Vertex.glsl index 870a4ee26e..2d35107bc9 100644 --- a/tests/out/skybox.Vertex.glsl +++ b/tests/out/skybox.Vertex.glsl @@ -11,7 +11,6 @@ uniform Data_block_0 { mat4x4 proj_inv; mat4x4 view; } _group_0_binding_0; - smooth layout(location = 0) out vec3 _vs2fs_location0; void main() {