MSL: Handle constant construct of block-like array types.

Need this to be context sensitive, since array of block-like struct is
template, but struct of block-like array is C-style.

Also, test a mix and match, so we have constant array of block-like
struct with array inside. :v

2 files changed, 154 insertions, 0 deletions

diff --git a/reference/shaders-msl-no-opt/asm/comp/block-like-array-type-construct-2.asm.comp b/reference/shaders-msl-no-opt/asm/comp/block-like-array-type-construct-2.asm.comp
new file mode 100644
index 00000000..734a6687
--- /dev/null
+++ b/reference/shaders-msl-no-opt/asm/comp/block-like-array-type-construct-2.asm.comp
@@ -0,0 +1,77 @@
+#pragma clang diagnostic ignored "-Wmissing-prototypes"
+#pragma clang diagnostic ignored "-Wmissing-braces"
+
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+template<typename T, size_t Num>
+struct spvUnsafeArray
+{
+    T elements[Num ? Num : 1];
+    
+    thread T& operator [] (size_t pos) thread
+    {
+        return elements[pos];
+    }
+    constexpr const thread T& operator [] (size_t pos) const thread
+    {
+        return elements[pos];
+    }
+    
+    device T& operator [] (size_t pos) device
+    {
+        return elements[pos];
+    }
+    constexpr const device T& operator [] (size_t pos) const device
+    {
+        return elements[pos];
+    }
+    
+    constexpr const constant T& operator [] (size_t pos) const constant
+    {
+        return elements[pos];
+    }
+    
+    threadgroup T& operator [] (size_t pos) threadgroup
+    {
+        return elements[pos];
+    }
+    constexpr const threadgroup T& operator [] (size_t pos) const threadgroup
+    {
+        return elements[pos];
+    }
+};
+
+struct type_CommonConstants
+{
+    uint g_count;
+    packed_uint3 g_padding4;
+};
+
+struct MyStruct
+{
+    float4 m_coefficients[4];
+};
+
+struct type_RWStructuredBuffer_MyStruct
+{
+    MyStruct _m0[1];
+};
+
+constant spvUnsafeArray<float4, 4> _27 = spvUnsafeArray<float4, 4>({ float4(0.0), float4(0.0), float4(0.0), float4(0.0) });
+
+kernel void main0(constant type_CommonConstants& CommonConstants [[buffer(0)]], device type_RWStructuredBuffer_MyStruct& g_data [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
+{
+    do
+    {
+        if (gl_GlobalInvocationID.x >= CommonConstants.g_count)
+        {
+            break;
+        }
+        g_data._m0[gl_GlobalInvocationID.x] = MyStruct{ { float4(0.0), float4(0.0), float4(0.0), float4(0.0) } };
+        break;
+    } while(false);
+}
+
diff --git a/reference/shaders-msl-no-opt/asm/comp/block-like-array-type-construct.asm.comp b/reference/shaders-msl-no-opt/asm/comp/block-like-array-type-construct.asm.comp
new file mode 100644
index 00000000..66550535
--- /dev/null
+++ b/reference/shaders-msl-no-opt/asm/comp/block-like-array-type-construct.asm.comp
@@ -0,0 +1,77 @@
+#pragma clang diagnostic ignored "-Wmissing-prototypes"
+#pragma clang diagnostic ignored "-Wmissing-braces"
+
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+template<typename T, size_t Num>
+struct spvUnsafeArray
+{
+    T elements[Num ? Num : 1];
+    
+    thread T& operator [] (size_t pos) thread
+    {
+        return elements[pos];
+    }
+    constexpr const thread T& operator [] (size_t pos) const thread
+    {
+        return elements[pos];
+    }
+    
+    device T& operator [] (size_t pos) device
+    {
+        return elements[pos];
+    }
+    constexpr const device T& operator [] (size_t pos) const device
+    {
+        return elements[pos];
+    }
+    
+    constexpr const constant T& operator [] (size_t pos) const constant
+    {
+        return elements[pos];
+    }
+    
+    threadgroup T& operator [] (size_t pos) threadgroup
+    {
+        return elements[pos];
+    }
+    constexpr const threadgroup T& operator [] (size_t pos) const threadgroup
+    {
+        return elements[pos];
+    }
+};
+
+struct _10
+{
+    float _m0[4];
+    float _m1[4];
+};
+
+constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u);
+
+struct SSBO
+{
+    uint a;
+    int b;
+};
+
+constant spvUnsafeArray<float, 4> _31 = spvUnsafeArray<float, 4>({ 1.0, 2.0, 3.0, 4.0 });
+
+kernel void main0()
+{
+    spvUnsafeArray<_10, 2> _34 = spvUnsafeArray<_10, 2>({ _10{ { 1.0, 2.0, 3.0, 4.0 }, { 1.0, 2.0, 3.0, 4.0 } }, _10{ { 1.0, 2.0, 3.0, 4.0 }, { 1.0, 2.0, 3.0, 4.0 } } });
+    
+    spvUnsafeArray<float, 4> foo;
+    foo[0] = 1.0;
+    foo = _31;
+    foo[1] = 2.0;
+    foo[2] = 3.0;
+    foo[3] = 4.0;
+    spvUnsafeArray<float, 4> foo2;
+    foo2 = foo;
+    _10 _37 = _10{ { foo[0], foo[1], foo[2], foo[3] }, { foo2[0], foo2[1], foo2[2], foo2[3] } };
+}
+