MSL: Support array-of-arrays composite construction.

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

@@ -6,10 +6,13 @@ using namespace metal;
 struct BUF
 {
     int a;
+    float b;
+    float c;
 };
 
 kernel void main0(device BUF& o [[buffer(0)]])
 {
     o.a = 4;
+    o.b = o.c;
 }
 

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

@@ -8,6 +8,8 @@ using namespace metal;
 struct BUF
 {
     int a;
+    float b;
+    float c;
 };
 
 constant float _16[2] = { 1.0, 2.0 };
@@ -69,5 +71,23 @@ kernel void main0(device BUF& o [[buffer(0)]])
     float c[2][2][2];
     spvArrayCopyFromConstant3(c, _21);
     o.a = int(c[1][1][1]);
+    float _43[2] = { o.b, o.c };
+    float _48[2] = { o.b, o.b };
+    float _49[2][2];
+    spvArrayCopyFromStack1(_49[0], _43);
+    spvArrayCopyFromStack1(_49[1], _48);
+    float _54[2] = { o.c, o.c };
+    float _59[2] = { o.c, o.b };
+    float _60[2][2];
+    spvArrayCopyFromStack1(_60[0], _54);
+    spvArrayCopyFromStack1(_60[1], _59);
+    float _61[2][2][2];
+    spvArrayCopyFromStack2(_61[0], _49);
+    spvArrayCopyFromStack2(_61[1], _60);
+    float d[2][2][2];
+    spvArrayCopyFromStack3(d, _61);
+    float e[2][2][2];
+    spvArrayCopyFromStack3(e, d);
+    o.b = e[1][0][1];
 }
 

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

@@ -4,6 +4,8 @@ layout(local_size_x = 1) in;
 layout(set = 0, binding = 0, std430) buffer BUF
 {
 	int a;
+	float b;
+	float c;
 } o;
 
 void main()
@@ -12,4 +14,8 @@ void main()
 	float b[2][2][2] = a;
 	float c[2][2][2] = b;
 	o.a = int(c[1][1][1]);
+
+	float d[2][2][2] = float[][][](float[][](float[](o.b, o.c), float[](o.b, o.b)), float[][](float[](o.c, o.c), float[](o.c, o.b)));
+	float e[2][2][2] = d;
+	o.b = e[1][0][1];
 }

spirv_glsl.cpp

@@ -6515,7 +6515,19 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction)
 			forward = false;
 
 		string constructor_op;
-		if (backend.use_initializer_list && composite)
+		if (!backend.array_is_value_type && out_type.array.size() > 1)
+		{
+			// We cannot construct array of arrays because we cannot treat the inputs
+			// as value types. Need to declare the array-of-arrays, and copy in elements one by one.
+			forced_temporaries.insert(id);
+			auto flags = meta[id].decoration.decoration_flags;
+			statement(flags_to_precision_qualifiers_glsl(out_type, flags),
+			          variable_decl(out_type, to_name(id)), ";");
+			set<SPIRExpression>(id, to_name(id), result_type, true);
+			for (uint32_t i = 0; i < length; i++)
+				emit_array_copy(join(to_expression(id), "[", i, "]"), elems[i]);
+		}
+		else if (backend.use_initializer_list && composite)
 		{
 			// Only use this path if we are building composites.
 			// This path cannot be used for arithmetic.
@@ -6546,9 +6558,12 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction)
 			constructor_op += ")";
 		}
 
-		emit_op(result_type, id, constructor_op, forward);
-		for (uint32_t i = 0; i < length; i++)
-			inherit_expression_dependencies(id, elems[i]);
+		if (!constructor_op.empty())
+		{
+			emit_op(result_type, id, constructor_op, forward);
+			for (uint32_t i = 0; i < length; i++)
+				inherit_expression_dependencies(id, elems[i]);
+		}
 		break;
 	}
 

spirv_glsl.hpp

@@ -363,6 +363,7 @@ protected:
 		bool allow_truncated_access_chain = false;
 		bool supports_extensions = false;
 		bool supports_empty_struct = false;
+		bool array_is_value_type = true;
 	} backend;
 
 	void emit_struct(SPIRType &type);

spirv_msl.cpp

@@ -276,6 +276,7 @@ string CompilerMSL::compile()
 	backend.can_return_array = false;
 	backend.boolean_mix_support = false;
 	backend.allow_truncated_access_chain = true;
+	backend.array_is_value_type = false;
 
 	is_rasterization_disabled = msl_options.disable_rasterization;
 
@@ -4514,6 +4515,14 @@ CompilerMSL::SPVFuncImpl CompilerMSL::OpCodePreprocessor::get_spv_func_impl(Op o
 		break;
 	}
 
+	case OpCompositeConstruct:
+	{
+		auto &type = compiler.get<SPIRType>(args[0]);
+		if (type.array.size() > 1) // We need to use copies to build the composite.
+			return static_cast<SPVFuncImpl>(SPVFuncImplArrayCopyMultidimBase + type.array.size() - 1);
+		break;
+	}
+
 	case OpExtInst:
 	{
 		uint32_t extension_set = args[2];