MSL: Add tests for array copies in and out of buffers.

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

@@ -81,6 +81,60 @@ inline void spvArrayCopyFromThreadGroupToThreadGroup1(threadgroup T (&dst)[A], t
     }
 }
 
+template<typename T, uint A>
+inline void spvArrayCopyFromDeviceToDevice1(device T (&dst)[A], device const T (&src)[A])
+{
+    for (uint i = 0; i < A; i++)
+    {
+        dst[i] = src[i];
+    }
+}
+
+template<typename T, uint A>
+inline void spvArrayCopyFromConstantToDevice1(device T (&dst)[A], constant T (&src)[A])
+{
+    for (uint i = 0; i < A; i++)
+    {
+        dst[i] = src[i];
+    }
+}
+
+template<typename T, uint A>
+inline void spvArrayCopyFromStackToDevice1(device T (&dst)[A], thread const T (&src)[A])
+{
+    for (uint i = 0; i < A; i++)
+    {
+        dst[i] = src[i];
+    }
+}
+
+template<typename T, uint A>
+inline void spvArrayCopyFromThreadGroupToDevice1(device T (&dst)[A], threadgroup const T (&src)[A])
+{
+    for (uint i = 0; i < A; i++)
+    {
+        dst[i] = src[i];
+    }
+}
+
+template<typename T, uint A>
+inline void spvArrayCopyFromDeviceToStack1(thread T (&dst)[A], device const T (&src)[A])
+{
+    for (uint i = 0; i < A; i++)
+    {
+        dst[i] = src[i];
+    }
+}
+
+template<typename T, uint A>
+inline void spvArrayCopyFromDeviceToThreadGroup1(threadgroup T (&dst)[A], device const T (&src)[A])
+{
+    for (uint i = 0; i < A; i++)
+    {
+        dst[i] = src[i];
+    }
+}
+
 kernel void main0(device SSBO& _53 [[buffer(0)]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]])
 {
     Data _25[2] = { Data{ 1.0, 2.0 }, Data{ 3.0, 4.0 } };

reference/shaders-msl-no-opt/asm/comp/device-constant-array-load-store.asm.comp

@@ -0,0 +1,179 @@
+#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 Block
+{
+    uint2 _m0[2];
+    uint2 _m1[2];
+};
+
+struct SSBO
+{
+    Block _m0[3];
+};
+
+template<typename T, uint A>
+inline void spvArrayCopyFromConstantToStack1(thread T (&dst)[A], constant T (&src)[A])
+{
+    for (uint i = 0; i < A; i++)
+    {
+        dst[i] = src[i];
+    }
+}
+
+template<typename T, uint A>
+inline void spvArrayCopyFromConstantToThreadGroup1(threadgroup T (&dst)[A], constant T (&src)[A])
+{
+    for (uint i = 0; i < A; i++)
+    {
+        dst[i] = src[i];
+    }
+}
+
+template<typename T, uint A>
+inline void spvArrayCopyFromStackToStack1(thread T (&dst)[A], thread const T (&src)[A])
+{
+    for (uint i = 0; i < A; i++)
+    {
+        dst[i] = src[i];
+    }
+}
+
+template<typename T, uint A>
+inline void spvArrayCopyFromStackToThreadGroup1(threadgroup T (&dst)[A], thread const T (&src)[A])
+{
+    for (uint i = 0; i < A; i++)
+    {
+        dst[i] = src[i];
+    }
+}
+
+template<typename T, uint A>
+inline void spvArrayCopyFromThreadGroupToStack1(thread T (&dst)[A], threadgroup const T (&src)[A])
+{
+    for (uint i = 0; i < A; i++)
+    {
+        dst[i] = src[i];
+    }
+}
+
+template<typename T, uint A>
+inline void spvArrayCopyFromThreadGroupToThreadGroup1(threadgroup T (&dst)[A], threadgroup const T (&src)[A])
+{
+    for (uint i = 0; i < A; i++)
+    {
+        dst[i] = src[i];
+    }
+}
+
+template<typename T, uint A>
+inline void spvArrayCopyFromDeviceToDevice1(device T (&dst)[A], device const T (&src)[A])
+{
+    for (uint i = 0; i < A; i++)
+    {
+        dst[i] = src[i];
+    }
+}
+
+template<typename T, uint A>
+inline void spvArrayCopyFromConstantToDevice1(device T (&dst)[A], constant T (&src)[A])
+{
+    for (uint i = 0; i < A; i++)
+    {
+        dst[i] = src[i];
+    }
+}
+
+template<typename T, uint A>
+inline void spvArrayCopyFromStackToDevice1(device T (&dst)[A], thread const T (&src)[A])
+{
+    for (uint i = 0; i < A; i++)
+    {
+        dst[i] = src[i];
+    }
+}
+
+template<typename T, uint A>
+inline void spvArrayCopyFromThreadGroupToDevice1(device T (&dst)[A], threadgroup const T (&src)[A])
+{
+    for (uint i = 0; i < A; i++)
+    {
+        dst[i] = src[i];
+    }
+}
+
+template<typename T, uint A>
+inline void spvArrayCopyFromDeviceToStack1(thread T (&dst)[A], device const T (&src)[A])
+{
+    for (uint i = 0; i < A; i++)
+    {
+        dst[i] = src[i];
+    }
+}
+
+template<typename T, uint A>
+inline void spvArrayCopyFromDeviceToThreadGroup1(threadgroup T (&dst)[A], device const T (&src)[A])
+{
+    for (uint i = 0; i < A; i++)
+    {
+        dst[i] = src[i];
+    }
+}
+
+kernel void main0(device SSBO& ssbo [[buffer(0)]], constant SSBO& ubo [[buffer(1)]])
+{
+    threadgroup uint2 _18[2];
+    spvArrayCopyFromDeviceToDevice1(ssbo._m0[0u]._m0, ssbo._m0[0u]._m1);
+    spvArrayCopyFromConstantToDevice1(ssbo._m0[0u]._m0, ubo._m0[0u]._m1);
+    spvUnsafeArray<uint2, 2> _24;
+    spvArrayCopyFromStackToDevice1(ssbo._m0[0u]._m0, _24.elements);
+    spvArrayCopyFromThreadGroupToDevice1(ssbo._m0[0u]._m0, _18);
+    spvArrayCopyFromDeviceToThreadGroup1(_18, ssbo._m0[0u]._m1);
+    spvArrayCopyFromDeviceToStack1(_24.elements, ssbo._m0[0u]._m1);
+    spvArrayCopyFromConstantToThreadGroup1(_18, ubo._m0[0u]._m1);
+    spvArrayCopyFromConstantToStack1(_24.elements, ubo._m0[0u]._m1);
+}
+

reference/shaders-msl-no-opt/asm/comp/device-constant-array-load-store.force-native-array.asm.comp

@@ -0,0 +1,140 @@
+#pragma clang diagnostic ignored "-Wmissing-prototypes"
+
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+struct Block
+{
+    uint2 _m0[2];
+    uint2 _m1[2];
+};
+
+struct SSBO
+{
+    Block _m0[3];
+};
+
+template<typename T, uint A>
+inline void spvArrayCopyFromConstantToStack1(thread T (&dst)[A], constant T (&src)[A])
+{
+    for (uint i = 0; i < A; i++)
+    {
+        dst[i] = src[i];
+    }
+}
+
+template<typename T, uint A>
+inline void spvArrayCopyFromConstantToThreadGroup1(threadgroup T (&dst)[A], constant T (&src)[A])
+{
+    for (uint i = 0; i < A; i++)
+    {
+        dst[i] = src[i];
+    }
+}
+
+template<typename T, uint A>
+inline void spvArrayCopyFromStackToStack1(thread T (&dst)[A], thread const T (&src)[A])
+{
+    for (uint i = 0; i < A; i++)
+    {
+        dst[i] = src[i];
+    }
+}
+
+template<typename T, uint A>
+inline void spvArrayCopyFromStackToThreadGroup1(threadgroup T (&dst)[A], thread const T (&src)[A])
+{
+    for (uint i = 0; i < A; i++)
+    {
+        dst[i] = src[i];
+    }
+}
+
+template<typename T, uint A>
+inline void spvArrayCopyFromThreadGroupToStack1(thread T (&dst)[A], threadgroup const T (&src)[A])
+{
+    for (uint i = 0; i < A; i++)
+    {
+        dst[i] = src[i];
+    }
+}
+
+template<typename T, uint A>
+inline void spvArrayCopyFromThreadGroupToThreadGroup1(threadgroup T (&dst)[A], threadgroup const T (&src)[A])
+{
+    for (uint i = 0; i < A; i++)
+    {
+        dst[i] = src[i];
+    }
+}
+
+template<typename T, uint A>
+inline void spvArrayCopyFromDeviceToDevice1(device T (&dst)[A], device const T (&src)[A])
+{
+    for (uint i = 0; i < A; i++)
+    {
+        dst[i] = src[i];
+    }
+}
+
+template<typename T, uint A>
+inline void spvArrayCopyFromConstantToDevice1(device T (&dst)[A], constant T (&src)[A])
+{
+    for (uint i = 0; i < A; i++)
+    {
+        dst[i] = src[i];
+    }
+}
+
+template<typename T, uint A>
+inline void spvArrayCopyFromStackToDevice1(device T (&dst)[A], thread const T (&src)[A])
+{
+    for (uint i = 0; i < A; i++)
+    {
+        dst[i] = src[i];
+    }
+}
+
+template<typename T, uint A>
+inline void spvArrayCopyFromThreadGroupToDevice1(device T (&dst)[A], threadgroup const T (&src)[A])
+{
+    for (uint i = 0; i < A; i++)
+    {
+        dst[i] = src[i];
+    }
+}
+
+template<typename T, uint A>
+inline void spvArrayCopyFromDeviceToStack1(thread T (&dst)[A], device const T (&src)[A])
+{
+    for (uint i = 0; i < A; i++)
+    {
+        dst[i] = src[i];
+    }
+}
+
+template<typename T, uint A>
+inline void spvArrayCopyFromDeviceToThreadGroup1(threadgroup T (&dst)[A], device const T (&src)[A])
+{
+    for (uint i = 0; i < A; i++)
+    {
+        dst[i] = src[i];
+    }
+}
+
+kernel void main0(device SSBO& ssbo [[buffer(0)]], constant SSBO& ubo [[buffer(1)]])
+{
+    threadgroup uint2 _18[2];
+    spvArrayCopyFromDeviceToDevice1(ssbo._m0[0u]._m0, ssbo._m0[0u]._m1);
+    spvArrayCopyFromConstantToDevice1(ssbo._m0[0u]._m0, ubo._m0[0u]._m1);
+    uint2 _24[2];
+    spvArrayCopyFromStackToDevice1(ssbo._m0[0u]._m0, _24);
+    spvArrayCopyFromThreadGroupToDevice1(ssbo._m0[0u]._m0, _18);
+    spvArrayCopyFromDeviceToThreadGroup1(_18, ssbo._m0[0u]._m1);
+    spvArrayCopyFromDeviceToStack1(_24, ssbo._m0[0u]._m1);
+    spvArrayCopyFromConstantToThreadGroup1(_18, ubo._m0[0u]._m1);
+    spvArrayCopyFromConstantToStack1(_24, ubo._m0[0u]._m1);
+}
+

reference/shaders-msl-no-opt/comp/array-copy-threadgroup-memory.comp

@@ -100,6 +100,60 @@ inline void spvArrayCopyFromThreadGroupToThreadGroup1(threadgroup T (&dst)[A], t
     }
 }
 
