fix(msl-out): emit and init struct member padding always

Previously, implicit padding members of `struct`s were suppressed from
structure definitions in Metal output if they had a binding specified
for them (i.e., `@location(0)`). This padding is, however, is necessary
for correct access of member fields passed into shaders by uniform and
storage buffers. Unconditionally emit padding members for `struct`s.

Resolves
[`gfx-rs/wgpu`#4701](#4701).

CHANGELOG.md

@@ -93,6 +93,8 @@ Passing an owned value `window` to `Surface` will return a `Surface<'static>`. S
 - Implement WGSL abstract types (by @jimblandy):
   - Add a new `naga::Literal` variant, `I64`, for signed 64-bit literals. [#4711](https://github.com/gfx-rs/wgpu/pull/4711)
 
+- Emit and init `struct` member padding always. By @ErichDonGubler in [#4701](https://github.com/gfx-rs/wgpu/pull/4701).
+
 ### Bug Fixes
 
 #### WGL

naga/src/back/msl/writer.rs

@@ -357,11 +357,6 @@ fn should_pack_struct_member(
     module: &crate::Module,
 ) -> Option<crate::Scalar> {
     let member = &members[index];
-    //Note: this is imperfect - the same structure can be used for host-shared
-    // things, where packed float would matter.
-    if member.binding.is_some() {
-        return None;
-    }
 
     let ty_inner = &module.types[member.ty].inner;
     let last_offset = member.offset + ty_inner.size(module.to_ctx());
@@ -3308,7 +3303,7 @@ impl<W: Write> Writer<W> {
                     let mut last_offset = 0;
                     for (index, member) in members.iter().enumerate() {
                         // quick and dirty way to figure out if we need this...
-                        if member.binding.is_none() && member.offset > last_offset {
+                        if member.offset > last_offset {
                             self.struct_member_pads.insert((handle, index as u32));
                             let pad = member.offset - last_offset;
                             writeln!(self.out, "{}char _pad{}[{}];", back::INDENT, index, pad)?;
@@ -4275,6 +4270,13 @@ impl<W: Write> Writer<W> {
                             if member_index != 0 {
                                 write!(self.out, ", ")?;
                             }
+                            // insert padding initialization, if needed
+                            if self
+                                .struct_member_pads
+                                .contains(&(arg.ty, member_index as u32))
+                            {
+                                write!(self.out, "{{}}, ")?;
+                            }
                             if let Some(crate::Binding::Location { .. }) = member.binding {
                                 write!(self.out, "{varyings_member_name}.")?;
                             }

naga/tests/out/msl/quad.msl

@@ -6,6 +6,7 @@ using metal::uint;
 
 struct VertexOutput {
     metal::float2 uv;
+    char _pad1[8];
     metal::float4 position;
 };
 constant float c_scale = 1.2;
@@ -23,7 +24,7 @@ vertex vert_mainOutput vert_main(
 ) {
     const auto pos = varyings.pos;
     const auto uv = varyings.uv;
-    const auto _tmp = VertexOutput {uv, metal::float4(c_scale * pos, 0.0, 1.0)};
+    const auto _tmp = VertexOutput {uv, {}, metal::float4(c_scale * pos, 0.0, 1.0)};
     return vert_mainOutput { _tmp.uv, _tmp.position };
 }