[spv-out] Emit required decorations on wrapper struct types, too.

src/back/spv/helpers.rs

@@ -62,7 +62,21 @@ impl crate::AddressSpace {
     }
 }
 
-/// Return true if the global requires a type decorated with "Block".
+/// Return true if the global requires a type decorated with `Block`.
+///
+/// Vulkan spec v1.3 §15.6.2, "Descriptor Set Interface", says:
+///
+/// > Variables identified with the `Uniform` storage class are used to
+/// > access transparent buffer backed resources. Such variables must
+/// > be:
+/// >
+/// > -   typed as `OpTypeStruct`, or an array of this type,
+/// >
+/// > -   identified with a `Block` or `BufferBlock` decoration, and
+/// >
+/// > -   laid out explicitly using the `Offset`, `ArrayStride`, and
+/// >     `MatrixStride` decorations as specified in §15.6.4, "Offset
+/// >     and Stride Assignment."
 // See `back::spv::GlobalVariable::access_id` for details.
 pub fn global_needs_wrapper(ir_module: &crate::Module, var: &crate::GlobalVariable) -> bool {
     match var.space {

src/back/spv/writer.rs

@@ -799,8 +799,6 @@ impl Writer {
         handle: Handle<crate::Type>,
     ) -> Result<Word, Error> {
         let ty = &arena[handle];
-        let decorate_layout = true; //TODO?
-
         let id = if let Some(local) = make_local(&ty.inner) {
             // This type can be represented as a `LocalType`, so check if we've
             // already written an instruction for it. If not, do so now, with
@@ -829,9 +827,7 @@ impl Writer {
             let id = self.id_gen.next();
             let instruction = match ty.inner {
                 crate::TypeInner::Array { base, size, stride } => {
-                    if decorate_layout {
-                        self.decorate(id, Decoration::ArrayStride, &[stride]);
-                    }
+                    self.decorate(id, Decoration::ArrayStride, &[stride]);
 
                     let type_id = self.get_type_id(LookupType::Handle(base));
                     match size {
@@ -848,52 +844,7 @@ impl Writer {
                 } => {
                     let mut member_ids = Vec::with_capacity(members.len());
                     for (index, member) in members.iter().enumerate() {
-                        if decorate_layout {
-                            self.annotations.push(Instruction::member_decorate(
-                                id,
-                                index as u32,
-                                Decoration::Offset,
-                                &[member.offset],
-                            ));
-                        }
-
-                        if self.flags.contains(WriterFlags::DEBUG) {
-                            if let Some(ref name) = member.name {
-                                self.debugs
-                                    .push(Instruction::member_name(id, index as u32, name));
-                            }
-                        }
-
-                        // The matrix decorations also go on arrays of matrices,
-                        // so lets check this first.
-                        let member_array_subty_inner = match arena[member.ty].inner {
-                            crate::TypeInner::Array { base, .. } => &arena[base].inner,
-                            ref other => other,
-                        };
-                        if let crate::TypeInner::Matrix {
-                            columns: _,
-                            rows,
-                            width,
-                        } = *member_array_subty_inner
-                        {
-                            let byte_stride = match rows {
-                                crate::VectorSize::Bi => 2 * width,
-                                crate::VectorSize::Tri | crate::VectorSize::Quad => 4 * width,
-                            };
-                            self.annotations.push(Instruction::member_decorate(
-                                id,
-                                index as u32,
-                                Decoration::ColMajor,
-                                &[],
-                            ));
-                            self.annotations.push(Instruction::member_decorate(
-                                id,
-                                index as u32,
-                                Decoration::MatrixStride,
-                                &[byte_stride as u32],
-                            ));
-                        }
-
+                        self.decorate_struct_member(id, index, member, arena)?;
                         let member_id = self.get_type_id(LookupType::Handle(member.ty));
                         member_ids.push(member_id);
                     }
@@ -1287,12 +1238,14 @@ impl Writer {
             let wrapper_type_id = self.id_gen.next();
 
             self.decorate(wrapper_type_id, Decoration::Block, &[]);
-            self.annotations.push(Instruction::member_decorate(
-                wrapper_type_id,
-                0,
-                Decoration::Offset,
-                &[0],
-            ));
+            let member = crate::StructMember {
+                name: None,
+                ty: global_variable.ty,
+                binding: None,
+                offset: 0,
+            };
+            self.decorate_struct_member(wrapper_type_id, 0, &member, &ir_module.types)?;
+
             Instruction::type_struct(wrapper_type_id, &[inner_type_id])
                 .to_words(&mut self.logical_layout.declarations);
 
@@ -1317,6 +1270,66 @@ impl Writer {
         Ok(id)
     }
 
+    /// Write the necessary decorations for a struct member.
+    ///
+    /// Emit decorations for the `index`'th member of the struct type
+    /// designated by `struct_id`, described by `member`.
+    fn decorate_struct_member(
+        &mut self,
+        struct_id: Word,
+        index: usize,
+        member: &crate::StructMember,
+        arena: &UniqueArena<crate::Type>,
+    ) -> Result<(), Error> {
+        use spirv::Decoration;
+
+        self.annotations.push(Instruction::member_decorate(
+            struct_id,
+            index as u32,
+            Decoration::Offset,
+            &[member.offset],
+        ));
+
+        if self.flags.contains(WriterFlags::DEBUG) {
+            if let Some(ref name) = member.name {
+                self.debugs
+                    .push(Instruction::member_name(struct_id, index as u32, name));
+            }
+        }
+
+        // Matrices and arrays of matrices both require decorations,
+        // so "see through" an array to determine if they're needed.
+        let member_array_subty_inner = match arena[member.ty].inner {
+            crate::TypeInner::Array { base, .. } => &arena[base].inner,
+            ref other => other,
+        };
+        if let crate::TypeInner::Matrix {
+            columns: _,
+            rows,
+            width,
+        } = *member_array_subty_inner
+        {
+            let byte_stride = match rows {
+                crate::VectorSize::Bi => 2 * width,
+                crate::VectorSize::Tri | crate::VectorSize::Quad => 4 * width,
+            };
+            self.annotations.push(Instruction::member_decorate(
+                struct_id,
+                index as u32,
+                Decoration::ColMajor,
+                &[],
+            ));
+            self.annotations.push(Instruction::member_decorate(
+                struct_id,
+                index as u32,
+                Decoration::MatrixStride,
+                &[byte_stride as u32],
+            ));
+        }
+
+        Ok(())
+    }
+
     fn get_function_type(&mut self, lookup_function_type: LookupFunctionType) -> Word {
         match self
             .lookup_function_type

tests/in/globals.wgsl

@@ -22,6 +22,9 @@ var<uniform> float_vecs: array<vec4<f32>, 20>;
 @group(0) @binding(4)
 var<uniform> global_vec: vec4<f32>;
 
+@group(0) @binding(5)
+var<uniform> global_mat: mat4x4<f32>;
+
 fn test_msl_packed_vec3_as_arg(arg: vec3<f32>) {}
 
 fn test_msl_packed_vec3() {
@@ -53,7 +56,7 @@ fn test_msl_packed_vec3() {
 fn main() {
     test_msl_packed_vec3();
 
-    wg[6] = global_vec.x;
+    wg[6] = (global_mat * global_vec).x;
     wg[5] = dummy[1].y;
     wg[4] = float_vecs[0].w;
     wg[3] = alignment.v1;

tests/out/glsl/globals.main.Compute.glsl

@@ -21,6 +21,8 @@ uniform type_8_block_2Compute { vec4 _group_0_binding_3_cs[20]; };
 
 uniform type_7_block_3Compute { vec4 _group_0_binding_4_cs; };
 
+uniform type_9_block_4Compute { mat4x4 _group_0_binding_5_cs; };
+
 
 void test_msl_packed_vec3_as_arg(vec3 arg) {
     return;
@@ -31,8 +33,8 @@ void test_msl_packed_vec3_() {
     _group_0_binding_1_cs.v3_ = vec3(1.0);
     _group_0_binding_1_cs.v3_.x = 1.0;
     _group_0_binding_1_cs.v3_.x = 2.0;
-    int _e20 = idx;
-    _group_0_binding_1_cs.v3_[_e20] = 3.0;
+    int _e21 = idx;
+    _group_0_binding_1_cs.v3_[_e21] = 3.0;
     Foo data = _group_0_binding_1_cs;
     vec3 unnamed = data.v3_;
     vec2 unnamed_1 = data.v3_.zx;
@@ -47,16 +49,17 @@ void main() {
     float Foo_1 = 1.0;
     bool at = true;
     test_msl_packed_vec3_();
-    float _e10 = _group_0_binding_4_cs.x;
-    wg[6] = _e10;
-    float _e16 = _group_0_binding_2_cs[1].y;
-    wg[5] = _e16;
-    float _e22 = _group_0_binding_3_cs[0].w;
-    wg[4] = _e22;
-    float _e26 = _group_0_binding_1_cs.v1_;
-    wg[3] = _e26;
-    float _e31 = _group_0_binding_1_cs.v3_.x;
-    wg[2] = _e31;
+    mat4x4 _e10 = _group_0_binding_5_cs;
+    vec4 _e11 = _group_0_binding_4_cs;
+    wg[6] = (_e10 * _e11).x;
+    float _e19 = _group_0_binding_2_cs[1].y;
+    wg[5] = _e19;
+    float _e25 = _group_0_binding_3_cs[0].w;
+    wg[4] = _e25;
+    float _e29 = _group_0_binding_1_cs.v1_;
+    wg[3] = _e29;
+    float _e34 = _group_0_binding_1_cs.v3_.x;
+    wg[2] = _e34;
     _group_0_binding_1_cs.v1_ = 4.0;
     wg[1] = float(uint(_group_0_binding_2_cs.length()));
     at_1 = 2u;

tests/out/hlsl/globals.hlsl

@@ -11,6 +11,7 @@ RWByteAddressBuffer alignment : register(u1);
 ByteAddressBuffer dummy : register(t2);
 cbuffer float_vecs : register(b3) { float4 float_vecs[20]; }
 cbuffer global_vec : register(b4) { float4 global_vec; }
+cbuffer global_mat : register(b5) { float4x4 global_mat; }
 
 void test_msl_packed_vec3_as_arg(float3 arg)
 {
@@ -24,8 +25,8 @@ void test_msl_packed_vec3_()
     alignment.Store3(0, asuint(float3(1.0.xxx)));
     alignment.Store(0+0, asuint(1.0));
     alignment.Store(0+0, asuint(2.0));
-    int _expr20 = idx;
-    alignment.Store(_expr20*4+0, asuint(3.0));
+    int _expr21 = idx;
+    alignment.Store(_expr21*4+0, asuint(3.0));
     Foo data = {asfloat(alignment.Load3(0)), asfloat(alignment.Load(12))};
     float3 unnamed = data.v3_;
     float2 unnamed_1 = data.v3_.zx;
@@ -50,16 +51,17 @@ void main()
     bool at = true;
 
     test_msl_packed_vec3_();
-    float _expr10 = global_vec.x;
-    wg[6] = _expr10;
-    float _expr16 = asfloat(dummy.Load(4+8));
-    wg[5] = _expr16;
-    float _expr22 = float_vecs[0].w;
-    wg[4] = _expr22;
-    float _expr26 = asfloat(alignment.Load(12));
-    wg[3] = _expr26;
-    float _expr31 = asfloat(alignment.Load(0+0));
-    wg[2] = _expr31;
+    float4x4 _expr10 = global_mat;
+    float4 _expr11 = global_vec;
+    wg[6] = mul(_expr11, _expr10).x;
+    float _expr19 = asfloat(dummy.Load(4+8));
+    wg[5] = _expr19;
+    float _expr25 = float_vecs[0].w;
+    wg[4] = _expr25;
+    float _expr29 = asfloat(alignment.Load(12));
+    wg[3] = _expr29;
+    float _expr34 = asfloat(alignment.Load(0+0));
+    wg[2] = _expr34;
     alignment.Store(12, asuint(4.0));
     wg[1] = float(((NagaBufferLength(dummy) - 0) / 8));
     at_1 = 2u;

tests/out/msl/globals.msl

@@ -21,7 +21,7 @@ struct type_8 {
     metal::float4 inner[20];
 };
 constant metal::float3 const_type_4_ = {0.0, 0.0, 0.0};
-constant metal::float3x3 const_type_10_ = {const_type_4_, const_type_4_, const_type_4_};
+constant metal::float3x3 const_type_11_ = {const_type_4_, const_type_4_, const_type_4_};
 
 void test_msl_packed_vec3_as_arg(
     metal::float3 arg
@@ -36,14 +36,14 @@ void test_msl_packed_vec3_(
     alignment.v3_ = metal::float3(1.0);
     alignment.v3_[0] = 1.0;
     alignment.v3_[0] = 2.0;
-    int _e20 = idx;
-    alignment.v3_[_e20] = 3.0;
+    int _e21 = idx;
+    alignment.v3_[_e21] = 3.0;
     Foo data = alignment;
     metal::float3 unnamed = data.v3_;
     metal::float2 unnamed_1 = metal::float3(data.v3_).zx;
     test_msl_packed_vec3_as_arg(data.v3_);
-    metal::float3 unnamed_2 = metal::float3(data.v3_) * const_type_10_;
-    metal::float3 unnamed_3 = const_type_10_ * metal::float3(data.v3_);
+    metal::float3 unnamed_2 = metal::float3(data.v3_) * const_type_11_;
+    metal::float3 unnamed_3 = const_type_11_ * metal::float3(data.v3_);
     metal::float3 unnamed_4 = data.v3_ * 2.0;
     metal::float3 unnamed_5 = 2.0 * data.v3_;
 }
@@ -55,21 +55,23 @@ kernel void main_(
 , device type_6 const& dummy [[user(fake0)]]
 , constant type_8& float_vecs [[user(fake0)]]
 , constant metal::float4& global_vec [[user(fake0)]]
+, constant metal::float4x4& global_mat [[user(fake0)]]
 , constant _mslBufferSizes& _buffer_sizes [[user(fake0)]]
 ) {
     float Foo_1 = 1.0;
     bool at = true;
     test_msl_packed_vec3_(alignment);
-    float _e10 = global_vec.x;
-    wg.inner[6] = _e10;
-    float _e16 = dummy[1].y;
-    wg.inner[5] = _e16;
-    float _e22 = float_vecs.inner[0].w;
-    wg.inner[4] = _e22;
-    float _e26 = alignment.v1_;
-    wg.inner[3] = _e26;
-    float _e31 = alignment.v3_[0];
-    wg.inner[2] = _e31;
+    metal::float4x4 _e10 = global_mat;
+    metal::float4 _e11 = global_vec;
+    wg.inner[6] = (_e10 * _e11).x;
+    float _e19 = dummy[1].y;
+    wg.inner[5] = _e19;
+    float _e25 = float_vecs.inner[0].w;
+    wg.inner[4] = _e25;
+    float _e29 = alignment.v1_;
+    wg.inner[3] = _e29;
+    float _e34 = alignment.v3_[0];
+    wg.inner[2] = _e34;
     alignment.v1_ = 4.0;
     wg.inner[1] = static_cast<float>(1 + (_buffer_sizes.size3 - 0 - 8) / 8);
     metal::atomic_store_explicit(&at_1, 2u, metal::memory_order_relaxed);