+template<typename T, uint A>
+inline void spvArrayCopyFromDeviceToDevice1(device T (&dst)[A], device const T (&src)[A])
+{
+    for (uint i = 0; i < A; i++)
+    {
+        dst[i] = src[i];
+    }
+}
+
+template<typename T, uint A>
+inline void spvArrayCopyFromConstantToDevice1(device T (&dst)[A], constant T (&src)[A])
+{
+    for (uint i = 0; i < A; i++)
+    {
+        dst[i] = src[i];
+    }
+}
+
+template<typename T, uint A>
+inline void spvArrayCopyFromStackToDevice1(device T (&dst)[A], thread const T (&src)[A])
+{
+    for (uint i = 0; i < A; i++)
+    {
+        dst[i] = src[i];
+    }
+}
+
+template<typename T, uint A>
+inline void spvArrayCopyFromThreadGroupToDevice1(device T (&dst)[A], threadgroup const T (&src)[A])
+{
+    for (uint i = 0; i < A; i++)
+    {
+        dst[i] = src[i];
+    }
+}
+
+template<typename T, uint A>
+inline void spvArrayCopyFromDeviceToStack1(thread T (&dst)[A], device const T (&src)[A])
+{
+    for (uint i = 0; i < A; i++)
+    {
+        dst[i] = src[i];
+    }
+}
+
+template<typename T, uint A>
+inline void spvArrayCopyFromDeviceToThreadGroup1(threadgroup T (&dst)[A], device const T (&src)[A])
+{
+    for (uint i = 0; i < A; i++)
+    {
+        dst[i] = src[i];
+    }
+}
+
 kernel void main0(uint gl_LocalInvocationIndex [[thread_index_in_threadgroup]])
 {
     threadgroup float shared_group[8][8];

reference/shaders-msl-no-opt/vert/pass-array-by-value.force-native-array.vert

@@ -72,6 +72,60 @@ inline void spvArrayCopyFromThreadGroupToThreadGroup1(threadgroup T (&dst)[A], t
     }
 }
 
