Merge pull request #559 from KhronosGroup/fix-558

Fix OpCompositeConstruct with arrays in MSL.

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

@@ -0,0 +1,35 @@
+#pragma clang diagnostic ignored "-Wmissing-prototypes"
+
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+struct SSBO0
+{
+    float4 as[1];
+};
+
+// 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])
+{
+    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])
+{
+    for (uint i = 0; i < N; dst[i] = src[i], i++);
+}
+
+kernel void main0(device SSBO0& _16 [[buffer(0)]], device SSBO0& _32 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint gl_LocalInvocationIndex [[thread_index_in_threadgroup]])
+{
+    float4 _37[2] = { _16.as[gl_GlobalInvocationID.x], _32.as[gl_GlobalInvocationID.x] };
+    float4 values[2];
+    spvArrayCopy(values, _37);
+    _16.as[0] = values[gl_LocalInvocationIndex];
+    _32.as[1] = float4(40.0);
+}
+

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

@@ -0,0 +1,50 @@
+#pragma clang diagnostic ignored "-Wmissing-prototypes"
+
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+struct SSBO0
+{
+    float4 as[1];
+};
+
+struct SSBO1
+{
+    float4 bs[1];
+};
+
+struct Composite
+{
+    float4 a;
+    float4 b;
+};
+
+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])
+{
+    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])
+{
+    for (uint i = 0; i < N; dst[i] = src[i], i++);
+}
+
+kernel void main0(device SSBO0& _16 [[buffer(0)]], device SSBO1& _32 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint gl_LocalInvocationIndex [[thread_index_in_threadgroup]])
+{
+    float4 _37[2] = { _16.as[gl_GlobalInvocationID.x], _32.bs[gl_GlobalInvocationID.x] };
+    float4 values[2];
+    spvArrayCopy(values, _37);
+    float4 copy_values[2] = {float4(20.0), float4(40.0)};
+    Composite c = Composite{ values[0], copy_values[1] };
+    _16.as[0] = values[gl_LocalInvocationIndex];
+    _32.bs[1] = c.b;
+}
+

shaders-msl/comp/composite-construct.comp

@@ -0,0 +1,31 @@
+#version 310 es
+layout(local_size_x = 1) in;
+
+layout(std430, binding = 0) buffer SSBO0
+{
+   vec4 as[];
+};
+
+layout(std430, binding = 1) buffer SSBO1
+{
+   vec4 bs[];
+};
+
+struct Composite
+{
+   vec4 a;
+   vec4 b;
+};
+
+const vec4 const_values[2] = vec4[](vec4(20.0), vec4(40.0));
+
+void main()
+{
+   vec4 values[2] = vec4[](as[gl_GlobalInvocationID.x], bs[gl_GlobalInvocationID.x]);
+   vec4 copy_values[2];
+   copy_values = const_values;
+   Composite c = Composite(values[0], copy_values[1]);
+
+   as[0] = values[gl_LocalInvocationIndex];
+   bs[1] = c.b;
+}

spirv_glsl.cpp

@@ -6289,7 +6289,7 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction)
 		{
 			// Only use this path if we are building composites.
 			// This path cannot be used for arithmetic.
-			if (backend.use_typed_initializer_list)
+			if (backend.use_typed_initializer_list && out_type.basetype == SPIRType::Struct)
 				constructor_op += type_to_glsl_constructor(get<SPIRType>(result_type));
 			constructor_op += "{ ";
 			if (type_is_empty(out_type) && !backend.supports_empty_struct)