diff --git a/src/back/glsl/mod.rs b/src/back/glsl/mod.rs index 72b31d4598..239cb0a159 100644 --- a/src/back/glsl/mod.rs +++ b/src/back/glsl/mod.rs @@ -102,21 +102,30 @@ impl crate::StorageClass { } } -//Note: similar to `back/spv/helpers.rs` -fn global_needs_wrapper(ir_module: &crate::Module, global_ty: Handle) -> bool { - match ir_module.types[global_ty].inner { - crate::TypeInner::Struct { - ref members, - span: _, - } => match ir_module.types[members.last().unwrap().ty].inner { - // Structs with dynamically sized arrays can't be copied and can't be wrapped. - crate::TypeInner::Array { - size: crate::ArraySize::Dynamic, - .. - } => false, - _ => true, - }, - _ => false, +#[derive(PartialEq)] +enum GlobalTypeKind<'a> { + WrappedStruct, + Unsized(&'a [crate::StructMember]), + Other, +} + +impl<'a> GlobalTypeKind<'a> { + //Note: similar to `back/spv/helpers.rs` + fn new(ir_module: &'a crate::Module, global_ty: Handle) -> Self { + match ir_module.types[global_ty].inner { + crate::TypeInner::Struct { + ref members, + span: _, + } => match ir_module.types[members.last().unwrap().ty].inner { + // Structs with dynamically sized arrays can't be copied and can't be wrapped. + crate::TypeInner::Array { + size: crate::ArraySize::Dynamic, + .. + } => Self::Unsized(members), + _ => Self::WrappedStruct, + }, + _ => Self::Other, + } } } @@ -545,14 +554,21 @@ impl<'a, W: Write> Writer<'a, W> { // struct without adding all of it's members first for (handle, ty) in self.module.types.iter() { if let TypeInner::Struct { ref members, .. } = ty.inner { - let used_by_global = self.module.global_variables.iter().any(|(vh, var)| { - !ep_info[vh].is_empty() && var.class.is_buffer() && var.ty == handle - }); - - let is_wrapped = global_needs_wrapper(self.module, handle); - // If it's a global non-wrapped struct, it will be printed - // with the corresponding global variable. - if !used_by_global || is_wrapped { + let generate_struct = match GlobalTypeKind::new(self.module, handle) { + GlobalTypeKind::WrappedStruct => true, + // If it's a global non-wrapped struct, it will be printed + // with the corresponding global variable. + GlobalTypeKind::Unsized(_) => false, + GlobalTypeKind::Other => { + let used_by_global = + self.module.global_variables.iter().any(|(vh, var)| { + !ep_info[vh].is_empty() && var.class.is_buffer() && var.ty == handle + }); + // If not used by a global, it's safe to just spew it here + !used_by_global + } + }; + if generate_struct { let name = &self.names[&NameKey::Type(handle)]; write!(self.out, "struct {} ", name)?; self.write_struct_body(handle, members)?; @@ -925,18 +941,22 @@ impl<'a, W: Write> Writer<'a, W> { write!(self.out, "{} ", block_name)?; self.reflection_names_globals.insert(handle, block_name); - let needs_wrapper = global_needs_wrapper(self.module, global.ty); - if needs_wrapper { - write!(self.out, "{{ ")?; - // Write the type - // `write_type` adds no leading or trailing spaces - self.write_type(global.ty)?; - } else if let crate::TypeInner::Struct { ref members, .. } = - self.module.types[global.ty].inner - { - self.write_struct_body(global.ty, members)?; + match GlobalTypeKind::new(self.module, global.ty) { + GlobalTypeKind::WrappedStruct => { + write!(self.out, "{{ ")?; + // Write the type + // `write_type` adds no leading or trailing spaces + self.write_type(global.ty)?; + true + } + GlobalTypeKind::Unsized(members) => { + self.write_struct_body(global.ty, members)?; + false + } + GlobalTypeKind::Other => { + return Err(Error::Custom("Non-struct type of a buffer".to_string())); + } } - needs_wrapper } else { self.write_type(global.ty)?; false diff --git a/tests/in/globals.wgsl b/tests/in/globals.wgsl index 03f87ddfb9..6c9fc1c162 100644 --- a/tests/in/globals.wgsl +++ b/tests/in/globals.wgsl @@ -13,6 +13,13 @@ struct Foo { [[group(0), binding(1)]] var alignment: Foo; +struct Dummy { + arr: array>; +}; + +[[group(0), binding(2)]] +var dummy: Dummy; + [[stage(compute), workgroup_size(1)]] fn main() { wg[3] = alignment.v1; diff --git a/tests/out/glsl/globals.main.Compute.glsl b/tests/out/glsl/globals.main.Compute.glsl index cdc5e7260a..4b53779517 100644 --- a/tests/out/glsl/globals.main.Compute.glsl +++ b/tests/out/glsl/globals.main.Compute.glsl @@ -19,10 +19,10 @@ layout(std430) readonly buffer Foo_block_0Compute { Foo _group_0_binding_1_cs; } void main() { float Foo_1 = 1.0; bool at = true; - float _e7 = _group_0_binding_1_cs.v1_; - wg[3] = _e7; - float _e12 = _group_0_binding_1_cs.v3_.x; - wg[2] = _e12; + float _e8 = _group_0_binding_1_cs.v1_; + wg[3] = _e8; + float _e13 = _group_0_binding_1_cs.v3_.x; + wg[2] = _e13; at_1 = 2u; return; } diff --git a/tests/out/hlsl/globals.hlsl b/tests/out/hlsl/globals.hlsl index 1e05d9fa7a..41a114f5b8 100644 --- a/tests/out/hlsl/globals.hlsl +++ b/tests/out/hlsl/globals.hlsl @@ -8,6 +8,7 @@ struct Foo { groupshared float wg[10]; groupshared uint at_1; ByteAddressBuffer alignment : register(t1); +ByteAddressBuffer dummy : register(t2); [numthreads(1, 1, 1)] void main() @@ -15,10 +16,10 @@ void main() float Foo_1 = 1.0; bool at = true; - float _expr7 = asfloat(alignment.Load(12)); - wg[3] = _expr7; - float _expr12 = asfloat(alignment.Load(0+0)); - wg[2] = _expr12; + float _expr8 = asfloat(alignment.Load(12)); + wg[3] = _expr8; + float _expr13 = asfloat(alignment.Load(0+0)); + wg[2] = _expr13; at_1 = 2u; return; } diff --git a/tests/out/msl/globals.msl b/tests/out/msl/globals.msl index 636aa69c8d..f2c37b5ab1 100644 --- a/tests/out/msl/globals.msl +++ b/tests/out/msl/globals.msl @@ -2,6 +2,10 @@ #include #include +struct _mslBufferSizes { + metal::uint size3; +}; + constexpr constant bool Foo_2 = true; struct type_2 { float inner[10u]; @@ -10,6 +14,10 @@ struct Foo { metal::packed_float3 v3_; float v1_; }; +typedef metal::float2 type_6[1]; +struct Dummy { + type_6 arr; +}; kernel void main_( threadgroup type_2& wg @@ -18,10 +26,10 @@ kernel void main_( ) { float Foo_1 = 1.0; bool at = true; - float _e7 = alignment.v1_; - wg.inner[3] = _e7; - float _e12 = metal::float3(alignment.v3_).x; - wg.inner[2] = _e12; + float _e8 = alignment.v1_; + wg.inner[3] = _e8; + float _e13 = metal::float3(alignment.v3_).x; + wg.inner[2] = _e13; metal::atomic_store_explicit(&at_1, 2u, metal::memory_order_relaxed); return; } diff --git a/tests/out/spv/globals.spvasm b/tests/out/spv/globals.spvasm index c9ec4c2203..a11785e8c4 100644 --- a/tests/out/spv/globals.spvasm +++ b/tests/out/spv/globals.spvasm @@ -1,21 +1,27 @@ ; SPIR-V ; Version: 1.1 ; Generator: rspirv -; Bound: 48 +; Bound: 53 OpCapability Shader OpExtension "SPV_KHR_storage_buffer_storage_class" %1 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 -OpEntryPoint GLCompute %29 "main" -OpExecutionMode %29 LocalSize 1 1 1 +OpEntryPoint GLCompute %34 "main" +OpExecutionMode %34 LocalSize 1 1 1 OpDecorate %14 ArrayStride 4 OpMemberDecorate %16 0 Offset 0 OpMemberDecorate %16 1 Offset 12 -OpDecorate %21 NonWritable -OpDecorate %21 DescriptorSet 0 -OpDecorate %21 Binding 1 -OpDecorate %22 Block -OpMemberDecorate %22 0 Offset 0 +OpDecorate %18 ArrayStride 8 +OpMemberDecorate %19 0 Offset 0 +OpDecorate %24 NonWritable +OpDecorate %24 DescriptorSet 0 +OpDecorate %24 Binding 1 +OpDecorate %25 Block +OpMemberDecorate %25 0 Offset 0 +OpDecorate %27 NonWritable +OpDecorate %27 DescriptorSet 0 +OpDecorate %27 Binding 2 +OpDecorate %19 Block %2 = OpTypeVoid %4 = OpTypeBool %3 = OpConstantTrue %4 @@ -31,40 +37,45 @@ OpMemberDecorate %22 0 Offset 0 %14 = OpTypeArray %12 %5 %15 = OpTypeVector %12 3 %16 = OpTypeStruct %15 %12 -%18 = OpTypePointer Workgroup %14 -%17 = OpVariable %18 Workgroup -%20 = OpTypePointer Workgroup %6 -%19 = OpVariable %20 Workgroup -%22 = OpTypeStruct %16 -%23 = OpTypePointer StorageBuffer %22 -%21 = OpVariable %23 StorageBuffer -%25 = OpTypePointer Function %12 -%27 = OpTypePointer Function %4 -%30 = OpTypeFunction %2 -%31 = OpTypePointer StorageBuffer %16 -%32 = OpConstant %6 0 -%35 = OpTypePointer Workgroup %12 -%36 = OpTypePointer StorageBuffer %12 -%37 = OpConstant %6 1 -%40 = OpConstant %6 3 -%42 = OpTypePointer StorageBuffer %15 -%43 = OpTypePointer StorageBuffer %12 -%47 = OpConstant %6 256 -%29 = OpFunction %2 None %30 -%28 = OpLabel -%24 = OpVariable %25 Function %11 -%26 = OpVariable %27 Function %13 -%33 = OpAccessChain %31 %21 %32 -OpBranch %34 -%34 = OpLabel -%38 = OpAccessChain %36 %33 %37 -%39 = OpLoad %12 %38 -%41 = OpAccessChain %35 %17 %40 -OpStore %41 %39 -%44 = OpAccessChain %43 %33 %32 %32 -%45 = OpLoad %12 %44 -%46 = OpAccessChain %35 %17 %10 -OpStore %46 %45 -OpAtomicStore %19 %9 %47 %10 +%17 = OpTypeVector %12 2 +%18 = OpTypeRuntimeArray %17 +%19 = OpTypeStruct %18 +%21 = OpTypePointer Workgroup %14 +%20 = OpVariable %21 Workgroup +%23 = OpTypePointer Workgroup %6 +%22 = OpVariable %23 Workgroup +%25 = OpTypeStruct %16 +%26 = OpTypePointer StorageBuffer %25 +%24 = OpVariable %26 StorageBuffer +%28 = OpTypePointer StorageBuffer %19 +%27 = OpVariable %28 StorageBuffer +%30 = OpTypePointer Function %12 +%32 = OpTypePointer Function %4 +%35 = OpTypeFunction %2 +%36 = OpTypePointer StorageBuffer %16 +%37 = OpConstant %6 0 +%40 = OpTypePointer Workgroup %12 +%41 = OpTypePointer StorageBuffer %12 +%42 = OpConstant %6 1 +%45 = OpConstant %6 3 +%47 = OpTypePointer StorageBuffer %15 +%48 = OpTypePointer StorageBuffer %12 +%52 = OpConstant %6 256 +%34 = OpFunction %2 None %35 +%33 = OpLabel +%29 = OpVariable %30 Function %11 +%31 = OpVariable %32 Function %13 +%38 = OpAccessChain %36 %24 %37 +OpBranch %39 +%39 = OpLabel +%43 = OpAccessChain %41 %38 %42 +%44 = OpLoad %12 %43 +%46 = OpAccessChain %40 %20 %45 +OpStore %46 %44 +%49 = OpAccessChain %48 %38 %37 %37 +%50 = OpLoad %12 %49 +%51 = OpAccessChain %40 %20 %10 +OpStore %51 %50 +OpAtomicStore %22 %9 %52 %10 OpReturn OpFunctionEnd \ No newline at end of file diff --git a/tests/out/wgsl/globals.wgsl b/tests/out/wgsl/globals.wgsl index a79623904b..c332963333 100644 --- a/tests/out/wgsl/globals.wgsl +++ b/tests/out/wgsl/globals.wgsl @@ -3,22 +3,28 @@ struct Foo { v1_: f32; }; +struct Dummy { + arr: [[stride(8)]] array>; +}; + let Foo_2: bool = true; var wg: array; var at_1: atomic; [[group(0), binding(1)]] var alignment: Foo; +[[group(0), binding(2)]] +var dummy: Dummy; [[stage(compute), workgroup_size(1, 1, 1)]] fn main() { var Foo_1: f32 = 1.0; var at: bool = true; - let _e7 = alignment.v1_; - wg[3] = _e7; - let _e12 = alignment.v3_.x; - wg[2] = _e12; + let _e8 = alignment.v1_; + wg[3] = _e8; + let _e13 = alignment.v3_.x; + wg[2] = _e13; atomicStore((&at_1), 2u); return; }