+template<typename T, uint A>
+inline void spvArrayCopyFromDeviceToDevice1(device T (&dst)[A], device const T (&src)[A])
+{
+    for (uint i = 0; i < A; i++)
+    {
+        dst[i] = src[i];
+    }
+}
+
+template<typename T, uint A>
+inline void spvArrayCopyFromConstantToDevice1(device T (&dst)[A], constant T (&src)[A])
+{
+    for (uint i = 0; i < A; i++)
+    {
+        dst[i] = src[i];
+    }
+}
+
+template<typename T, uint A>
+inline void spvArrayCopyFromStackToDevice1(device T (&dst)[A], thread const T (&src)[A])
+{
+    for (uint i = 0; i < A; i++)
+    {
+        dst[i] = src[i];
+    }
+}
+
+template<typename T, uint A>
+inline void spvArrayCopyFromThreadGroupToDevice1(device T (&dst)[A], threadgroup const T (&src)[A])
+{
+    for (uint i = 0; i < A; i++)
+    {
+        dst[i] = src[i];
+    }
+}
+
+template<typename T, uint A>
+inline void spvArrayCopyFromDeviceToStack1(thread T (&dst)[A], device const T (&src)[A])
+{
+    for (uint i = 0; i < A; i++)
+    {
+        dst[i] = src[i];
+    }
+}
+
+template<typename T, uint A>
+inline void spvArrayCopyFromDeviceToThreadGroup1(threadgroup T (&dst)[A], device const T (&src)[A])
+{
+    for (uint i = 0; i < A; i++)
+    {
+        dst[i] = src[i];
+    }
+}
+
 static inline __attribute__((always_inline))
 float4 consume_constant_arrays2(thread const float4 (&positions)[4], thread const float4 (&positions2)[4], thread int& Index1, thread int& Index2)
 {

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

@@ -83,6 +83,60 @@ inline void spvArrayCopyFromThreadGroupToThreadGroup1(threadgroup T (&dst)[A], t
     }
 }
 
