MSL: Support copying array of arrays.

reference/opt/shaders-msl/asm/frag/lut-promotion-initializer.asm.frag

@@ -21,14 +21,13 @@ struct main0_in
 
 // Implementation of an array copy function to cover GLSL's ability to copy an array via assignment.
 template<typename T, uint N>
-void spvArrayCopy(thread T (&dst)[N], thread const T (&src)[N])
+void spvArrayCopyFromStack1(thread T (&dst)[N], thread const T (&src)[N])
 {
     for (uint i = 0; i < N; dst[i] = src[i], i++);
 }
 
-// An overload for constant arrays.
 template<typename T, uint N>
-void spvArrayCopyConstant(thread T (&dst)[N], constant T (&src)[N])
+void spvArrayCopyFromConstant1(thread T (&dst)[N], constant T (&src)[N])
 {
     for (uint i = 0; i < N; dst[i] = src[i], i++);
 }
@@ -62,7 +61,7 @@ fragment main0_out main0(main0_in in [[stage_in]])
     }
     int _37 = in.index & 3;
     out.FragColor += foobar[_37].z;
-    spvArrayCopyConstant(baz, _90);
+    spvArrayCopyFromConstant1(baz, _90);
     out.FragColor += baz[_37].z;
     return out;
 }

reference/opt/shaders-msl/comp/composite-array-initialization.comp

@@ -30,14 +30,13 @@ constant Data _25[2] = { Data{ 1.0, 2.0 }, Data{ 3.0, 4.0 } };
 
 // Implementation of an array copy function to cover GLSL's ability to copy an array via assignment.
 template<typename T, uint N>
-void spvArrayCopy(thread T (&dst)[N], thread const T (&src)[N])
+void spvArrayCopyFromStack1(thread T (&dst)[N], thread const T (&src)[N])
 {
     for (uint i = 0; i < N; dst[i] = src[i], i++);
 }
 
-// An overload for constant arrays.
 template<typename T, uint N>
-void spvArrayCopyConstant(thread T (&dst)[N], constant T (&src)[N])
+void spvArrayCopyFromConstant1(thread T (&dst)[N], constant T (&src)[N])
 {
     for (uint i = 0; i < N; dst[i] = src[i], i++);
 }
