diff --git a/reference/shaders-msl-no-opt/packing/load-store-col-rows.comp b/reference/shaders-msl-no-opt/packing/load-store-col-rows.comp
new file mode 100644
index 00000000..c7024354
--- /dev/null
+++ b/reference/shaders-msl-no-opt/packing/load-store-col-rows.comp
@@ -0,0 +1,76 @@
+#pragma clang diagnostic ignored "-Wmissing-prototypes"
+
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+typedef packed_float3 packed_float2x3[2];
+typedef packed_float3 packed_rm_float3x2[2];
+
+struct SSBO1
+{
+    float2x4 a;
+    float2x4 a2;
+};
+
+struct SSBO2
+{
+    packed_float2x3 b;
+    packed_rm_float3x2 b2;
+};
+
+constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u);
+
+static inline __attribute__((always_inline))
+void load_store_column(device SSBO1& v_21)
+{
+    float2 u = v_21.a[0].xy;
+    float2 v = v_21.a[1].xy;
+    u += v;
+    v_21.a[0].xy = u;
+    v_21.a[1].xy = v;
+}
+
+static inline __attribute__((always_inline))
+void load_store_row(device SSBO1& v_21)
+{
+    float2 u = float2(v_21.a2[0][0], v_21.a2[1][0]);
+    float2 v = float2(v_21.a2[0][1], v_21.a2[1][1]);
+    u += v;
+    v_21.a2[0][0] = u.x;
+    v_21.a2[1][0] = u.y;
+    v_21.a2[0][1] = v.x;
+    v_21.a2[1][1] = v.y;
+}
+
+static inline __attribute__((always_inline))
+void load_store_packed_column(device SSBO2& v_58)
+{
+    float3 u = float3(v_58.b[0]);
+    float3 v = float3(v_58.b[1]);
+    u += v;
+    v_58.b[0] = u;
+    v_58.b[1] = v;
+}
+
+static inline __attribute__((always_inline))
+void load_store_packed_row(device SSBO2& v_58)
+{
+    float2 u = float2(v_58.b2[0][0], v_58.b2[1][0]);
+    float2 v = float2(v_58.b2[0][1], v_58.b2[1][1]);
+    u += v;
+    v_58.b2[0][0] = u.x;
+    v_58.b2[1][0] = u.y;
+    v_58.b2[0][1] = v.x;
+    v_58.b2[1][1] = v.y;
+}
+
+kernel void main0(device SSBO1& v_21 [[buffer(0)]], device SSBO2& v_58 [[buffer(1)]])
+{
+    load_store_column(v_21);
+    load_store_row(v_21);
+    load_store_packed_column(v_58);
+    load_store_packed_row(v_58);
+}
+
diff --git a/shaders-msl-no-opt/packing/load-store-col-rows.comp b/shaders-msl-no-opt/packing/load-store-col-rows.comp
new file mode 100644
index 00000000..b3f28970
--- /dev/null
+++ b/shaders-msl-no-opt/packing/load-store-col-rows.comp
@@ -0,0 +1,59 @@
+#version 450
+#extension GL_EXT_scalar_block_layout : require
+layout(local_size_x = 1) in;
+
+layout(binding = 0, std140) buffer SSBO1
+{
+	mat2 a;
+	layout(row_major) mat2 a2;
+};
+
+layout(scalar, binding = 1) buffer SSBO2
+{
+	mat2x3 b;
+	layout(row_major) mat3x2 b2;
+};
+
+void load_store_column()
+{
+	vec2 u = a[0];
+	vec2 v = a[1];
+	u += v;
+	a[0] = u;
+	a[1] = v;
+}
+
+void load_store_row()
+{
+	vec2 u = a2[0];
+	vec2 v = a2[1];
+	u += v;
+	a2[0] = u;
+	a2[1] = v;
+}
+
+void load_store_packed_column()
+{
+	vec3 u = b[0];
+	vec3 v = b[1];
+	u += v;
+	b[0] = u;
+	b[1] = v;
+}
+
+void load_store_packed_row()
+{
+	vec2 u = b2[0];
+	vec2 v = b2[1];
+	u += v;
+	b2[0] = u;
+	b2[1] = v;
+}
+
+void main()
+{
+	load_store_column();
+	load_store_row();
+	load_store_packed_column();
+	load_store_packed_row();
+}
diff --git a/spirv_msl.cpp b/spirv_msl.cpp
index bc79e55b..44a78748 100644
--- a/spirv_msl.cpp
+++ b/spirv_msl.cpp
@@ -3193,10 +3193,17 @@ string CompilerMSL::unpack_expression_type(string expr_str, const SPIRType &type
 		".xyz",
 	};
 
-	// std140 array cases for vectors.
 	if (physical_type && is_vector(*physical_type) && is_array(*physical_type) &&
 	    physical_type->vecsize > type.vecsize && !expression_ends_with(expr_str, swizzle_lut[type.vecsize - 1]))
 	{
+		// std140 array cases for vectors.
+		assert(type.vecsize >= 1 && type.vecsize <= 3);
+		return enclose_expression(expr_str) + swizzle_lut[type.vecsize - 1];
+	}
+	else if (physical_type && is_matrix(*physical_type) && is_vector(type) &&
+	         physical_type->vecsize > type.vecsize)
+	{
+		// Extract column from padded matrix.
 		assert(type.vecsize >= 1 && type.vecsize <= 3);
 		return enclose_expression(expr_str) + swizzle_lut[type.vecsize - 1];
 	}