+template<typename T, uint A>
+inline void spvArrayCopyFromDeviceToDevice1(device T (&dst)[A], device const T (&src)[A])
+{
+    for (uint i = 0; i < A; i++)
+    {
+        dst[i] = src[i];
+    }
+}
+
+template<typename T, uint A>
+inline void spvArrayCopyFromConstantToDevice1(device T (&dst)[A], constant T (&src)[A])
+{
+    for (uint i = 0; i < A; i++)
+    {
+        dst[i] = src[i];
+    }
+}
+
+template<typename T, uint A>
+inline void spvArrayCopyFromStackToDevice1(device T (&dst)[A], thread const T (&src)[A])
+{
+    for (uint i = 0; i < A; i++)
+    {
+        dst[i] = src[i];
+    }
+}
+
+template<typename T, uint A>
+inline void spvArrayCopyFromThreadGroupToDevice1(device T (&dst)[A], threadgroup const T (&src)[A])
+{
+    for (uint i = 0; i < A; i++)
+    {
+        dst[i] = src[i];
+    }
+}
+
+template<typename T, uint A>
+inline void spvArrayCopyFromDeviceToStack1(thread T (&dst)[A], device const T (&src)[A])
+{
+    for (uint i = 0; i < A; i++)
+    {
+        dst[i] = src[i];
+    }
+}
+
+template<typename T, uint A>
+inline void spvArrayCopyFromDeviceToThreadGroup1(threadgroup T (&dst)[A], device const T (&src)[A])
+{
+    for (uint i = 0; i < A; i++)
+    {
+        dst[i] = src[i];
+    }
+}
+
 static inline __attribute__((always_inline))
 Data combine(thread const Data& a, thread const Data& b)
 {

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

@@ -73,6 +73,60 @@ inline void spvArrayCopyFromThreadGroupToThreadGroup1(threadgroup T (&dst)[A], t
     }
 }
 
