Merge pull request #2328 from KhronosGroup/dynamic-buffer-cleanup

Dynamic buffer cleanup

reference/shaders-msl-no-opt/comp/basic.dynamic-buffer.msl2.comp

@@ -0,0 +1,66 @@
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+struct Baz
+{
+    int e;
+    int f;
+};
+
+struct Foo
+{
+    int a;
+    int b;
+};
+
+struct Bar
+{
+    int c;
+    int d;
+};
+
+constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(3u, 3u, 2u);
+
+struct spvDescriptorSetBuffer0
+{
+    constant Foo* m_34 [[id(0)]];
+    constant Bar* m_40 [[id(1)]];
+};
+
+struct spvDescriptorSetBuffer1
+{
+    device Baz* baz [[id(0)]][18];
+};
+
+kernel void main0(constant spvDescriptorSetBuffer0& spvDescriptorSet0 [[buffer(0)]], constant spvDescriptorSetBuffer1& spvDescriptorSet1 [[buffer(1)]], constant uint* spvDynamicOffsets [[buffer(23)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
+{
+    constant auto& _34 = *(constant Foo* )((constant char* )spvDescriptorSet0.m_34 + spvDynamicOffsets[0]);
+    device Baz* baz[18] =
+    {
+        (device Baz* )((device char* )spvDescriptorSet1.baz[0] + spvDynamicOffsets[1]),
+        (device Baz* )((device char* )spvDescriptorSet1.baz[1] + spvDynamicOffsets[2]),
+        (device Baz* )((device char* )spvDescriptorSet1.baz[2] + spvDynamicOffsets[3]),
+        (device Baz* )((device char* )spvDescriptorSet1.baz[3] + spvDynamicOffsets[4]),
+        (device Baz* )((device char* )spvDescriptorSet1.baz[4] + spvDynamicOffsets[5]),
+        (device Baz* )((device char* )spvDescriptorSet1.baz[5] + spvDynamicOffsets[6]),
+        (device Baz* )((device char* )spvDescriptorSet1.baz[6] + spvDynamicOffsets[7]),
+        (device Baz* )((device char* )spvDescriptorSet1.baz[7] + spvDynamicOffsets[8]),
+        (device Baz* )((device char* )spvDescriptorSet1.baz[8] + spvDynamicOffsets[9]),
+        (device Baz* )((device char* )spvDescriptorSet1.baz[9] + spvDynamicOffsets[10]),
+        (device Baz* )((device char* )spvDescriptorSet1.baz[10] + spvDynamicOffsets[11]),
+        (device Baz* )((device char* )spvDescriptorSet1.baz[11] + spvDynamicOffsets[12]),
+        (device Baz* )((device char* )spvDescriptorSet1.baz[12] + spvDynamicOffsets[13]),
+        (device Baz* )((device char* )spvDescriptorSet1.baz[13] + spvDynamicOffsets[14]),
+        (device Baz* )((device char* )spvDescriptorSet1.baz[14] + spvDynamicOffsets[15]),
+        (device Baz* )((device char* )spvDescriptorSet1.baz[15] + spvDynamicOffsets[16]),
+        (device Baz* )((device char* )spvDescriptorSet1.baz[16] + spvDynamicOffsets[17]),
+        (device Baz* )((device char* )spvDescriptorSet1.baz[17] + spvDynamicOffsets[18]),
+    };
+
+    uint3 coords = gl_GlobalInvocationID;
+    baz[(coords.x + coords.y) + coords.z]->e = _34.a + (*spvDescriptorSet0.m_40).c;
+    baz[(coords.x + coords.y) + coords.z]->f = _34.b * (*spvDescriptorSet0.m_40).d;
+}
+

shaders-msl-no-opt/comp/basic.dynamic-buffer.msl2.invalid.comp → shaders-msl-no-opt/comp/basic.dynamic-buffer.msl2.comp

@@ -17,11 +17,11 @@ layout(set = 1, binding = 2) buffer Baz
 {
 	int e;
 	int f;
-} baz[3][3][2];
+} baz[3 * 3 * 2];
 
 void main()
 {
 	uvec3 coords = gl_GlobalInvocationID;
-	baz[coords.x][coords.y][coords.z].e = a + c;
-	baz[coords.x][coords.y][coords.z].f = b * d;
+	baz[coords.x + coords.y + coords.z].e = a + c;
+	baz[coords.x + coords.y + coords.z].f = b * d;
 }

spirv_msl.cpp

@@ -1314,48 +1314,29 @@ void CompilerMSL::emit_entry_point_declarations()
 		uint32_t arg_id = argument_buffer_ids[desc_set];
 		uint32_t base_index = dynamic_buffer.second.first;
 
-		if (!type.array.empty())
+		if (is_array(type))
 		{
-			// This is complicated, because we need to support arrays of arrays.
-			// And it's even worse if the outermost dimension is a runtime array, because now
-			// all this complicated goop has to go into the shader itself. (FIXME)
 			if (!type.array[type.array.size() - 1])
 				SPIRV_CROSS_THROW("Runtime arrays with dynamic offsets are not supported yet.");
-			else
-			{
-				is_using_builtin_array = true;
-				statement(get_argument_address_space(var), " ", type_to_glsl(type), "* ", to_restrict(var_id, true), name,
-				          type_to_array_glsl(type), " =");
 
-				uint32_t dim = uint32_t(type.array.size());
-				uint32_t j = 0;
-				for (SmallVector<uint32_t> indices(type.array.size());
-				     indices[type.array.size() - 1] < to_array_size_literal(type); j++)
-				{
-					while (dim > 0)
-					{
-						begin_scope();
-						--dim;
-					}
+			is_using_builtin_array = true;
+			statement(get_argument_address_space(var), " ", type_to_glsl(type), "* ", to_restrict(var_id, true), name,
+			          type_to_array_glsl(type), " =");
 
-					string arrays;
-					for (uint32_t i = uint32_t(type.array.size()); i; --i)
-						arrays += join("[", indices[i - 1], "]");
-					statement("(", get_argument_address_space(var), " ", type_to_glsl(type), "* ",
-					          to_restrict(var_id, false), ")((", get_argument_address_space(var), " char* ",
-					          to_restrict(var_id, false), ")", to_name(arg_id), ".", ensure_valid_name(name, "m"),
-					          arrays, " + ", to_name(dynamic_offsets_buffer_id), "[", base_index + j, "]),");
+			uint32_t array_size = to_array_size_literal(type);
+			begin_scope();
 
-					while (++indices[dim] >= to_array_size_literal(type, dim) && dim < type.array.size() - 1)
-					{
-						end_scope(",");
-						indices[dim++] = 0;
-					}
-				}
-				end_scope_decl();
-				statement_no_indent("");
-				is_using_builtin_array = false;
+			for (uint32_t i = 0; i < array_size; i++)
+			{
+				statement("(", get_argument_address_space(var), " ", type_to_glsl(type), "* ",
+				          to_restrict(var_id, false), ")((", get_argument_address_space(var), " char* ",
+				          to_restrict(var_id, false), ")", to_name(arg_id), ".", ensure_valid_name(name, "m"),
+				          "[", i, "]", " + ", to_name(dynamic_offsets_buffer_id), "[", base_index + i, "]),");
 			}
+
+			end_scope_decl();
+			statement_no_indent("");
+			is_using_builtin_array = false;
 		}
 		else
 		{