gl-out: skip unsized types if unused

src/back/glsl/mod.rs

@@ -102,8 +102,16 @@ impl crate::StorageClass {
     }
 }
 
-//Note: similar to `back/spv/helpers.rs`
+#[derive(PartialEq)]
-fn global_needs_wrapper(ir_module: &crate::Module, global_ty: Handle<crate::Type>) -> bool {
+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<crate::Type>) -> Self {
         match ir_module.types[global_ty].inner {
             crate::TypeInner::Struct {
                 ref members,
@@ -113,10 +121,11 @@ fn global_needs_wrapper(ir_module: &crate::Module, global_ty: Handle<crate::Type
                 crate::TypeInner::Array {
                     size: crate::ArraySize::Dynamic,
                     ..
-            } => false,
+                } => Self::Unsized(members),
-            _ => true,
+                _ => Self::WrappedStruct,
             },
-        _ => false,
+            _ => 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)| {
+                let generate_struct = match GlobalTypeKind::new(self.module, handle) {
-                    !ep_info[vh].is_empty() && var.class.is_buffer() && var.ty == handle
+                    GlobalTypeKind::WrappedStruct => true,
-                });
-
-                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 {
+                    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);
+            match GlobalTypeKind::new(self.module, global.ty) {
-            if needs_wrapper {
+                GlobalTypeKind::WrappedStruct => {
                     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, .. } =
+                    true
-                self.module.types[global.ty].inner
+                }
-            {
+                GlobalTypeKind::Unsized(members) => {
-                self.write_struct_body(global.ty, 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

tests/in/globals.wgsl

@@ -13,6 +13,13 @@ struct Foo {
 [[group(0), binding(1)]]
 var<storage> alignment: Foo;
 
+struct Dummy {
+    arr: array<vec2<f32>>;
+};
+
+[[group(0), binding(2)]]
+var<storage> dummy: Dummy;
+
 [[stage(compute), workgroup_size(1)]]
 fn main() {
     wg[3] = alignment.v1;

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_;
+    float _e8 = _group_0_binding_1_cs.v1_;
-    wg[3] = _e7;
+    wg[3] = _e8;
-    float _e12 = _group_0_binding_1_cs.v3_.x;
+    float _e13 = _group_0_binding_1_cs.v3_.x;
-    wg[2] = _e12;
+    wg[2] = _e13;
     at_1 = 2u;
     return;
 }

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));
+    float _expr8 = asfloat(alignment.Load(12));
-    wg[3] = _expr7;
+    wg[3] = _expr8;
-    float _expr12 = asfloat(alignment.Load(0+0));
+    float _expr13 = asfloat(alignment.Load(0+0));
-    wg[2] = _expr12;
+    wg[2] = _expr13;
     at_1 = 2u;
     return;
 }

tests/out/msl/globals.msl

@@ -2,6 +2,10 @@
 #include <metal_stdlib>
 #include <simd/simd.h>
 
+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_;
+    float _e8 = alignment.v1_;
-    wg.inner[3] = _e7;
+    wg.inner[3] = _e8;
-    float _e12 = metal::float3(alignment.v3_).x;
+    float _e13 = metal::float3(alignment.v3_).x;
-    wg.inner[2] = _e12;
+    wg.inner[2] = _e13;
     metal::atomic_store_explicit(&at_1, 2u, metal::memory_order_relaxed);
     return;
 }

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"
+OpEntryPoint GLCompute %34 "main"
-OpExecutionMode %29 LocalSize 1 1 1
+OpExecutionMode %34 LocalSize 1 1 1
 OpDecorate %14 ArrayStride 4
 OpMemberDecorate %16 0 Offset 0
 OpMemberDecorate %16 1 Offset 12
-OpDecorate %21 NonWritable
+OpDecorate %18 ArrayStride 8
-OpDecorate %21 DescriptorSet 0
+OpMemberDecorate %19 0 Offset 0
-OpDecorate %21 Binding 1
+OpDecorate %24 NonWritable
-OpDecorate %22 Block
+OpDecorate %24 DescriptorSet 0
-OpMemberDecorate %22 0 Offset 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 = OpTypeVector %12 2
-%17 = OpVariable  %18  Workgroup
+%18 = OpTypeRuntimeArray %17
-%20 = OpTypePointer Workgroup %6
+%19 = OpTypeStruct %18
-%19 = OpVariable  %20  Workgroup
+%21 = OpTypePointer Workgroup %14
-%22 = OpTypeStruct %16
+%20 = OpVariable  %21  Workgroup
-%23 = OpTypePointer StorageBuffer %22
+%23 = OpTypePointer Workgroup %6
-%21 = OpVariable  %23  StorageBuffer
+%22 = OpVariable  %23  Workgroup
-%25 = OpTypePointer Function %12
+%25 = OpTypeStruct %16
-%27 = OpTypePointer Function %4
+%26 = OpTypePointer StorageBuffer %25
-%30 = OpTypeFunction %2
+%24 = OpVariable  %26  StorageBuffer
-%31 = OpTypePointer StorageBuffer %16
+%28 = OpTypePointer StorageBuffer %19
-%32 = OpConstant  %6  0
+%27 = OpVariable  %28  StorageBuffer
-%35 = OpTypePointer Workgroup %12
+%30 = OpTypePointer Function %12
-%36 = OpTypePointer StorageBuffer %12
+%32 = OpTypePointer Function %4
-%37 = OpConstant  %6  1
+%35 = OpTypeFunction %2
-%40 = OpConstant  %6  3
+%36 = OpTypePointer StorageBuffer %16
-%42 = OpTypePointer StorageBuffer %15
+%37 = OpConstant  %6  0
-%43 = OpTypePointer StorageBuffer %12
+%40 = OpTypePointer Workgroup %12
-%47 = OpConstant  %6  256
+%41 = OpTypePointer StorageBuffer %12
-%29 = OpFunction  %2  None %30
+%42 = OpConstant  %6  1
-%28 = OpLabel
+%45 = OpConstant  %6  3
-%24 = OpVariable  %25  Function %11
+%47 = OpTypePointer StorageBuffer %15
-%26 = OpVariable  %27  Function %13
+%48 = OpTypePointer StorageBuffer %12
-%33 = OpAccessChain  %31  %21 %32
+%52 = OpConstant  %6  256
-OpBranch %34
+%34 = OpFunction  %2  None %35
-%34 = OpLabel
+%33 = OpLabel
-%38 = OpAccessChain  %36  %33 %37
+%29 = OpVariable  %30  Function %11
-%39 = OpLoad  %12  %38
+%31 = OpVariable  %32  Function %13
-%41 = OpAccessChain  %35  %17 %40
+%38 = OpAccessChain  %36  %24 %37
-OpStore %41 %39
+OpBranch %39
-%44 = OpAccessChain  %43  %33 %32 %32
+%39 = OpLabel
-%45 = OpLoad  %12  %44
+%43 = OpAccessChain  %41  %38 %42
-%46 = OpAccessChain  %35  %17 %10
+%44 = OpLoad  %12  %43
-OpStore %46 %45
+%46 = OpAccessChain  %40  %20 %45
-OpAtomicStore %19 %9 %47 %10
+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

tests/out/wgsl/globals.wgsl

@@ -3,22 +3,28 @@ struct Foo {
     v1_: f32;
 };
 
+struct Dummy {
+    arr: [[stride(8)]] array<vec2<f32>>;
+};
+
 let Foo_2: bool = true;
 
 var<workgroup> wg: array<f32,10u>;
 var<workgroup> at_1: atomic<u32>;
 [[group(0), binding(1)]]
 var<storage> alignment: Foo;
+[[group(0), binding(2)]]
+var<storage> 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_;
+    let _e8 = alignment.v1_;
-    wg[3] = _e7;
+    wg[3] = _e8;
-    let _e12 = alignment.v3_.x;
+    let _e13 = alignment.v3_.x;
-    wg[2] = _e12;
+    wg[2] = _e13;
     atomicStore((&at_1), 2u);
     return;
 }