+template<typename T, uint A>
+inline void spvArrayCopyFromDeviceToDevice1(device T (&dst)[A], device const T (&src)[A])
+{
+    for (uint i = 0; i < A; i++)
+    {
+        dst[i] = src[i];
+    }
+}
+
+template<typename T, uint A>
+inline void spvArrayCopyFromConstantToDevice1(device T (&dst)[A], constant T (&src)[A])
+{
+    for (uint i = 0; i < A; i++)
+    {
+        dst[i] = src[i];
+    }
+}
+
+template<typename T, uint A>
+inline void spvArrayCopyFromStackToDevice1(device T (&dst)[A], thread const T (&src)[A])
+{
+    for (uint i = 0; i < A; i++)
+    {
+        dst[i] = src[i];
+    }
+}
+
+template<typename T, uint A>
+inline void spvArrayCopyFromThreadGroupToDevice1(device T (&dst)[A], threadgroup const T (&src)[A])
+{
+    for (uint i = 0; i < A; i++)
+    {
+        dst[i] = src[i];
+    }
+}
+
+template<typename T, uint A>
+inline void spvArrayCopyFromDeviceToStack1(thread T (&dst)[A], device const T (&src)[A])
+{
+    for (uint i = 0; i < A; i++)
+    {
+        dst[i] = src[i];
+    }
+}
+
+template<typename T, uint A>
+inline void spvArrayCopyFromDeviceToThreadGroup1(threadgroup T (&dst)[A], device const T (&src)[A])
+{
+    for (uint i = 0; i < A; i++)
+    {
+        dst[i] = src[i];
+    }
+}
+
 template<typename T, uint A, uint B>
 inline void spvArrayCopyFromConstantToStack2(thread T (&dst)[A][B], constant T (&src)[A][B])
 {
@@ -127,6 +181,60 @@ inline void spvArrayCopyFromThreadGroupToThreadGroup2(threadgroup T (&dst)[A][B]
     }
 }
 
+template<typename T, uint A, uint B>
+inline void spvArrayCopyFromDeviceToDevice2(device T (&dst)[A][B], device const T (&src)[A][B])
+{
+    for (uint i = 0; i < A; i++)
+    {
+        spvArrayCopyFromDeviceToDevice1(dst[i], src[i]);
+    }
+}
+
+template<typename T, uint A, uint B>
+inline void spvArrayCopyFromConstantToDevice2(device T (&dst)[A][B], constant T (&src)[A][B])
+{
+    for (uint i = 0; i < A; i++)
+    {
+        spvArrayCopyFromConstantToDevice1(dst[i], src[i]);
+    }
+}
+
+template<typename T, uint A, uint B>
+inline void spvArrayCopyFromStackToDevice2(device T (&dst)[A][B], thread const T (&src)[A][B])
+{
+    for (uint i = 0; i < A; i++)
+    {
+        spvArrayCopyFromStackToDevice1(dst[i], src[i]);
+    }
+}
+
+template<typename T, uint A, uint B>
+inline void spvArrayCopyFromThreadGroupToDevice2(device T (&dst)[A][B], threadgroup const T (&src)[A][B])
+{
+    for (uint i = 0; i < A; i++)
+    {
+        spvArrayCopyFromThreadGroupToDevice1(dst[i], src[i]);
+    }
+}
+
+template<typename T, uint A, uint B>
+inline void spvArrayCopyFromDeviceToStack2(thread T (&dst)[A][B], device const T (&src)[A][B])
+{
+    for (uint i = 0; i < A; i++)
+    {
+        spvArrayCopyFromDeviceToStack1(dst[i], src[i]);
+    }
+}
+
+template<typename T, uint A, uint B>
+inline void spvArrayCopyFromDeviceToThreadGroup2(threadgroup T (&dst)[A][B], device const T (&src)[A][B])
+{
+    for (uint i = 0; i < A; i++)
+    {
+        spvArrayCopyFromDeviceToThreadGroup1(dst[i], src[i]);
+    }
+}
+
 template<typename T, uint A, uint B, uint C>
 inline void spvArrayCopyFromConstantToStack3(thread T (&dst)[A][B][C], constant T (&src)[A][B][C])
 {
@@ -181,6 +289,60 @@ inline void spvArrayCopyFromThreadGroupToThreadGroup3(threadgroup T (&dst)[A][B]
     }
 }
 
