[naga msl-out] Emit and init struct member padding always

CHANGELOG.md

@@ -83,6 +83,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());
@@ -375,7 +370,7 @@ fn should_pack_struct_member(
         crate::TypeInner::Vector {
             size: crate::VectorSize::Tri,
             scalar: scalar @ crate::Scalar { width: 4, .. },
-        } if member.offset & 0xF != 0 || is_tight => Some(scalar),
+        } if is_tight => Some(scalar),
         _ => None,
     }
 }
@@ -3307,8 +3302,7 @@ impl<W: Write> Writer<W> {
                     writeln!(self.out, "struct {name} {{")?;
                     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 +4269,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/in/struct-layout.wgsl

@@ -0,0 +1,50 @@
+// Create several type definitions to test `align` and `size` layout.
+
+struct NoPadding {
+	@location(0)
+	v3: vec3f, // align 16, size 12; no start padding needed
+	@location(1)
+	f3: f32, // align 4, size 4; no start padding needed
+}
+@fragment
+fn no_padding_frag(input: NoPadding) -> @location(0) vec4f {
+	_ = input;
+	return vec4f(0.0);
+}
+@vertex
+fn no_padding_vert(input: NoPadding) -> @builtin(position) vec4f {
+	_ = input;
+	return vec4f(0.0);
+}
+@group(0) @binding(0) var<uniform> no_padding_uniform: NoPadding;
+@group(0) @binding(1) var<storage, read_write> no_padding_storage: NoPadding;
+@compute @workgroup_size(16,1,1)
+fn no_padding_comp() {
+	var x: NoPadding;
+	x = no_padding_uniform;
+	x = no_padding_storage;
+}
+
+struct NeedsPadding {
+	@location(0) f3_forces_padding: f32, // align 4, size 4; no start padding needed
+	@location(1) v3_needs_padding: vec3f, // align 16, size 12; needs 12 bytes of padding
+	@location(2) f3: f32, // align 4, size 4; no start padding needed
+}
+@fragment
+fn needs_padding_frag(input: NeedsPadding) -> @location(0) vec4f {
+	_ = input;
+	return vec4f(0.0);
+}
+@vertex
+fn needs_padding_vert(input: NeedsPadding) -> @builtin(position) vec4f {
+	_ = input;
+	return vec4f(0.0);
+}
+@group(0) @binding(2) var<uniform> needs_padding_uniform: NeedsPadding;
+@group(0) @binding(3) var<storage, read_write> needs_padding_storage: NeedsPadding;
+@compute @workgroup_size(16,1,1)
+fn needs_padding_comp() {
+	var x: NeedsPadding;
+	x = needs_padding_uniform;
+	x = needs_padding_storage;
+}

naga/tests/out/glsl/struct-layout.needs_padding_comp.Compute.glsl

@@ -0,0 +1,30 @@
+#version 310 es
+
+precision highp float;
+precision highp int;
+
+layout(local_size_x = 16, local_size_y = 1, local_size_z = 1) in;
+
+struct NoPadding {
+    vec3 v3_;
+    float f3_;
+};
+struct NeedsPadding {
+    float f3_forces_padding;
+    vec3 v3_needs_padding;
+    float f3_;
+};
+uniform NeedsPadding_block_0Compute { NeedsPadding _group_0_binding_2_cs; };
+
+layout(std430) buffer NeedsPadding_block_1Compute { NeedsPadding _group_0_binding_3_cs; };
+
+
+void main() {
+    NeedsPadding x_1 = NeedsPadding(0.0, vec3(0.0), 0.0);
+    NeedsPadding _e2 = _group_0_binding_2_cs;
+    x_1 = _e2;
+    NeedsPadding _e4 = _group_0_binding_3_cs;
+    x_1 = _e4;
+    return;
+}
+

naga/tests/out/glsl/struct-layout.needs_padding_frag.Fragment.glsl

@@ -0,0 +1,25 @@
+#version 310 es
+
+precision highp float;
+precision highp int;
+
+struct NoPadding {
+    vec3 v3_;
+    float f3_;
+};
+struct NeedsPadding {
+    float f3_forces_padding;
+    vec3 v3_needs_padding;
+    float f3_;
+};
+layout(location = 0) smooth in float _vs2fs_location0;
+layout(location = 1) smooth in vec3 _vs2fs_location1;
+layout(location = 2) smooth in float _vs2fs_location2;
+layout(location = 0) out vec4 _fs2p_location0;
+
+void main() {
+    NeedsPadding input_2 = NeedsPadding(_vs2fs_location0, _vs2fs_location1, _vs2fs_location2);
+    _fs2p_location0 = vec4(0.0);
+    return;
+}
+

naga/tests/out/glsl/struct-layout.needs_padding_vert.Vertex.glsl

@@ -0,0 +1,25 @@
+#version 310 es
+
+precision highp float;
+precision highp int;
+
+struct NoPadding {
+    vec3 v3_;
+    float f3_;
+};
+struct NeedsPadding {
+    float f3_forces_padding;
+    vec3 v3_needs_padding;
+    float f3_;
+};
+layout(location = 0) in float _p2vs_location0;
+layout(location = 1) in vec3 _p2vs_location1;
+layout(location = 2) in float _p2vs_location2;
+
+void main() {
+    NeedsPadding input_3 = NeedsPadding(_p2vs_location0, _p2vs_location1, _p2vs_location2);
+    gl_Position = vec4(0.0);
+    gl_Position.yz = vec2(-gl_Position.y, gl_Position.z * 2.0 - gl_Position.w);
+    return;
+}
+

naga/tests/out/glsl/struct-layout.no_padding_comp.Compute.glsl

@@ -0,0 +1,30 @@
+#version 310 es
+
+precision highp float;
+precision highp int;
+
+layout(local_size_x = 16, local_size_y = 1, local_size_z = 1) in;
+
+struct NoPadding {
+    vec3 v3_;
+    float f3_;
+};
+struct NeedsPadding {
+    float f3_forces_padding;
+    vec3 v3_needs_padding;
+    float f3_;
+};
+uniform NoPadding_block_0Compute { NoPadding _group_0_binding_0_cs; };
+
+layout(std430) buffer NoPadding_block_1Compute { NoPadding _group_0_binding_1_cs; };
+
+
+void main() {
+    NoPadding x = NoPadding(vec3(0.0), 0.0);
+    NoPadding _e2 = _group_0_binding_0_cs;
+    x = _e2;
+    NoPadding _e4 = _group_0_binding_1_cs;
+    x = _e4;
+    return;
+}
+

naga/tests/out/glsl/struct-layout.no_padding_frag.Fragment.glsl

@@ -0,0 +1,24 @@
+#version 310 es
+
+precision highp float;
+precision highp int;
+
+struct NoPadding {
+    vec3 v3_;
+    float f3_;
+};
+struct NeedsPadding {
+    float f3_forces_padding;
+    vec3 v3_needs_padding;
+    float f3_;
+};
+layout(location = 0) smooth in vec3 _vs2fs_location0;
+layout(location = 1) smooth in float _vs2fs_location1;
+layout(location = 0) out vec4 _fs2p_location0;
+
+void main() {
+    NoPadding input_ = NoPadding(_vs2fs_location0, _vs2fs_location1);
+    _fs2p_location0 = vec4(0.0);
+    return;
+}
+

naga/tests/out/glsl/struct-layout.no_padding_vert.Vertex.glsl

@@ -0,0 +1,24 @@
+#version 310 es
+
+precision highp float;
+precision highp int;
+
+struct NoPadding {
+    vec3 v3_;
+    float f3_;
+};
+struct NeedsPadding {
+    float f3_forces_padding;
+    vec3 v3_needs_padding;
+    float f3_;
+};
+layout(location = 0) in vec3 _p2vs_location0;
+layout(location = 1) in float _p2vs_location1;
+
+void main() {
+    NoPadding input_1 = NoPadding(_p2vs_location0, _p2vs_location1);
+    gl_Position = vec4(0.0);
+    gl_Position.yz = vec2(-gl_Position.y, gl_Position.z * 2.0 - gl_Position.w);
+    return;
+}
+

naga/tests/out/hlsl/struct-layout.hlsl

@@ -0,0 +1,87 @@
+struct NoPadding {
+    float3 v3_ : LOC0;
+    float f3_ : LOC1;
+};
+
+struct NeedsPadding {
+    float f3_forces_padding : LOC0;
+    float3 v3_needs_padding : LOC1;
+    float f3_ : LOC2;
+};
+
+cbuffer no_padding_uniform : register(b0) { NoPadding no_padding_uniform; }
+RWByteAddressBuffer no_padding_storage : register(u1);
+cbuffer needs_padding_uniform : register(b2) { NeedsPadding needs_padding_uniform; }
+RWByteAddressBuffer needs_padding_storage : register(u3);
+
+struct FragmentInput_no_padding_frag {
+    float3 v3_ : LOC0;
+    float f3_ : LOC1;
+};
+
+struct FragmentInput_needs_padding_frag {
+    float f3_forces_padding : LOC0;
+    float3 v3_needs_padding : LOC1;
+    float f3_1 : LOC2;
+};
+
+float4 no_padding_frag(FragmentInput_no_padding_frag fragmentinput_no_padding_frag) : SV_Target0
+{
+    NoPadding input = { fragmentinput_no_padding_frag.v3_, fragmentinput_no_padding_frag.f3_ };
+    return (0.0).xxxx;
+}
+
+float4 no_padding_vert(NoPadding input_1) : SV_Position
+{
+    return (0.0).xxxx;
+}
+
+NoPadding ConstructNoPadding(float3 arg0, float arg1) {
+    NoPadding ret = (NoPadding)0;
+    ret.v3_ = arg0;
+    ret.f3_ = arg1;
+    return ret;
+}
+
+[numthreads(16, 1, 1)]
+void no_padding_comp()
+{
+    NoPadding x = (NoPadding)0;
+
+    NoPadding _expr2 = no_padding_uniform;
+    x = _expr2;
+    NoPadding _expr4 = ConstructNoPadding(asfloat(no_padding_storage.Load3(0)), asfloat(no_padding_storage.Load(12)));
+    x = _expr4;
+    return;
+}
+
+float4 needs_padding_frag(FragmentInput_needs_padding_frag fragmentinput_needs_padding_frag) : SV_Target0
+{
+    NeedsPadding input_2 = { fragmentinput_needs_padding_frag.f3_forces_padding, fragmentinput_needs_padding_frag.v3_needs_padding, fragmentinput_needs_padding_frag.f3_1 };
+    return (0.0).xxxx;
+}
+
+float4 needs_padding_vert(NeedsPadding input_3) : SV_Position
+{
+    return (0.0).xxxx;
+}
+
+NeedsPadding ConstructNeedsPadding(float arg0, float3 arg1, float arg2) {
+    NeedsPadding ret = (NeedsPadding)0;
+    ret.f3_forces_padding = arg0;
+    ret.v3_needs_padding = arg1;
+    ret.f3_ = arg2;
+    return ret;
+}
+
+[numthreads(16, 1, 1)]
+void needs_padding_comp()
+{
+    NeedsPadding x_1 = (NeedsPadding)0;
+
+    NeedsPadding _expr2 = needs_padding_uniform;
+    x_1 = _expr2;
+    NeedsPadding _expr4 = ConstructNeedsPadding(asfloat(needs_padding_storage.Load(0)), asfloat(needs_padding_storage.Load3(16)), asfloat(needs_padding_storage.Load(28)));
+    x_1 = _expr4;
+    return;
+}

naga/tests/out/hlsl/struct-layout.ron

@@ -0,0 +1,32 @@
+(
+    vertex:[
+        (
+            entry_point:"no_padding_vert",
+            target_profile:"vs_5_1",
+        ),
+        (
+            entry_point:"needs_padding_vert",
+            target_profile:"vs_5_1",
+        ),
+    ],
+    fragment:[
+        (
+            entry_point:"no_padding_frag",
+            target_profile:"ps_5_1",
+        ),
+        (
+            entry_point:"needs_padding_frag",
+            target_profile:"ps_5_1",
+        ),
+    ],
+    compute:[
+        (
+            entry_point:"no_padding_comp",
+            target_profile:"cs_5_1",
+        ),
+        (
+            entry_point:"needs_padding_comp",
+            target_profile:"cs_5_1",
+        ),
+    ],
+)