@@ -47,7 +46,7 @@ kernel void main0(device SSBO& _53 [[buffer(0)]], uint3 gl_WorkGroupID [[threadg
     Data data[2] = { Data{ 1.0, 2.0 }, Data{ 3.0, 4.0 } };
     Data _31[2] = { Data{ X, 2.0 }, Data{ 3.0, 5.0 } };
     Data data2[2];
-    spvArrayCopy(data2, _31);
+    spvArrayCopyFromStack1(data2, _31);
     _53.outdata[gl_WorkGroupID.x].a = data[gl_LocalInvocationID.x].a + data2[gl_LocalInvocationID.x].a;
     _53.outdata[gl_WorkGroupID.x].b = data[gl_LocalInvocationID.x].b + data2[gl_LocalInvocationID.x].b;
 }

reference/opt/shaders-msl/comp/composite-construct.comp

@@ -12,14 +12,13 @@ struct SSBO0
 
 // Implementation of an array copy function to cover GLSL's ability to copy an array via assignment.
 template<typename T, uint N>
-void spvArrayCopy(thread T (&dst)[N], thread const T (&src)[N])
+void spvArrayCopyFromStack1(thread T (&dst)[N], thread const T (&src)[N])
 {
     for (uint i = 0; i < N; dst[i] = src[i], i++);
 }
 
-// An overload for constant arrays.
 template<typename T, uint N>
-void spvArrayCopyConstant(thread T (&dst)[N], constant T (&src)[N])
+void spvArrayCopyFromConstant1(thread T (&dst)[N], constant T (&src)[N])
 {
     for (uint i = 0; i < N; dst[i] = src[i], i++);
 }
@@ -28,7 +27,7 @@ kernel void main0(device SSBO0& _16 [[buffer(0)]], device SSBO0& _32 [[buffer(1)
 {
     float4 _37[2] = { _16.as[gl_GlobalInvocationID.x], _32.as[gl_GlobalInvocationID.x] };
     float4 values[2];
-    spvArrayCopy(values, _37);
+    spvArrayCopyFromStack1(values, _37);
     _16.as[0] = values[gl_LocalInvocationIndex];
     _32.as[1] = float4(40.0);
 }

reference/opt/shaders-msl/comp/copy-array-of-arrays.comp

@@ -0,0 +1,15 @@
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+struct BUF
+{
+    int a;
+};
+
+kernel void main0(device BUF& o [[buffer(0)]])
+{
+    o.a = 4;
+}
+

reference/opt/shaders-msl/frag/constant-array.frag

@@ -29,14 +29,13 @@ struct main0_in
 
 // Implementation of an array copy function to cover GLSL's ability to copy an array via assignment.
 template<typename T, uint N>
-void spvArrayCopy(thread T (&dst)[N], thread const T (&src)[N])
+void spvArrayCopyFromStack1(thread T (&dst)[N], thread const T (&src)[N])
 {
     for (uint i = 0; i < N; dst[i] = src[i], i++);
 }
 
-// An overload for constant arrays.
 template<typename T, uint N>
-void spvArrayCopyConstant(thread T (&dst)[N], constant T (&src)[N])
+void spvArrayCopyFromConstant1(thread T (&dst)[N], constant T (&src)[N])
 {
     for (uint i = 0; i < N; dst[i] = src[i], i++);
 }

reference/opt/shaders-msl/frag/constant-composites.frag

@@ -26,14 +26,13 @@ struct main0_in
 
 // Implementation of an array copy function to cover GLSL's ability to copy an array via assignment.
 template<typename T, uint N>
-void spvArrayCopy(thread T (&dst)[N], thread const T (&src)[N])
+void spvArrayCopyFromStack1(thread T (&dst)[N], thread const T (&src)[N])
 {
     for (uint i = 0; i < N; dst[i] = src[i], i++);
 }
 
-// An overload for constant arrays.
 template<typename T, uint N>
-void spvArrayCopyConstant(thread T (&dst)[N], constant T (&src)[N])
+void spvArrayCopyFromConstant1(thread T (&dst)[N], constant T (&src)[N])
 {
     for (uint i = 0; i < N; dst[i] = src[i], i++);
 }

reference/opt/shaders-msl/frag/lut-promotion.frag

@@ -21,14 +21,13 @@ struct main0_in
 
 // Implementation of an array copy function to cover GLSL's ability to copy an array via assignment.
 template<typename T, uint N>
-void spvArrayCopy(thread T (&dst)[N], thread const T (&src)[N])
+void spvArrayCopyFromStack1(thread T (&dst)[N], thread const T (&src)[N])
 {
     for (uint i = 0; i < N; dst[i] = src[i], i++);
 }
 
-// An overload for constant arrays.
 template<typename T, uint N>
-void spvArrayCopyConstant(thread T (&dst)[N], constant T (&src)[N])
+void spvArrayCopyFromConstant1(thread T (&dst)[N], constant T (&src)[N])
 {
     for (uint i = 0; i < N; dst[i] = src[i], i++);
 }
@@ -62,7 +61,7 @@ fragment main0_out main0(main0_in in [[stage_in]])
     int _91 = in.index & 3;
     out.FragColor += foobar[_91].z;
     float4 baz[4] = { float4(0.0), float4(1.0), float4(8.0), float4(5.0) };
-    spvArrayCopyConstant(baz, _104);
+    spvArrayCopyFromConstant1(baz, _104);
     out.FragColor += baz[_91].z;
     return out;
 }

reference/shaders-msl/asm/frag/lut-promotion-initializer.asm.frag

@@ -21,14 +21,13 @@ struct main0_in
 
 // Implementation of an array copy function to cover GLSL's ability to copy an array via assignment.
 template<typename T, uint N>
-void spvArrayCopy(thread T (&dst)[N], thread const T (&src)[N])
+void spvArrayCopyFromStack1(thread T (&dst)[N], thread const T (&src)[N])
 {
     for (uint i = 0; i < N; dst[i] = src[i], i++);
 }
 
-// An overload for constant arrays.
 template<typename T, uint N>
-void spvArrayCopyConstant(thread T (&dst)[N], constant T (&src)[N])
+void spvArrayCopyFromConstant1(thread T (&dst)[N], constant T (&src)[N])
 {
     for (uint i = 0; i < N; dst[i] = src[i], i++);
 }
@@ -60,7 +59,7 @@ fragment main0_out main0(main0_in in [[stage_in]])
         foobar[1].z = 20.0;
     }
     out.FragColor += foobar[in.index & 3].z;
-    spvArrayCopyConstant(baz, _90);
+    spvArrayCopyFromConstant1(baz, _90);
     out.FragColor += baz[in.index & 3].z;
     return out;
 }

reference/shaders-msl/comp/composite-array-initialization.comp

@@ -30,14 +30,13 @@ constant Data _25[2] = { Data{ 1.0, 2.0 }, Data{ 3.0, 4.0 } };
 
 // Implementation of an array copy function to cover GLSL's ability to copy an array via assignment.
 template<typename T, uint N>
-void spvArrayCopy(thread T (&dst)[N], thread const T (&src)[N])
+void spvArrayCopyFromStack1(thread T (&dst)[N], thread const T (&src)[N])
 {
     for (uint i = 0; i < N; dst[i] = src[i], i++);
 }
 
-// An overload for constant arrays.
 template<typename T, uint N>
-void spvArrayCopyConstant(thread T (&dst)[N], constant T (&src)[N])
+void spvArrayCopyFromConstant1(thread T (&dst)[N], constant T (&src)[N])
 {
     for (uint i = 0; i < N; dst[i] = src[i], i++);
 }
@@ -52,7 +51,7 @@ kernel void main0(device SSBO& _53 [[buffer(0)]], uint3 gl_WorkGroupID [[threadg
     Data data[2] = { Data{ 1.0, 2.0 }, Data{ 3.0, 4.0 } };
     Data _31[2] = { Data{ X, 2.0 }, Data{ 3.0, 5.0 } };
     Data data2[2];
-    spvArrayCopy(data2, _31);
+    spvArrayCopyFromStack1(data2, _31);
     Data param = data[gl_LocalInvocationID.x];
     Data param_1 = data2[gl_LocalInvocationID.x];
     Data _73 = combine(param, param_1);

reference/shaders-msl/comp/composite-construct.comp

@@ -25,14 +25,13 @@ constant float4 _43[2] = { float4(20.0), float4(40.0) };
 
 // Implementation of an array copy function to cover GLSL's ability to copy an array via assignment.
 template<typename T, uint N>
-void spvArrayCopy(thread T (&dst)[N], thread const T (&src)[N])
+void spvArrayCopyFromStack1(thread T (&dst)[N], thread const T (&src)[N])
 {
     for (uint i = 0; i < N; dst[i] = src[i], i++);
 }
 
-// An overload for constant arrays.
 template<typename T, uint N>
-void spvArrayCopyConstant(thread T (&dst)[N], constant T (&src)[N])
+void spvArrayCopyFromConstant1(thread T (&dst)[N], constant T (&src)[N])
 {
     for (uint i = 0; i < N; dst[i] = src[i], i++);
 }
@@ -41,7 +40,7 @@ kernel void main0(device SSBO0& _16 [[buffer(0)]], device SSBO1& _32 [[buffer(1)
 {
     float4 _37[2] = { _16.as[gl_GlobalInvocationID.x], _32.bs[gl_GlobalInvocationID.x] };
     float4 values[2];
-    spvArrayCopy(values, _37);
+    spvArrayCopyFromStack1(values, _37);
     Composite c = Composite{ values[0], _43[1] };
     _16.as[0] = values[gl_LocalInvocationIndex];
     _32.bs[1] = c.b;

reference/shaders-msl/comp/copy-array-of-arrays.comp

@@ -0,0 +1,73 @@
+#pragma clang diagnostic ignored "-Wmissing-prototypes"
+
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+struct BUF
+{
+    int a;
+};
+
+constant float _16[2] = { 1.0, 2.0 };
+constant float _19[2] = { 3.0, 4.0 };
+constant float _20[2][2] = { { 1.0, 2.0 }, { 3.0, 4.0 } };
+constant float _21[2][2][2] = { { { 1.0, 2.0 }, { 3.0, 4.0 } }, { { 1.0, 2.0 }, { 3.0, 4.0 } } };
+
+// Implementation of an array copy function to cover GLSL's ability to copy an array via assignment.
+template<typename T, uint N>
+void spvArrayCopyFromStack1(thread T (&dst)[N], thread const T (&src)[N])
+{
+    for (uint i = 0; i < N; dst[i] = src[i], i++);
+}
+
+template<typename T, uint N>
+void spvArrayCopyFromConstant1(thread T (&dst)[N], constant T (&src)[N])
+{
+    for (uint i = 0; i < N; dst[i] = src[i], i++);
+}
+
+template<typename T, uint A, uint B>
+void spvArrayCopyFromStack2(thread T (&dst)[A][B], thread const T (&src)[A][B])
+{
+    for (uint i = 0; i < A; i++)
+    {
+        spvArrayCopyFromStack1(dst[i], src[i]);
+    }
+}
+
+template<typename T, uint A, uint B>
+void spvArrayCopyFromConstant2(thread T (&dst)[A][B], constant T (&src)[A][B])
+{
+    for (uint i = 0; i < A; i++)
+    {
+        spvArrayCopyFromConstant1(dst[i], src[i]);
+    }
+}
+
+template<typename T, uint A, uint B, uint C>
+void spvArrayCopyFromStack3(thread T (&dst)[A][B][C], thread const T (&src)[A][B][C])
+{
+    for (uint i = 0; i < A; i++)
+    {
+        spvArrayCopyFromStack2(dst[i], src[i]);
+    }
+}
+
+template<typename T, uint A, uint B, uint C>
+void spvArrayCopyFromConstant3(thread T (&dst)[A][B][C], constant T (&src)[A][B][C])
+{
+    for (uint i = 0; i < A; i++)
+    {
+        spvArrayCopyFromConstant2(dst[i], src[i]);
+    }
+}
+
+kernel void main0(device BUF& o [[buffer(0)]])
+{
+    float c[2][2][2];
+    spvArrayCopyFromConstant3(c, _21);
+    o.a = int(c[1][1][1]);
+}
+

reference/shaders-msl/frag/constant-array.frag

@@ -29,14 +29,13 @@ struct main0_in
 
 // Implementation of an array copy function to cover GLSL's ability to copy an array via assignment.
 template<typename T, uint N>
-void spvArrayCopy(thread T (&dst)[N], thread const T (&src)[N])
+void spvArrayCopyFromStack1(thread T (&dst)[N], thread const T (&src)[N])
 {
     for (uint i = 0; i < N; dst[i] = src[i], i++);
 }
 
-// An overload for constant arrays.
 template<typename T, uint N>
-void spvArrayCopyConstant(thread T (&dst)[N], constant T (&src)[N])
+void spvArrayCopyFromConstant1(thread T (&dst)[N], constant T (&src)[N])
 {
     for (uint i = 0; i < N; dst[i] = src[i], i++);
 }

reference/shaders-msl/frag/constant-composites.frag

@@ -26,14 +26,13 @@ struct main0_in
 
 // Implementation of an array copy function to cover GLSL's ability to copy an array via assignment.
 template<typename T, uint N>
-void spvArrayCopy(thread T (&dst)[N], thread const T (&src)[N])
+void spvArrayCopyFromStack1(thread T (&dst)[N], thread const T (&src)[N])
 {
     for (uint i = 0; i < N; dst[i] = src[i], i++);
 }
 
-// An overload for constant arrays.
 template<typename T, uint N>
-void spvArrayCopyConstant(thread T (&dst)[N], constant T (&src)[N])
+void spvArrayCopyFromConstant1(thread T (&dst)[N], constant T (&src)[N])
 {
     for (uint i = 0; i < N; dst[i] = src[i], i++);
 }

reference/shaders-msl/frag/lut-promotion.frag

@@ -21,14 +21,13 @@ struct main0_in
 
 // Implementation of an array copy function to cover GLSL's ability to copy an array via assignment.
 template<typename T, uint N>
-void spvArrayCopy(thread T (&dst)[N], thread const T (&src)[N])
+void spvArrayCopyFromStack1(thread T (&dst)[N], thread const T (&src)[N])
 {
     for (uint i = 0; i < N; dst[i] = src[i], i++);
 }
 
-// An overload for constant arrays.
 template<typename T, uint N>
-void spvArrayCopyConstant(thread T (&dst)[N], constant T (&src)[N])
+void spvArrayCopyFromConstant1(thread T (&dst)[N], constant T (&src)[N])
 {
     for (uint i = 0; i < N; dst[i] = src[i], i++);
 }
@@ -60,7 +59,7 @@ fragment main0_out main0(main0_in in [[stage_in]])
     }
     out.FragColor += foobar[in.index & 3].z;
     float4 baz[4] = { float4(0.0), float4(1.0), float4(8.0), float4(5.0) };
-    spvArrayCopyConstant(baz, _104);
+    spvArrayCopyFromConstant1(baz, _104);
     out.FragColor += baz[in.index & 3].z;
     return out;
 }

reference/shaders-msl/vert/return-array.vert

@@ -20,21 +20,20 @@ struct main0_in
 
 // Implementation of an array copy function to cover GLSL's ability to copy an array via assignment.
 template<typename T, uint N>
-void spvArrayCopy(thread T (&dst)[N], thread const T (&src)[N])
+void spvArrayCopyFromStack1(thread T (&dst)[N], thread const T (&src)[N])
 {
     for (uint i = 0; i < N; dst[i] = src[i], i++);
 }
 
-// An overload for constant arrays.
 template<typename T, uint N>
-void spvArrayCopyConstant(thread T (&dst)[N], constant T (&src)[N])
+void spvArrayCopyFromConstant1(thread T (&dst)[N], constant T (&src)[N])
 {
     for (uint i = 0; i < N; dst[i] = src[i], i++);
 }
 
 void test(thread float4 (&SPIRV_Cross_return_value)[2])
 {
-    spvArrayCopyConstant(SPIRV_Cross_return_value, _20);
+    spvArrayCopyFromConstant1(SPIRV_Cross_return_value, _20);
 }
 
 void test2(thread float4 (&SPIRV_Cross_return_value)[2], thread float4& vInput0, thread float4& vInput1)
@@ -42,7 +41,7 @@ void test2(thread float4 (&SPIRV_Cross_return_value)[2], thread float4& vInput0,
     float4 foobar[2];
     foobar[0] = vInput0;
     foobar[1] = vInput1;
-    spvArrayCopy(SPIRV_Cross_return_value, foobar);
+    spvArrayCopyFromStack1(SPIRV_Cross_return_value, foobar);
 }
 
 vertex main0_out main0(main0_in in [[stage_in]])

shaders-msl/comp/copy-array-of-arrays.comp

@@ -0,0 +1,15 @@
+#version 450
+layout(local_size_x = 1) in;
+
+layout(set = 0, binding = 0, std430) buffer BUF
+{
+	int a;
+} o;
+
+void main()
+{
+	const float a[2][2][2] = float[][][](float[][](float[](1.0, 2.0), float[](3.0, 4.0)), float[][](float[](1.0, 2.0), float[](3.0, 4.0)));
+	float b[2][2][2] = a;
+	float c[2][2][2] = b;
+	o.a = int(c[1][1][1]);
+}

spirv_msl.cpp

@@ -1168,6 +1168,10 @@ void CompilerMSL::add_typedef_line(const string &line)
 // Emits any needed custom function bodies.
 void CompilerMSL::emit_custom_functions()
 {
+	for (uint32_t i = SPVFuncImplArrayCopyMultidimMax; i >= 2; i--)
+		if (spv_function_implementations.count(static_cast<SPVFuncImpl>(SPVFuncImplArrayCopyMultidimBase + i)))
+			spv_function_implementations.insert(static_cast<SPVFuncImpl>(SPVFuncImplArrayCopyMultidimBase + i - 1));
+
 	for (auto &spv_func : spv_function_implementations)
 	{
 		switch (spv_func)
@@ -1237,21 +1241,72 @@ void CompilerMSL::emit_custom_functions()
 			statement("// Implementation of an array copy function to cover GLSL's ability to copy an array via "
 			          "assignment.");
 			statement("template<typename T, uint N>");
-			statement("void spvArrayCopy(thread T (&dst)[N], thread const T (&src)[N])");
+			statement("void spvArrayCopyFromStack1(thread T (&dst)[N], thread const T (&src)[N])");
 			begin_scope();
 			statement("for (uint i = 0; i < N; dst[i] = src[i], i++);");
 			end_scope();
 			statement("");
 
-			statement("// An overload for constant arrays.");
 			statement("template<typename T, uint N>");
-			statement("void spvArrayCopyConstant(thread T (&dst)[N], constant T (&src)[N])");
+			statement("void spvArrayCopyFromConstant1(thread T (&dst)[N], constant T (&src)[N])");
 			begin_scope();
 			statement("for (uint i = 0; i < N; dst[i] = src[i], i++);");
 			end_scope();
 			statement("");
 			break;
 
+		case SPVFuncImplArrayOfArrayCopy2Dim:
+		case SPVFuncImplArrayOfArrayCopy3Dim:
+		case SPVFuncImplArrayOfArrayCopy4Dim:
+		case SPVFuncImplArrayOfArrayCopy5Dim:
+		case SPVFuncImplArrayOfArrayCopy6Dim:
+		{
+			static const char *function_name_tags[] = {
+				"FromStack",
+				"FromConstant",
+			};
+
+			static const char *src_address_space[] = {
+				"thread const",
+				"constant",
+			};
+
+			for (uint32_t variant = 0; variant < 2; variant++)
+			{
+				uint32_t dimensions = spv_func - SPVFuncImplArrayCopyMultidimBase;
+				string tmp = "template<typename T";
+				for (uint32_t i = 0; i < dimensions; i++)
+				{
+					tmp += ", uint ";
+					tmp += 'A' + i;
+				}
+				tmp += ">";
+				statement(tmp);
+
+				string array_arg;
+				for (uint32_t i = 0; i < dimensions; i++)
+				{
+					array_arg += "[";
+					array_arg += 'A' + i;
+					array_arg += "]";
+				}
+
+				statement("void spvArrayCopy", function_name_tags[variant], dimensions,
+				          "(thread T (&dst)", array_arg,
+				          ", ", src_address_space[variant], " T (&src)", array_arg,
+				          ")");
+
+				begin_scope();
+				statement("for (uint i = 0; i < A; i++)");
+				begin_scope();
+				statement("spvArrayCopy", function_name_tags[variant], dimensions - 1, "(dst[i], src[i]);");
+				end_scope();
+				end_scope();
+				statement("");
+			}
+			break;
+		}
+
 		case SPVFuncImplTexelBufferCoords:
 		{
 			string tex_width_str = convert_to_string(msl_options.texel_buffer_texture_width);
@@ -2212,10 +2267,24 @@ bool CompilerMSL::maybe_emit_input_struct_assignment(uint32_t id_lhs, uint32_t i
 void CompilerMSL::emit_array_copy(const string &lhs, uint32_t rhs_id)
 {
 	// Assignment from an array initializer is fine.
+	auto &type = expression_type(rhs_id);
+	auto *var = maybe_get_backing_variable(rhs_id);
+
+	// Unfortunately, we cannot template on address space in MSL,
+	// so explicit address space redirection it is ...
+	bool is_constant = false;
 	if (ids[rhs_id].get_type() == TypeConstant)
-		statement("spvArrayCopyConstant(", lhs, ", ", to_expression(rhs_id), ");");
-	else
-		statement("spvArrayCopy(", lhs, ", ", to_expression(rhs_id), ");");
+	{
+		is_constant = true;
+	}
+	else if (var && var->remapped_variable && var->statically_assigned &&
+	         ids[var->static_expression].get_type() == TypeConstant)
+	{
+		is_constant = true;
+	}
+
+	const char *tag = is_constant ? "FromConstant" : "FromStack";
+	statement("spvArrayCopy", tag, type.array.size(), "(", lhs, ", ", to_expression(rhs_id), ");");
 }
 
 // Since MSL does not allow arrays to be copied via simple variable assignment,
@@ -4380,7 +4449,13 @@ CompilerMSL::SPVFuncImpl CompilerMSL::OpCodePreprocessor::get_spv_func_impl(Op o
 	case OpFunctionCall:
 	{
 		auto &return_type = compiler.get<SPIRType>(args[0]);
-		if (!return_type.array.empty())
+		if (return_type.array.size() > 1)
+		{
+			if (return_type.array.size() > SPVFuncImplArrayCopyMultidimMax)
+				SPIRV_CROSS_THROW("Cannot support this many dimensions for arrays of arrays.");
+			return static_cast<SPVFuncImpl>(SPVFuncImplArrayCopyMultidimBase + return_type.array.size());
+		}
+		else if (return_type.array.size() > 0)
 			return SPVFuncImplArrayCopy;
 
 		break;
@@ -4414,7 +4489,16 @@ CompilerMSL::SPVFuncImpl CompilerMSL::OpCodePreprocessor::get_spv_func_impl(Op o
 		bool static_expression_lhs =
 		    var && var->storage == StorageClassFunction && var->statically_assigned && var->remapped_variable;
 		if (type && compiler.is_array(*type) && !static_expression_lhs)
+		{
+			if (type->array.size() > 1)
+			{
+				if (type->array.size() > SPVFuncImplArrayCopyMultidimMax)
+					SPIRV_CROSS_THROW("Cannot support this many dimensions for arrays of arrays.");
+				return static_cast<SPVFuncImpl>(SPVFuncImplArrayCopyMultidimBase + type->array.size());
+			}
+			else
 				return SPVFuncImplArrayCopy;
+		}
 
 		break;
 	}

spirv_msl.hpp

@@ -222,7 +222,15 @@ public:
 		SPVFuncImplFindILsb,
 		SPVFuncImplFindSMsb,
 		SPVFuncImplFindUMsb,
-		SPVFuncImplArrayCopy,
+		SPVFuncImplArrayCopyMultidimBase,
+		// Unfortunately, we cannot use recursive templates in the MSL compiler properly,
+		// so stamp out variants up to some arbitrary maximum.
+		SPVFuncImplArrayCopy = SPVFuncImplArrayCopyMultidimBase + 1,
+		SPVFuncImplArrayOfArrayCopy2Dim = SPVFuncImplArrayCopyMultidimBase + 2,
+		SPVFuncImplArrayOfArrayCopy3Dim = SPVFuncImplArrayCopyMultidimBase + 3,
+		SPVFuncImplArrayOfArrayCopy4Dim = SPVFuncImplArrayCopyMultidimBase + 4,
+		SPVFuncImplArrayOfArrayCopy5Dim = SPVFuncImplArrayCopyMultidimBase + 5,
+		SPVFuncImplArrayOfArrayCopy6Dim = SPVFuncImplArrayCopyMultidimBase + 6,
 		SPVFuncImplTexelBufferCoords,
 		SPVFuncImplInverse4x4,
 		SPVFuncImplInverse3x3,
@@ -233,6 +241,7 @@ public:
 		SPVFuncImplRowMajor3x4,
 		SPVFuncImplRowMajor4x2,
 		SPVFuncImplRowMajor4x3,
+		SPVFuncImplArrayCopyMultidimMax = 6
 	};
 
 	// Constructs an instance to compile the SPIR-V code into Metal Shading Language,