+template<typename T, uint A, uint B, uint C>
+inline void spvArrayCopyFromDeviceToDevice3(device T (&dst)[A][B][C], device const T (&src)[A][B][C])
+{
+    for (uint i = 0; i < A; i++)
+    {
+        spvArrayCopyFromDeviceToDevice2(dst[i], src[i]);
+    }
+}
+
+template<typename T, uint A, uint B, uint C>
+inline void spvArrayCopyFromConstantToDevice3(device T (&dst)[A][B][C], constant T (&src)[A][B][C])
+{
+    for (uint i = 0; i < A; i++)
+    {
+        spvArrayCopyFromConstantToDevice2(dst[i], src[i]);
+    }
+}
+
+template<typename T, uint A, uint B, uint C>
+inline void spvArrayCopyFromStackToDevice3(device T (&dst)[A][B][C], thread const T (&src)[A][B][C])
+{
+    for (uint i = 0; i < A; i++)
+    {
+        spvArrayCopyFromStackToDevice2(dst[i], src[i]);
+    }
+}
+
+template<typename T, uint A, uint B, uint C>
+inline void spvArrayCopyFromThreadGroupToDevice3(device T (&dst)[A][B][C], threadgroup const T (&src)[A][B][C])
+{
+    for (uint i = 0; i < A; i++)
+    {
+        spvArrayCopyFromThreadGroupToDevice2(dst[i], src[i]);
+    }
+}
+
+template<typename T, uint A, uint B, uint C>
+inline void spvArrayCopyFromDeviceToStack3(thread T (&dst)[A][B][C], device const T (&src)[A][B][C])
+{
+    for (uint i = 0; i < A; i++)
+    {
+        spvArrayCopyFromDeviceToStack2(dst[i], src[i]);
+    }
+}
+
+template<typename T, uint A, uint B, uint C>
+inline void spvArrayCopyFromDeviceToThreadGroup3(threadgroup T (&dst)[A][B][C], device const T (&src)[A][B][C])
+{
+    for (uint i = 0; i < A; i++)
+    {
+        spvArrayCopyFromDeviceToThreadGroup2(dst[i], src[i]);
+    }
+}
+
 kernel void main0(device BUF& o [[buffer(0)]])
 {
     float c[2][2][2];

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

@@ -72,6 +72,60 @@ inline void spvArrayCopyFromThreadGroupToThreadGroup1(threadgroup T (&dst)[A], t
     }
 }
 
+template<typename T, uint A>
+inline void spvArrayCopyFromDeviceToDevice1(device T (&dst)[A], device const T (&src)[A])
+{
+    for (uint i = 0; i < A; i++)
+    {
+        dst[i] = src[i];
+    }
+}
+
+template<typename T, uint A>
+inline void spvArrayCopyFromConstantToDevice1(device T (&dst)[A], constant T (&src)[A])
+{
+    for (uint i = 0; i < A; i++)
+    {
+        dst[i] = src[i];
+    }
+}
+
+template<typename T, uint A>
+inline void spvArrayCopyFromStackToDevice1(device T (&dst)[A], thread const T (&src)[A])
+{
+    for (uint i = 0; i < A; i++)
+    {
+        dst[i] = src[i];
+    }
+}
+
+template<typename T, uint A>
+inline void spvArrayCopyFromThreadGroupToDevice1(device T (&dst)[A], threadgroup const T (&src)[A])
+{
+    for (uint i = 0; i < A; i++)
+    {
+        dst[i] = src[i];
+    }
+}
+
+template<typename T, uint A>
+inline void spvArrayCopyFromDeviceToStack1(thread T (&dst)[A], device const T (&src)[A])
+{
+    for (uint i = 0; i < A; i++)
+    {
+        dst[i] = src[i];
+    }
+}
+
+template<typename T, uint A>
+inline void spvArrayCopyFromDeviceToThreadGroup1(threadgroup T (&dst)[A], device const T (&src)[A])
+{
+    for (uint i = 0; i < A; i++)
+    {
+        dst[i] = src[i];
+    }
+}
+
 static inline __attribute__((always_inline))
 void test(thread float4 (&SPIRV_Cross_return_value)[2])
 {

shaders-msl-no-opt/asm/comp/device-constant-array-load-store.asm.comp

@@ -0,0 +1,81 @@
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+			   OpName %Block "Block"
+			   OpName %SSBO "SSBO"
+			   OpName %SSBO_Var "ssbo"
+			   OpName %UBO_Var "ubo"
+			   OpDecorate %SSBO_Var Binding 0
+			   OpDecorate %SSBO_Var DescriptorSet 0
+			   OpDecorate %UBO_Var Binding 1
+			   OpDecorate %UBO_Var DescriptorSet 0
+			   OpMemberDecorate %SSBO 0 Offset 0
+			   OpMemberDecorate %Block 0 Offset 0
+			   OpMemberDecorate %Block 1 Offset 16
+			   OpDecorate %BlockArray ArrayStride 32
+			   OpDecorate %arr_uvec2_2 ArrayStride 8
+			   OpDecorate %SSBO Block
+       %uint = OpTypeInt 32 0
+     %uint_0 = OpConstant %uint 0
+     %uint_1 = OpConstant %uint 1
+     %uint_2 = OpConstant %uint 2
+     %uint_3 = OpConstant %uint 3
+     %uvec2 = OpTypeVector %uint 2
+     %arr_uvec2_2 = OpTypeArray %uvec2 %uint_2
+	 %arr_uvec2_2_ptr = OpTypePointer StorageBuffer %arr_uvec2_2
+	 %arr_uvec2_2_ptr_const = OpTypePointer Uniform %arr_uvec2_2
+	 %arr_uvec2_2_ptr_func = OpTypePointer Function %arr_uvec2_2
+	 %arr_uvec2_2_ptr_workgroup = OpTypePointer Workgroup %arr_uvec2_2
+	 %wg = OpVariable %arr_uvec2_2_ptr_workgroup Workgroup
+   %Block = OpTypeStruct %arr_uvec2_2 %arr_uvec2_2
+%BlockArray = OpTypeArray %Block %uint_3
+%SSBO = OpTypeStruct %BlockArray
+%SSBO_Ptr = OpTypePointer StorageBuffer %SSBO
+%SSBO_Var = OpVariable %SSBO_Ptr StorageBuffer
+%UBO_Ptr = OpTypePointer Uniform %SSBO
+%UBO_Var = OpVariable %UBO_Ptr Uniform
+%void = OpTypeVoid
+%func_type = OpTypeFunction %void
+
+    %main = OpFunction %void None %func_type
+         %25 = OpLabel
+		 %func = OpVariable %arr_uvec2_2_ptr_func Function
+
+		 ; DeviceToDevice
+		 %ptr_arr_0 = OpAccessChain %arr_uvec2_2_ptr %SSBO_Var %uint_0 %uint_0 %uint_0
+		 %ptr_arr_1 = OpAccessChain %arr_uvec2_2_ptr %SSBO_Var %uint_0 %uint_0 %uint_1
+		 %loaded_array = OpLoad %arr_uvec2_2 %ptr_arr_1
+		 OpStore %ptr_arr_0 %loaded_array
+
+		 ; ConstantToDevice
+		 %ptr_arr_1_const = OpAccessChain %arr_uvec2_2_ptr_const %UBO_Var %uint_0 %uint_0 %uint_1
+		 %loaded_array_const = OpLoad %arr_uvec2_2 %ptr_arr_1_const
+		 OpStore %ptr_arr_0 %loaded_array_const
+
+		 ; StackToDevice
+		 %loaded_array_func = OpLoad %arr_uvec2_2 %func
+		 OpStore %ptr_arr_0 %loaded_array_func
+
+		 ; ThreadGroupToDevice
+		 %loaded_array_workgroup = OpLoad %arr_uvec2_2 %wg
+		 OpStore %ptr_arr_0 %loaded_array_workgroup
+
+		 ; DeviceToThreadGroup
+		 %loaded_array_2 = OpLoad %arr_uvec2_2 %ptr_arr_1
+		 OpStore %wg %loaded_array_2
+
+		 ; DeviceToStack
+		 %loaded_array_3 = OpLoad %arr_uvec2_2 %ptr_arr_1
+		 OpStore %func %loaded_array_3
+
+		 ; ConstantToThreadGroup
+		 %loaded_array_const_2 = OpLoad %arr_uvec2_2 %ptr_arr_1_const
+		 OpStore %wg %loaded_array_const_2
+
+		 ; ConstantToStack
+		 %loaded_array_const_3 = OpLoad %arr_uvec2_2 %ptr_arr_1_const
+		 OpStore %func %loaded_array_const_3
+
+               OpReturn
+               OpFunctionEnd

shaders-msl-no-opt/asm/comp/device-constant-array-load-store.force-native-array.asm.comp

@@ -0,0 +1,81 @@
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+			   OpName %Block "Block"
+			   OpName %SSBO "SSBO"
+			   OpName %SSBO_Var "ssbo"
+			   OpName %UBO_Var "ubo"
+			   OpDecorate %SSBO_Var Binding 0
+			   OpDecorate %SSBO_Var DescriptorSet 0
+			   OpDecorate %UBO_Var Binding 1
+			   OpDecorate %UBO_Var DescriptorSet 0
+			   OpMemberDecorate %SSBO 0 Offset 0
+			   OpMemberDecorate %Block 0 Offset 0
+			   OpMemberDecorate %Block 1 Offset 16
+			   OpDecorate %BlockArray ArrayStride 32
+			   OpDecorate %arr_uvec2_2 ArrayStride 8
+			   OpDecorate %SSBO Block
+       %uint = OpTypeInt 32 0
+     %uint_0 = OpConstant %uint 0
+     %uint_1 = OpConstant %uint 1
+     %uint_2 = OpConstant %uint 2
+     %uint_3 = OpConstant %uint 3
+     %uvec2 = OpTypeVector %uint 2
+     %arr_uvec2_2 = OpTypeArray %uvec2 %uint_2
+	 %arr_uvec2_2_ptr = OpTypePointer StorageBuffer %arr_uvec2_2
+	 %arr_uvec2_2_ptr_const = OpTypePointer Uniform %arr_uvec2_2
+	 %arr_uvec2_2_ptr_func = OpTypePointer Function %arr_uvec2_2
+	 %arr_uvec2_2_ptr_workgroup = OpTypePointer Workgroup %arr_uvec2_2
+	 %wg = OpVariable %arr_uvec2_2_ptr_workgroup Workgroup
+   %Block = OpTypeStruct %arr_uvec2_2 %arr_uvec2_2
+%BlockArray = OpTypeArray %Block %uint_3
+%SSBO = OpTypeStruct %BlockArray
+%SSBO_Ptr = OpTypePointer StorageBuffer %SSBO
+%SSBO_Var = OpVariable %SSBO_Ptr StorageBuffer
+%UBO_Ptr = OpTypePointer Uniform %SSBO
+%UBO_Var = OpVariable %UBO_Ptr Uniform
+%void = OpTypeVoid
+%func_type = OpTypeFunction %void
+
+    %main = OpFunction %void None %func_type
+         %25 = OpLabel
+		 %func = OpVariable %arr_uvec2_2_ptr_func Function
+
+		 ; DeviceToDevice
+		 %ptr_arr_0 = OpAccessChain %arr_uvec2_2_ptr %SSBO_Var %uint_0 %uint_0 %uint_0
+		 %ptr_arr_1 = OpAccessChain %arr_uvec2_2_ptr %SSBO_Var %uint_0 %uint_0 %uint_1
+		 %loaded_array = OpLoad %arr_uvec2_2 %ptr_arr_1
+		 OpStore %ptr_arr_0 %loaded_array
+
+		 ; ConstantToDevice
+		 %ptr_arr_1_const = OpAccessChain %arr_uvec2_2_ptr_const %UBO_Var %uint_0 %uint_0 %uint_1
+		 %loaded_array_const = OpLoad %arr_uvec2_2 %ptr_arr_1_const
+		 OpStore %ptr_arr_0 %loaded_array_const
+
+		 ; StackToDevice
+		 %loaded_array_func = OpLoad %arr_uvec2_2 %func
+		 OpStore %ptr_arr_0 %loaded_array_func
+
+		 ; ThreadGroupToDevice
+		 %loaded_array_workgroup = OpLoad %arr_uvec2_2 %wg
+		 OpStore %ptr_arr_0 %loaded_array_workgroup
+
+		 ; DeviceToThreadGroup
+		 %loaded_array_2 = OpLoad %arr_uvec2_2 %ptr_arr_1
+		 OpStore %wg %loaded_array_2
+
+		 ; DeviceToStack
+		 %loaded_array_3 = OpLoad %arr_uvec2_2 %ptr_arr_1
+		 OpStore %func %loaded_array_3
+
+		 ; ConstantToThreadGroup
+		 %loaded_array_const_2 = OpLoad %arr_uvec2_2 %ptr_arr_1_const
+		 OpStore %wg %loaded_array_const_2
+
+		 ; ConstantToStack
+		 %loaded_array_const_3 = OpLoad %arr_uvec2_2 %ptr_arr_1_const
+		 OpStore %func %loaded_array_const_3
+
+               OpReturn
+               OpFunctionEnd