diff --git a/reference/opt/shaders-msl/asm/comp/block-name-alias-global.asm.comp b/reference/opt/shaders-msl/asm/comp/block-name-alias-global.asm.comp
index 6dcc14ea8..0136d13b3 100644
--- a/reference/opt/shaders-msl/asm/comp/block-name-alias-global.asm.comp
+++ b/reference/opt/shaders-msl/asm/comp/block-name-alias-global.asm.comp
@@ -1,8 +1,13 @@
+#pragma clang diagnostic ignored "-Wmissing-prototypes"
+
 #include <metal_stdlib>
 #include <simd/simd.h>
 
 using namespace metal;
 
+template <typename T, int stride>
+struct spvPaddedArrayElement { T data; char padding[stride - sizeof(T)]; };
+
 struct A
 {
     int a;
@@ -18,12 +23,11 @@ struct A_2
 {
     int a;
     int b;
-    char _m0_final_padding[8];
 };
 
 struct A_3
 {
-    A_2 Data[1024];
+    spvPaddedArrayElement<A_2, 16> Data[1024];
 };
 
 struct B
@@ -33,14 +37,14 @@ struct B
 
 struct B_1
 {
-    A_2 Data[1024];
+    spvPaddedArrayElement<A_2, 16> Data[1024];
 };
 
 kernel void main0(device A_1& C1 [[buffer(0)]], constant A_3& C2 [[buffer(1)]], device B& C3 [[buffer(2)]], constant B_1& C4 [[buffer(3)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
 {
-    C1.Data[gl_GlobalInvocationID.x].a = C2.Data[gl_GlobalInvocationID.x].a;
-    C1.Data[gl_GlobalInvocationID.x].b = C2.Data[gl_GlobalInvocationID.x].b;
-    C3.Data[gl_GlobalInvocationID.x].a = C4.Data[gl_GlobalInvocationID.x].a;
-    C3.Data[gl_GlobalInvocationID.x].b = C4.Data[gl_GlobalInvocationID.x].b;
+    C1.Data[gl_GlobalInvocationID.x].a = C2.Data[gl_GlobalInvocationID.x].data.a;
+    C1.Data[gl_GlobalInvocationID.x].b = C2.Data[gl_GlobalInvocationID.x].data.b;
+    C3.Data[gl_GlobalInvocationID.x].a = C4.Data[gl_GlobalInvocationID.x].data.a;
+    C3.Data[gl_GlobalInvocationID.x].b = C4.Data[gl_GlobalInvocationID.x].data.b;
 }
 
diff --git a/reference/opt/shaders-msl/comp/struct-packing.comp b/reference/opt/shaders-msl/comp/struct-packing.comp
index dc1654399..4fe733a38 100644
--- a/reference/opt/shaders-msl/comp/struct-packing.comp
+++ b/reference/opt/shaders-msl/comp/struct-packing.comp
@@ -1,13 +1,17 @@
+#pragma clang diagnostic ignored "-Wmissing-prototypes"
+
 #include <metal_stdlib>
 #include <simd/simd.h>
 
 using namespace metal;
 
+template <typename T, int stride>
+struct spvPaddedArrayElement { T data; char padding[stride - sizeof(T)]; };
+
 struct S0
 {
     float2 a[1];
     float b;
-    char _m0_final_padding[4];
 };
 
 struct S1
@@ -20,7 +24,6 @@ struct S2
 {
     float3 a[1];
     float b;
-    char _m0_final_padding[12];
 };
 
 struct S3
@@ -45,7 +48,6 @@ struct Content
     S3 m3;
     float m4;
     S4 m3s[8];
-    char _m0_final_padding[8];
 };
 
 struct SSBO1
@@ -69,7 +71,6 @@ struct S0_1
     float2 a[1];
     char _m1_pad[8];
     float b;
-    char _m0_final_padding[12];
 };
 
 struct S1_1
@@ -82,7 +83,6 @@ struct S2_1
 {
     float3 a[1];
     float b;
-    char _m0_final_padding[12];
 };
 
 struct S3_1
@@ -94,21 +94,21 @@ struct S3_1
 struct S4_1
 {
     float2 c;
-    char _m0_final_padding[8];
 };
 
 struct Content_1
 {
-    S0_1 m0s[1];
+    spvPaddedArrayElement<S0_1, 32> m0s[1];
     S1_1 m1s[1];
     S2_1 m2s[1];
     S0_1 m0;
+    char _m4_pad[8];
     S1_1 m1;
     S2_1 m2;
     S3_1 m3;
     float m4;
     char _m8_pad[8];
-    S4_1 m3s[8];
+    spvPaddedArrayElement<S4_1, 16> m3s[8];
 };
 
 struct SSBO0
@@ -124,8 +124,8 @@ constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u);
 kernel void main0(device SSBO1& ssbo_430 [[buffer(0)]], device SSBO0& ssbo_140 [[buffer(1)]])
 {
     Content_1 _60 = ssbo_140.content;
-    ssbo_430.content.m0s[0].a[0] = _60.m0s[0].a[0];
-    ssbo_430.content.m0s[0].b = _60.m0s[0].b;
+    ssbo_430.content.m0s[0].a[0] = _60.m0s[0].data.a[0];
+    ssbo_430.content.m0s[0].b = _60.m0s[0].data.b;
     ssbo_430.content.m1s[0].a = float3(_60.m1s[0].a);
     ssbo_430.content.m1s[0].b = _60.m1s[0].b;
     ssbo_430.content.m2s[0].a[0] = _60.m2s[0].a[0];
@@ -139,14 +139,14 @@ kernel void main0(device SSBO1& ssbo_430 [[buffer(0)]], device SSBO0& ssbo_140 [
     ssbo_430.content.m3.a = _60.m3.a;
     ssbo_430.content.m3.b = _60.m3.b;
     ssbo_430.content.m4 = _60.m4;
-    ssbo_430.content.m3s[0].c = _60.m3s[0].c;
-    ssbo_430.content.m3s[1].c = _60.m3s[1].c;
-    ssbo_430.content.m3s[2].c = _60.m3s[2].c;
-    ssbo_430.content.m3s[3].c = _60.m3s[3].c;
-    ssbo_430.content.m3s[4].c = _60.m3s[4].c;
-    ssbo_430.content.m3s[5].c = _60.m3s[5].c;
-    ssbo_430.content.m3s[6].c = _60.m3s[6].c;
-    ssbo_430.content.m3s[7].c = _60.m3s[7].c;
+    ssbo_430.content.m3s[0].c = _60.m3s[0].data.c;
+    ssbo_430.content.m3s[1].c = _60.m3s[1].data.c;
+    ssbo_430.content.m3s[2].c = _60.m3s[2].data.c;
+    ssbo_430.content.m3s[3].c = _60.m3s[3].data.c;
+    ssbo_430.content.m3s[4].c = _60.m3s[4].data.c;
+    ssbo_430.content.m3s[5].c = _60.m3s[5].data.c;
+    ssbo_430.content.m3s[6].c = _60.m3s[6].data.c;
+    ssbo_430.content.m3s[7].c = _60.m3s[7].data.c;
     ssbo_430.content.m1.a = ssbo_430.content.m3.a * ssbo_430.m6[1][1];
 }
 
diff --git a/reference/shaders-msl-no-opt/asm/comp/aliased-struct-divergent-member-name.asm.comp b/reference/shaders-msl-no-opt/asm/comp/aliased-struct-divergent-member-name.asm.comp
index 4151832e8..4e38689ca 100644
--- a/reference/shaders-msl-no-opt/asm/comp/aliased-struct-divergent-member-name.asm.comp
+++ b/reference/shaders-msl-no-opt/asm/comp/aliased-struct-divergent-member-name.asm.comp
@@ -1,8 +1,13 @@
+#pragma clang diagnostic ignored "-Wmissing-prototypes"
+
 #include <metal_stdlib>
 #include <simd/simd.h>
 
 using namespace metal;
 
+template <typename T, int stride>
+struct spvPaddedArrayElement { T data; char padding[stride - sizeof(T)]; };
+
 struct T
 {
     float a;
@@ -21,18 +26,17 @@ struct SSBO1
 struct T_2
 {
     float c;
-    char _m0_final_padding[12];
 };
 
 struct SSBO2
 {
-    T_2 bar[1];
+    spvPaddedArrayElement<T_2, 16> bar[1];
 };
 
 kernel void main0(device SSBO1& _9 [[buffer(0)]], device SSBO2& _13 [[buffer(1)]])
 {
     T v = T{ 40.0 };
     _9.foo[10].b = v.a;
-    _13.bar[30].c = v.a;
+    _13.bar[30].data.c = v.a;
 }
 
diff --git a/reference/shaders-msl-no-opt/comp/struct-packing-scalar.nocompat.invalid.vk.comp b/reference/shaders-msl-no-opt/comp/struct-packing-scalar.nocompat.invalid.vk.comp
index a0bb9c10f..0e22b9f11 100644
--- a/reference/shaders-msl-no-opt/comp/struct-packing-scalar.nocompat.invalid.vk.comp
+++ b/reference/shaders-msl-no-opt/comp/struct-packing-scalar.nocompat.invalid.vk.comp
@@ -1,3 +1,5 @@
+#pragma clang diagnostic ignored "-Wmissing-prototypes"
+
 #include <metal_stdlib>
 #include <simd/simd.h>
 
@@ -7,6 +9,9 @@ typedef packed_float3 packed_float2x3[2];
 typedef packed_float3 packed_rm_float3x2[2];
 typedef packed_float2 packed_float2x2[2];
 
+template <typename T, int stride>
+struct spvPaddedArrayElement { T data; char padding[stride - sizeof(T)]; };
+
 struct S0
 {
     packed_float2 a[1];
@@ -64,7 +69,6 @@ struct S0_1
     float2 a[1];
     char _m1_pad[8];
     float b;
-    char _m0_final_padding[12];
 };
 
 struct S1_1
@@ -77,7 +81,6 @@ struct S2_1
 {
     float3 a[1];
     float b;
-    char _m0_final_padding[12];
 };
 
 struct S3_1
@@ -88,15 +91,15 @@ struct S3_1
 
 struct Content_1
 {
-    S0_1 m0s[1];
+    spvPaddedArrayElement<S0_1, 32> m0s[1];
     S1_1 m1s[1];
     S2_1 m2s[1];
     S0_1 m0;
+    char _m4_pad[8];
     S1_1 m1;
     S2_1 m2;
     S3_1 m3;
     float m4;
-    char _m0_final_padding[12];
 };
 
 struct SSBO0
@@ -126,8 +129,8 @@ constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u);
 
 kernel void main0(device SSBO1& __restrict ssbo_scalar [[buffer(0)]], device SSBO0& __restrict ssbo_140 [[buffer(1)]], device SSBO2& __restrict ssbo_scalar2 [[buffer(2)]])
 {
-    ssbo_scalar.content.m0s[0].a[0] = ssbo_140.content.m0s[0].a[0];
-    ssbo_scalar.content.m0s[0].b = ssbo_140.content.m0s[0].b;
+    ssbo_scalar.content.m0s[0].a[0] = ssbo_140.content.m0s[0].data.a[0];
+    ssbo_scalar.content.m0s[0].b = ssbo_140.content.m0s[0].data.b;
     ssbo_scalar.content.m1s[0].a = float3(ssbo_140.content.m1s[0].a);
     ssbo_scalar.content.m1s[0].b = ssbo_140.content.m1s[0].b;
     ssbo_scalar.content.m2s[0].a[0] = ssbo_140.content.m2s[0].a[0];
diff --git a/reference/shaders-msl-no-opt/packing/struct-size-padding-array-of-array.comp b/reference/shaders-msl-no-opt/packing/struct-size-padding-array-of-array.comp
index c30fd070e..e47bcad25 100644
--- a/reference/shaders-msl-no-opt/packing/struct-size-padding-array-of-array.comp
+++ b/reference/shaders-msl-no-opt/packing/struct-size-padding-array-of-array.comp
@@ -1,18 +1,21 @@
+#pragma clang diagnostic ignored "-Wmissing-prototypes"
+
 #include <metal_stdlib>
 #include <simd/simd.h>
 
 using namespace metal;
 
+template <typename T, int stride>
+struct spvPaddedArrayElement { T data; char padding[stride - sizeof(T)]; };
+
 struct A
 {
     float v;
-    char _m0_final_padding[12];
 };
 
 struct B
 {
     float2 v;
-    char _m0_final_padding[8];
 };
 
 struct C
@@ -29,13 +32,12 @@ struct E
 {
     float4 a;
     float2 b;
-    char _m0_final_padding[8];
 };
 
 struct SSBO
 {
-    A a[2][4];
-    B b[2][4];
+    spvPaddedArrayElement<A, 16> a[2][4];
+    spvPaddedArrayElement<B, 16> b[2][4];
     C c[2][4];
     D d[2][4];
     float2x4 e[2][4];
diff --git a/reference/shaders-msl-no-opt/packing/struct-size-padding.comp b/reference/shaders-msl-no-opt/packing/struct-size-padding.comp
index 98f039fc9..9da3dfbbe 100644
--- a/reference/shaders-msl-no-opt/packing/struct-size-padding.comp
+++ b/reference/shaders-msl-no-opt/packing/struct-size-padding.comp
@@ -1,18 +1,21 @@
+#pragma clang diagnostic ignored "-Wmissing-prototypes"
+
 #include <metal_stdlib>
 #include <simd/simd.h>
 
 using namespace metal;
 
+template <typename T, int stride>
+struct spvPaddedArrayElement { T data; char padding[stride - sizeof(T)]; };
+
 struct A
 {
     float v;
-    char _m0_final_padding[12];
 };
 
 struct B
 {
     float2 v;
-    char _m0_final_padding[8];
 };
 
 struct C
@@ -29,13 +32,12 @@ struct E
 {
     float4 a;
     float2 b;
-    char _m0_final_padding[8];
 };
 
 struct SSBO
 {
-    A a[4];
-    B b[4];
+    spvPaddedArrayElement<A, 16> a[4];
+    spvPaddedArrayElement<B, 16> b[4];
     C c[4];
     D d[4];
     float2x4 e[4];
diff --git a/reference/shaders-msl/asm/comp/block-name-alias-global.asm.comp b/reference/shaders-msl/asm/comp/block-name-alias-global.asm.comp
index 6dcc14ea8..0136d13b3 100644
--- a/reference/shaders-msl/asm/comp/block-name-alias-global.asm.comp
+++ b/reference/shaders-msl/asm/comp/block-name-alias-global.asm.comp
@@ -1,8 +1,13 @@
+#pragma clang diagnostic ignored "-Wmissing-prototypes"
+
 #include <metal_stdlib>
 #include <simd/simd.h>
 
 using namespace metal;
 
+template <typename T, int stride>
+struct spvPaddedArrayElement { T data; char padding[stride - sizeof(T)]; };
+
 struct A
 {
     int a;
@@ -18,12 +23,11 @@ struct A_2
 {
     int a;
     int b;
-    char _m0_final_padding[8];
 };
 
 struct A_3
 {
-    A_2 Data[1024];
+    spvPaddedArrayElement<A_2, 16> Data[1024];
 };
 
 struct B
@@ -33,14 +37,14 @@ struct B
 
 struct B_1
 {
-    A_2 Data[1024];
+    spvPaddedArrayElement<A_2, 16> Data[1024];
 };
 
 kernel void main0(device A_1& C1 [[buffer(0)]], constant A_3& C2 [[buffer(1)]], device B& C3 [[buffer(2)]], constant B_1& C4 [[buffer(3)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
 {
-    C1.Data[gl_GlobalInvocationID.x].a = C2.Data[gl_GlobalInvocationID.x].a;
-    C1.Data[gl_GlobalInvocationID.x].b = C2.Data[gl_GlobalInvocationID.x].b;
-    C3.Data[gl_GlobalInvocationID.x].a = C4.Data[gl_GlobalInvocationID.x].a;
-    C3.Data[gl_GlobalInvocationID.x].b = C4.Data[gl_GlobalInvocationID.x].b;
+    C1.Data[gl_GlobalInvocationID.x].a = C2.Data[gl_GlobalInvocationID.x].data.a;
+    C1.Data[gl_GlobalInvocationID.x].b = C2.Data[gl_GlobalInvocationID.x].data.b;
+    C3.Data[gl_GlobalInvocationID.x].a = C4.Data[gl_GlobalInvocationID.x].data.a;
+    C3.Data[gl_GlobalInvocationID.x].b = C4.Data[gl_GlobalInvocationID.x].data.b;
 }
 
diff --git a/reference/shaders-msl/comp/struct-packing.comp b/reference/shaders-msl/comp/struct-packing.comp
index dc1654399..4fe733a38 100644
--- a/reference/shaders-msl/comp/struct-packing.comp
+++ b/reference/shaders-msl/comp/struct-packing.comp
@@ -1,13 +1,17 @@
+#pragma clang diagnostic ignored "-Wmissing-prototypes"
+
 #include <metal_stdlib>
 #include <simd/simd.h>
 
 using namespace metal;
 
+template <typename T, int stride>
+struct spvPaddedArrayElement { T data; char padding[stride - sizeof(T)]; };
+
 struct S0
 {
     float2 a[1];
     float b;
-    char _m0_final_padding[4];
 };
 
 struct S1
@@ -20,7 +24,6 @@ struct S2
 {
     float3 a[1];
     float b;
-    char _m0_final_padding[12];
 };
 
 struct S3
@@ -45,7 +48,6 @@ struct Content
     S3 m3;
     float m4;
     S4 m3s[8];
-    char _m0_final_padding[8];
 };
 
 struct SSBO1
@@ -69,7 +71,6 @@ struct S0_1
     float2 a[1];
     char _m1_pad[8];
     float b;
-    char _m0_final_padding[12];
 };
 
 struct S1_1
@@ -82,7 +83,6 @@ struct S2_1
 {
     float3 a[1];
     float b;
-    char _m0_final_padding[12];
 };
 
 struct S3_1
@@ -94,21 +94,21 @@ struct S3_1
 struct S4_1
 {
     float2 c;
-    char _m0_final_padding[8];
 };
 
 struct Content_1
 {
-    S0_1 m0s[1];
+    spvPaddedArrayElement<S0_1, 32> m0s[1];
     S1_1 m1s[1];
     S2_1 m2s[1];
     S0_1 m0;
+    char _m4_pad[8];
     S1_1 m1;
     S2_1 m2;
     S3_1 m3;
     float m4;
     char _m8_pad[8];
-    S4_1 m3s[8];
+    spvPaddedArrayElement<S4_1, 16> m3s[8];
 };
 
 struct SSBO0
@@ -124,8 +124,8 @@ constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u);
 kernel void main0(device SSBO1& ssbo_430 [[buffer(0)]], device SSBO0& ssbo_140 [[buffer(1)]])
 {
     Content_1 _60 = ssbo_140.content;
-    ssbo_430.content.m0s[0].a[0] = _60.m0s[0].a[0];
-    ssbo_430.content.m0s[0].b = _60.m0s[0].b;
+    ssbo_430.content.m0s[0].a[0] = _60.m0s[0].data.a[0];
+    ssbo_430.content.m0s[0].b = _60.m0s[0].data.b;
     ssbo_430.content.m1s[0].a = float3(_60.m1s[0].a);
     ssbo_430.content.m1s[0].b = _60.m1s[0].b;
     ssbo_430.content.m2s[0].a[0] = _60.m2s[0].a[0];
@@ -139,14 +139,14 @@ kernel void main0(device SSBO1& ssbo_430 [[buffer(0)]], device SSBO0& ssbo_140 [
     ssbo_430.content.m3.a = _60.m3.a;
     ssbo_430.content.m3.b = _60.m3.b;
     ssbo_430.content.m4 = _60.m4;
-    ssbo_430.content.m3s[0].c = _60.m3s[0].c;
-    ssbo_430.content.m3s[1].c = _60.m3s[1].c;
-    ssbo_430.content.m3s[2].c = _60.m3s[2].c;
-    ssbo_430.content.m3s[3].c = _60.m3s[3].c;
-    ssbo_430.content.m3s[4].c = _60.m3s[4].c;
-    ssbo_430.content.m3s[5].c = _60.m3s[5].c;
-    ssbo_430.content.m3s[6].c = _60.m3s[6].c;
-    ssbo_430.content.m3s[7].c = _60.m3s[7].c;
+    ssbo_430.content.m3s[0].c = _60.m3s[0].data.c;
+    ssbo_430.content.m3s[1].c = _60.m3s[1].data.c;
+    ssbo_430.content.m3s[2].c = _60.m3s[2].data.c;
+    ssbo_430.content.m3s[3].c = _60.m3s[3].data.c;
+    ssbo_430.content.m3s[4].c = _60.m3s[4].data.c;
+    ssbo_430.content.m3s[5].c = _60.m3s[5].data.c;
+    ssbo_430.content.m3s[6].c = _60.m3s[6].data.c;
+    ssbo_430.content.m3s[7].c = _60.m3s[7].data.c;
     ssbo_430.content.m1.a = ssbo_430.content.m3.a * ssbo_430.m6[1][1];
 }
 
diff --git a/spirv_glsl.cpp b/spirv_glsl.cpp
index 90fa14be9..923636050 100644
--- a/spirv_glsl.cpp
+++ b/spirv_glsl.cpp
@@ -1335,9 +1335,6 @@ void CompilerGLSL::emit_struct(SPIRType &type)
 		emitted = true;
 	}
 
-	if (has_extended_decoration(type.self, SPIRVCrossDecorationPaddingTarget))
-		emit_struct_padding_target(type);
-
 	end_scope_decl();
 
 	if (emitted)
@@ -10872,6 +10869,15 @@ string CompilerGLSL::access_chain_internal(uint32_t base, const uint32_t *indice
 				access_meshlet_position_y = true;
 			}
 
+			if (get<SPIRType>(type->parent_type).op == OpTypeStruct &&
+			    has_decoration(type->parent_type, DecorationArrayStride))
+			{
+				uint32_t native_stride = get_decoration(type->parent_type, DecorationArrayStride);
+				uint32_t array_stride = get_decoration(type_id, DecorationArrayStride);
+				if (native_stride != array_stride)
+					expr += ".data";
+			}
+
 			type_id = type->parent_type;
 			type = &get<SPIRType>(type_id);
 
@@ -10957,6 +10963,7 @@ string CompilerGLSL::access_chain_internal(uint32_t base, const uint32_t *indice
 				physical_type = 0;
 
 			row_major_matrix_needs_conversion = member_is_non_native_row_major_matrix(*type, index);
+			type_id = type->member_types[index];
 			type = &get<SPIRType>(type->member_types[index]);
 		}
 		// Matrix -> Vector
@@ -11168,9 +11175,9 @@ string CompilerGLSL::to_flattened_struct_member(const string &basename, const SP
 	return ret;
 }
 
-uint32_t CompilerGLSL::get_physical_type_stride(const SPIRType &) const
+uint32_t CompilerGLSL::get_physical_type_id_stride(TypeID) const
 {
-	SPIRV_CROSS_THROW("Invalid to call get_physical_type_stride on a backend without native pointer support.");
+	SPIRV_CROSS_THROW("Invalid to call get_physical_type_id_stride on a backend without native pointer support.");
 }
 
 string CompilerGLSL::access_chain(uint32_t base, const uint32_t *indices, uint32_t count, const SPIRType &target_type,
@@ -11231,13 +11238,13 @@ string CompilerGLSL::access_chain(uint32_t base, const uint32_t *indices, uint32
 				// If there is a mismatch we have to go via 64-bit pointer arithmetic :'(
 				// Using packed hacks only gets us so far, and is not designed to deal with pointer to
 				// random values. It works for structs though.
-				auto &pointee_type = get_pointee_type(get<SPIRType>(type_id));
-				uint32_t physical_stride = get_physical_type_stride(pointee_type);
+				TypeID pointee_type_id = get_pointee_type_id(type_id);
+				uint32_t physical_stride = get_physical_type_id_stride(pointee_type_id);
 				uint32_t requested_stride = get_decoration(type_id, DecorationArrayStride);
 				if (physical_stride != requested_stride)
 				{
 					flags |= ACCESS_CHAIN_PTR_CHAIN_POINTER_ARITH_BIT;
-					if (is_vector(pointee_type))
+					if (is_vector(get<SPIRType>(pointee_type_id)))
 						flags |= ACCESS_CHAIN_PTR_CHAIN_CAST_TO_SCALAR_BIT;
 				}
 			}
@@ -16338,10 +16345,6 @@ void CompilerGLSL::emit_struct_member(const SPIRType &type, uint32_t member_type
 	          variable_decl(membertype, to_member_name(type, index)), ";");
 }
 
-void CompilerGLSL::emit_struct_padding_target(const SPIRType &)
-{
-}
-
 string CompilerGLSL::flags_to_qualifiers_glsl(const SPIRType &type, uint32_t id, const Bitset &flags)
 {
 	// GL_EXT_buffer_reference variables can be marked as restrict.
diff --git a/spirv_glsl.hpp b/spirv_glsl.hpp
index 98e93ee99..709b3ea2b 100644
--- a/spirv_glsl.hpp
+++ b/spirv_glsl.hpp
@@ -447,7 +447,6 @@ class CompilerGLSL : public Compiler
 	virtual std::string builtin_to_glsl(BuiltIn builtin, StorageClass storage);
 	virtual void emit_struct_member(const SPIRType &type, uint32_t member_type_id, uint32_t index,
 	                                const std::string &qualifier = "", uint32_t base_offset = 0);
-	virtual void emit_struct_padding_target(const SPIRType &type);
 	virtual std::string image_type_glsl(const SPIRType &type, uint32_t id = 0, bool member = false);
 	std::string constant_expression(const SPIRConstant &c,
 	                                bool inside_block_like_struct_scope = false,
@@ -773,7 +772,7 @@ class CompilerGLSL : public Compiler
 
 	// Only meaningful on backends with physical pointer support ala MSL.
 	// Relevant for PtrAccessChain / BDA.
-	virtual uint32_t get_physical_type_stride(const SPIRType &type) const;
+	virtual uint32_t get_physical_type_id_stride(TypeID type_id) const;
 
 	StorageClass get_expression_effective_storage_class(uint32_t ptr);
 	virtual bool access_chain_needs_stage_io_builtin_translation(uint32_t base);
diff --git a/spirv_msl.cpp b/spirv_msl.cpp
index 025427cc8..1291d5d22 100644
--- a/spirv_msl.cpp
+++ b/spirv_msl.cpp
@@ -5048,23 +5048,16 @@ void CompilerMSL::mark_scalar_layout_structs(const SPIRType &type)
 
 			if (struct_needs_explicit_padding)
 			{
-				msl_size = get_declared_struct_size_msl(*struct_type, true, true);
-				if (array_stride < msl_size)
+				msl_size = get_declared_struct_size_msl(*struct_type);
+
+				if (array_stride > msl_size)
 				{
-					SPIRV_CROSS_THROW("Cannot express an array stride smaller than size of struct type.");
-				}
-				else
-				{
-					if (has_extended_decoration(struct_type->self, SPIRVCrossDecorationPaddingTarget))
-					{
-						if (array_stride !=
-						    get_extended_decoration(struct_type->self, SPIRVCrossDecorationPaddingTarget))
-							SPIRV_CROSS_THROW(
-							    "A struct is used with different array strides. Cannot express this in MSL.");
-					}
-					else
-						set_extended_decoration(struct_type->self, SPIRVCrossDecorationPaddingTarget, array_stride);
+					set_decoration(struct_type->self, DecorationArrayStride, msl_size);
+					add_spv_func_and_recompile(SPVFuncImplPaddedArrayElement);
 				}
+
+				if (array_stride < msl_size)
+					SPIRV_CROSS_THROW("Cannot express an array stride smaller than size of struct type.");
 			}
 		}
 	}
@@ -5170,8 +5163,10 @@ bool CompilerMSL::validate_member_packing_rules_msl(const SPIRType &type, uint32
 		// If app tries to be cheeky and access the member out of bounds, this will not work, but this is the best we can do.
 		// In OpAccessChain with logical memory models, access chains must be in-bounds in SPIR-V specification.
 		bool relax_array_stride = mbr_type.array.back() == 1 && mbr_type.array_size_literal.back();
+		bool is_plain_struct = !mbr_type.pointer && mbr_type.basetype == SPIRType::Struct;
 
-		if (!relax_array_stride)
+		// Array of struct is padded on-demand.
+		if (!relax_array_stride && !is_plain_struct)
 		{
 			uint32_t spirv_array_stride = type_struct_member_array_stride(type, index);
 			uint32_t msl_array_stride = get_declared_struct_member_array_stride_msl(type, index);
@@ -8213,6 +8208,13 @@ void CompilerMSL::emit_custom_functions()
 			statement("");
 			break;
 
+		case SPVFuncImplPaddedArrayElement:
+			// .data is used in access chain.
+			statement("template <typename T, int stride>");
+			statement("struct spvPaddedArrayElement { T data; char padding[stride - sizeof(T)]; };");
+			statement("");
+			break;
+
 		case SPVFuncImplReduceAdd:
 			// Metal doesn't support __builtin_reduce_add or simd_reduce_add, so we need this.
 			// Metal also doesn't support the other vector builtins, which would have been useful to make this a single template.
@@ -13475,6 +13477,24 @@ string CompilerMSL::to_struct_member(const SPIRType &type, uint32_t member_type_
 	else
 		decl_type = type_to_glsl(*declared_type, orig_id, true);
 
+	if (physical_type.basetype == SPIRType::Struct &&
+	    has_decoration(physical_type.self, DecorationArrayStride) &&
+	    is_array(physical_type))
+	{
+		uint32_t native_stride = get_decoration(physical_type.self, DecorationArrayStride);
+		uint32_t array_stride = get_decoration(type.member_types[index], DecorationArrayStride);
+		auto *struct_array_type = &physical_type;
+
+		while (struct_array_type->parent_type && is_array(get<SPIRType>(struct_array_type->parent_type)))
+		{
+			array_stride = get_decoration(struct_array_type->parent_type, DecorationArrayStride);
+			struct_array_type = &get<SPIRType>(struct_array_type->parent_type);
+		}
+
+		if (array_stride != native_stride)
+			decl_type = join("spvPaddedArrayElement<", decl_type, ", ", array_stride, ">");
+	}
+
 	const char *overlapping_binding_tag =
 			has_extended_member_decoration(type.self, index, SPIRVCrossDecorationOverlappingBinding) ?
 			"// Overlapping binding: " : "";
@@ -13513,16 +13533,6 @@ void CompilerMSL::emit_struct_member(const SPIRType &type, uint32_t member_type_
 	builtin_declaration = false;
 }
 
-void CompilerMSL::emit_struct_padding_target(const SPIRType &type)
-{
-	uint32_t struct_size = get_declared_struct_size_msl(type, true, true);
-	uint32_t target_size = get_extended_decoration(type.self, SPIRVCrossDecorationPaddingTarget);
-	if (target_size < struct_size)
-		SPIRV_CROSS_THROW("Cannot pad with negative bytes.");
-	else if (target_size > struct_size)
-		statement("char _m0_final_padding[", target_size - struct_size, "];");
-}
-
 // Return a MSL qualifier for the specified function attribute member
 string CompilerMSL::member_attribute_qualifier(const SPIRType &type, uint32_t index)
 {
@@ -18273,17 +18283,17 @@ string CompilerMSL::built_in_func_arg(BuiltIn builtin, bool prefix_comma)
 	return bi_arg;
 }
 
-const SPIRType &CompilerMSL::get_physical_member_type(const SPIRType &type, uint32_t index) const
+TypeID CompilerMSL::get_physical_member_type_id(const SPIRType &type, uint32_t index) const
 {
 	if (member_is_remapped_physical_type(type, index))
-		return get<SPIRType>(get_extended_member_decoration(type.self, index, SPIRVCrossDecorationPhysicalTypeID));
+		return get_extended_member_decoration(type.self, index, SPIRVCrossDecorationPhysicalTypeID);
 	else
-		return get<SPIRType>(type.member_types[index]);
+		return type.member_types[index];
 }
 
 SPIRType CompilerMSL::get_presumed_input_type(const SPIRType &ib_type, uint32_t index) const
 {
-	SPIRType type = get_physical_member_type(ib_type, index);
+	SPIRType type = get<SPIRType>(get_physical_member_type_id(ib_type, index));
 	uint32_t loc = get_member_decoration(ib_type.self, index, DecorationLocation);
 	uint32_t cmp = get_member_decoration(ib_type.self, index, DecorationComponent);
 	auto p_va = inputs_by_location.find({loc, cmp});
@@ -18293,7 +18303,7 @@ SPIRType CompilerMSL::get_presumed_input_type(const SPIRType &ib_type, uint32_t
 	return type;
 }
 
-uint32_t CompilerMSL::get_declared_type_array_stride_msl(const SPIRType &type, bool is_packed, bool row_major) const
+uint32_t CompilerMSL::get_declared_type_array_stride_msl(TypeID type_id, const SPIRType *special_type, bool is_packed, bool row_major) const
 {
 	// Array stride in MSL is always size * array_size. sizeof(float3) == 16,
 	// unlike GLSL and HLSL where array stride would be 16 and size 12.
@@ -18302,11 +18312,42 @@ uint32_t CompilerMSL::get_declared_type_array_stride_msl(const SPIRType &type, b
 	// far more complicated. We'd rather just create the final type, and ignore having to create the entire type
 	// hierarchy in order to compute this value, so make a temporary type on the stack.
 
-	auto basic_type = type;
-	basic_type.array.clear();
-	basic_type.array_size_literal.clear();
-	uint32_t value_size = get_declared_type_size_msl(basic_type, is_packed, row_major);
+	uint32_t value_size;
 
+	// We don't always use proper type hierarchy for synthesized types, so be robust.
+	if (type_id && get<SPIRType>(type_id).parent_type)
+	{
+		bool uses_declared_array_stride = false;
+
+		uint32_t array_stride = 0;
+		TypeID basic_type_id = type_id;
+		while (is_array(get<SPIRType>(basic_type_id)))
+		{
+			array_stride = get_decoration(basic_type_id, DecorationArrayStride);
+			auto parent_type_id = get<SPIRType>(basic_type_id).parent_type;
+			// If the base struct itself has ArrayStride decoration, it will be padded on-demand.
+			uses_declared_array_stride = has_decoration(parent_type_id, DecorationArrayStride);
+			if (parent_type_id)
+				basic_type_id = parent_type_id;
+			else
+				break;
+		}
+
+		if (array_stride && uses_declared_array_stride)
+			value_size = array_stride;
+		else
+			value_size = get_declared_type_size_msl(basic_type_id, nullptr, is_packed, row_major);
+	}
+	else
+	{
+		// Old, broken path.
+		auto basic_type = type_id ? get<SPIRType>(type_id) : *special_type;
+		basic_type.array.clear();
+		basic_type.array_size_literal.clear();
+		value_size = get_declared_type_size_msl(0, &basic_type, is_packed, row_major);
+	}
+
+	auto &type = type_id ? get<SPIRType>(type_id) : *special_type;
 	uint32_t dimensions = uint32_t(type.array.size());
 	assert(dimensions > 0);
 	dimensions--;
@@ -18323,47 +18364,47 @@ uint32_t CompilerMSL::get_declared_type_array_stride_msl(const SPIRType &type, b
 
 uint32_t CompilerMSL::get_declared_struct_member_array_stride_msl(const SPIRType &type, uint32_t index) const
 {
-	return get_declared_type_array_stride_msl(get_physical_member_type(type, index),
+	return get_declared_type_array_stride_msl(get_physical_member_type_id(type, index), nullptr,
 	                                          member_is_packed_physical_type(type, index),
 	                                          has_member_decoration(type.self, index, DecorationRowMajor));
 }
 
 uint32_t CompilerMSL::get_declared_input_array_stride_msl(const SPIRType &type, uint32_t index) const
 {
-	return get_declared_type_array_stride_msl(get_presumed_input_type(type, index), false,
+	auto presumed_type = get_presumed_input_type(type, index);
+	return get_declared_type_array_stride_msl(0, &presumed_type, false,
 	                                          has_member_decoration(type.self, index, DecorationRowMajor));
 }
 
-uint32_t CompilerMSL::get_declared_type_matrix_stride_msl(const SPIRType &type, bool packed, bool row_major) const
+uint32_t CompilerMSL::get_declared_type_matrix_stride_msl(TypeID type_id, const SPIRType *special_type,
+                                                          bool packed, bool row_major) const
 {
+	auto &type = type_id ? get<SPIRType>(type_id) : *special_type;
+
 	// For packed matrices, we just use the size of the vector type.
 	// Otherwise, MatrixStride == alignment, which is the size of the underlying vector type.
 	if (packed)
 		return (type.width / 8) * ((row_major && type.columns > 1) ? type.columns : type.vecsize);
 	else
-		return get_declared_type_alignment_msl(type, false, row_major);
+		return get_declared_type_alignment_msl(type_id, special_type, false, row_major);
 }
 
 uint32_t CompilerMSL::get_declared_struct_member_matrix_stride_msl(const SPIRType &type, uint32_t index) const
 {
-	return get_declared_type_matrix_stride_msl(get_physical_member_type(type, index),
+	return get_declared_type_matrix_stride_msl(get_physical_member_type_id(type, index), nullptr,
 	                                           member_is_packed_physical_type(type, index),
 	                                           has_member_decoration(type.self, index, DecorationRowMajor));
 }
 
 uint32_t CompilerMSL::get_declared_input_matrix_stride_msl(const SPIRType &type, uint32_t index) const
 {
-	return get_declared_type_matrix_stride_msl(get_presumed_input_type(type, index), false,
+	auto presumed_type = get_presumed_input_type(type, index);
+	return get_declared_type_matrix_stride_msl(0, &presumed_type, false,
 	                                           has_member_decoration(type.self, index, DecorationRowMajor));
 }
 
-uint32_t CompilerMSL::get_declared_struct_size_msl(const SPIRType &struct_type, bool ignore_alignment,
-                                                   bool ignore_padding) const
+uint32_t CompilerMSL::get_declared_struct_size_msl(const SPIRType &struct_type) const
 {
-	// If we have a target size, that is the declared size as well.
-	if (!ignore_padding && has_extended_decoration(struct_type.self, SPIRVCrossDecorationPaddingTarget))
-		return get_extended_decoration(struct_type.self, SPIRVCrossDecorationPaddingTarget);
-
 	if (struct_type.member_types.empty())
 		return 0;
 
@@ -18372,13 +18413,10 @@ uint32_t CompilerMSL::get_declared_struct_size_msl(const SPIRType &struct_type,
 	// In MSL, a struct's alignment is equal to the maximum alignment of any of its members.
 	uint32_t alignment = 1;
 
-	if (!ignore_alignment)
+	for (uint32_t i = 0; i < mbr_cnt; i++)
 	{
-		for (uint32_t i = 0; i < mbr_cnt; i++)
-		{
-			uint32_t mbr_alignment = get_declared_struct_member_alignment_msl(struct_type, i);
-			alignment = max(alignment, mbr_alignment);
-		}
+		uint32_t mbr_alignment = get_declared_struct_member_alignment_msl(struct_type, i);
+		alignment = max(alignment, mbr_alignment);
 	}
 
 	// Last member will always be matched to the final Offset decoration, but size of struct in MSL now depends
@@ -18389,16 +18427,19 @@ uint32_t CompilerMSL::get_declared_struct_size_msl(const SPIRType &struct_type,
 	return msl_size;
 }
 
-uint32_t CompilerMSL::get_physical_type_stride(const SPIRType &type) const
+uint32_t CompilerMSL::get_physical_type_id_stride(TypeID type_id) const
 {
 	// This should only be relevant for plain types such as scalars and vectors?
 	// If we're pointing to a struct, it will recursively pick up packed/row-major state.
-	return get_declared_type_size_msl(type, false, false);
+	return get_declared_type_size_msl(type_id, nullptr, false, false);
 }
 
 // Returns the byte size of a struct member.
-uint32_t CompilerMSL::get_declared_type_size_msl(const SPIRType &type, bool is_packed, bool row_major) const
+uint32_t CompilerMSL::get_declared_type_size_msl(TypeID type_id, const SPIRType *special_type,
+                                                 bool is_packed, bool row_major) const
 {
+	auto &type = type_id ? get<SPIRType>(type_id) : *special_type;
+
 	// Pointers take 8 bytes each
 	// Match both pointer and array-of-pointer here.
 	if (type.pointer && type.storage == StorageClassPhysicalStorageBuffer)
@@ -18431,10 +18472,27 @@ uint32_t CompilerMSL::get_declared_type_size_msl(const SPIRType &type, bool is_p
 
 	default:
 	{
-		if (!type.array.empty())
+		if ((!type.parent_type || special_type) && !type.array.empty())
 		{
+			// Special case where the type hierarchy is not set up properly.
+			// Don't want to have to allocate a bunch of dummy type IDs just to make it work.
 			uint32_t array_size = to_array_size_literal(type);
-			return get_declared_type_array_stride_msl(type, is_packed, row_major) * max<uint32_t>(array_size, 1u);
+			return get_declared_type_array_stride_msl(type_id, special_type, is_packed, row_major) * max<uint32_t>(array_size, 1u);
+		}
+		else if (is_array(type) && type.parent_type)
+		{
+			// For the proper case. Ideally all code paths should go through here, but
+			// would need a lot of cleanup to make that work ...
+			auto &parent_type = get<SPIRType>(type.parent_type);
+			uint32_t effective_stride;
+
+			if (parent_type.op == OpTypeStruct && has_decoration(parent_type.self, DecorationArrayStride))
+				effective_stride = get_decoration(type_id, DecorationArrayStride);
+			else
+				effective_stride = get_declared_type_array_stride_msl(type_id, special_type, is_packed, row_major);
+
+			uint32_t array_size = to_array_size_literal(type);
+			return effective_stride * max<uint32_t>(array_size, 1u);
 		}
 
 		if (type.basetype == SPIRType::Struct)
@@ -18464,20 +18522,24 @@ uint32_t CompilerMSL::get_declared_type_size_msl(const SPIRType &type, bool is_p
 
 uint32_t CompilerMSL::get_declared_struct_member_size_msl(const SPIRType &type, uint32_t index) const
 {
-	return get_declared_type_size_msl(get_physical_member_type(type, index),
+	return get_declared_type_size_msl(get_physical_member_type_id(type, index), nullptr,
 	                                  member_is_packed_physical_type(type, index),
 	                                  has_member_decoration(type.self, index, DecorationRowMajor));
 }
 
 uint32_t CompilerMSL::get_declared_input_size_msl(const SPIRType &type, uint32_t index) const
 {
-	return get_declared_type_size_msl(get_presumed_input_type(type, index), false,
+	auto presumed_type = get_presumed_input_type(type, index);
+	return get_declared_type_size_msl(0, &presumed_type, false,
 	                                  has_member_decoration(type.self, index, DecorationRowMajor));
 }
 
 // Returns the byte alignment of a type.
-uint32_t CompilerMSL::get_declared_type_alignment_msl(const SPIRType &type, bool is_packed, bool row_major) const
+uint32_t CompilerMSL::get_declared_type_alignment_msl(TypeID type_id, const SPIRType *special_type,
+                                                      bool is_packed, bool row_major) const
 {
+	auto &type = type_id ? get<SPIRType>(type_id) : *special_type;
+
 	// Pointers align on multiples of 8 bytes.
 	// Deliberately ignore array-ness here. It's not relevant for alignment.
 	if (type.pointer && type.storage == StorageClassPhysicalStorageBuffer)
@@ -18531,14 +18593,15 @@ uint32_t CompilerMSL::get_declared_type_alignment_msl(const SPIRType &type, bool
 
 uint32_t CompilerMSL::get_declared_struct_member_alignment_msl(const SPIRType &type, uint32_t index) const
 {
-	return get_declared_type_alignment_msl(get_physical_member_type(type, index),
+	return get_declared_type_alignment_msl(get_physical_member_type_id(type, index), nullptr,
 	                                       member_is_packed_physical_type(type, index),
 	                                       has_member_decoration(type.self, index, DecorationRowMajor));
 }
 
 uint32_t CompilerMSL::get_declared_input_alignment_msl(const SPIRType &type, uint32_t index) const
 {
-	return get_declared_type_alignment_msl(get_presumed_input_type(type, index), false,
+	auto presumed_type = get_presumed_input_type(type, index);
+	return get_declared_type_alignment_msl(0, &presumed_type, false,
 	                                       has_member_decoration(type.self, index, DecorationRowMajor));
 }
 
diff --git a/spirv_msl.hpp b/spirv_msl.hpp
index 033cb903b..d6bd8faab 100644
--- a/spirv_msl.hpp
+++ b/spirv_msl.hpp
@@ -886,6 +886,7 @@ class CompilerMSL : public CompilerGLSL
 		SPVFuncImplVariableSizedDescriptor,
 		SPVFuncImplVariableDescriptorArray,
 		SPVFuncImplPaddedStd140,
+		SPVFuncImplPaddedArrayElement,
 		SPVFuncImplReduceAdd,
 		SPVFuncImplImageFence,
 		SPVFuncImplTextureCast,
@@ -921,7 +922,6 @@ class CompilerMSL : public CompilerGLSL
 	                             const std::string &qualifier = "");
 	void emit_struct_member(const SPIRType &type, uint32_t member_type_id, uint32_t index,
 	                        const std::string &qualifier = "", uint32_t base_offset = 0) override;
-	void emit_struct_padding_target(const SPIRType &type) override;
 	std::string type_to_glsl(const SPIRType &type, uint32_t id, bool member);
 	std::string type_to_glsl(const SPIRType &type, uint32_t id = 0) override;
 	void emit_block_hints(const SPIRBlock &block) override;
@@ -1096,15 +1096,15 @@ class CompilerMSL : public CompilerGLSL
 
 	uint32_t get_physical_tess_level_array_size(BuiltIn builtin) const;
 
-	uint32_t get_physical_type_stride(const SPIRType &type) const override;
+	uint32_t get_physical_type_id_stride(TypeID type_id) const override;
 
 	// MSL packing rules. These compute the effective packing rules as observed by the MSL compiler in the MSL output.
 	// These values can change depending on various extended decorations which control packing rules.
 	// We need to make these rules match up with SPIR-V declared rules.
-	uint32_t get_declared_type_size_msl(const SPIRType &type, bool packed, bool row_major) const;
-	uint32_t get_declared_type_array_stride_msl(const SPIRType &type, bool packed, bool row_major) const;
-	uint32_t get_declared_type_matrix_stride_msl(const SPIRType &type, bool packed, bool row_major) const;
-	uint32_t get_declared_type_alignment_msl(const SPIRType &type, bool packed, bool row_major) const;
+	uint32_t get_declared_type_size_msl(TypeID type_id, const SPIRType *special_type, bool packed, bool row_major) const;
+	uint32_t get_declared_type_array_stride_msl(TypeID type_id, const SPIRType *special_type, bool packed, bool row_major) const;
+	uint32_t get_declared_type_matrix_stride_msl(TypeID type_id, const SPIRType *special_type, bool packed, bool row_major) const;
+	uint32_t get_declared_type_alignment_msl(TypeID type_id, const SPIRType *special_type, bool packed, bool row_major) const;
 
 	uint32_t get_declared_struct_member_size_msl(const SPIRType &struct_type, uint32_t index) const;
 	uint32_t get_declared_struct_member_array_stride_msl(const SPIRType &struct_type, uint32_t index) const;
@@ -1116,11 +1116,10 @@ class CompilerMSL : public CompilerGLSL
 	uint32_t get_declared_input_matrix_stride_msl(const SPIRType &struct_type, uint32_t index) const;
 	uint32_t get_declared_input_alignment_msl(const SPIRType &struct_type, uint32_t index) const;
 
-	const SPIRType &get_physical_member_type(const SPIRType &struct_type, uint32_t index) const;
+	TypeID get_physical_member_type_id(const SPIRType &struct_type, uint32_t index) const;
 	SPIRType get_presumed_input_type(const SPIRType &struct_type, uint32_t index) const;
 
-	uint32_t get_declared_struct_size_msl(const SPIRType &struct_type, bool ignore_alignment = false,
-	                                      bool ignore_padding = false) const;
+	uint32_t get_declared_struct_size_msl(const SPIRType &struct_type) const;
 
 	std::string to_component_argument(uint32_t id);
 	void align_struct(SPIRType &ib_type, std::unordered_set<uint32_t> &aligned_structs);
diff --git a/test_shaders.py b/test_shaders.py
index 9ac4e14a7..2e019f091 100755
--- a/test_shaders.py
+++ b/test_shaders.py
@@ -438,6 +438,8 @@ def shader_model_hlsl(shader):
             return '-Tps_5_1'
     elif '.comp' in shader:
         return '-Tcs_5_1'
+    elif '.geom' in shader:
+        return '-Tgs_5_1'
     elif '.mesh' in shader:
         return '-Tms_6_5'
     elif '.task' in shader: