MSL: Handle array with component when we cannot rely on user() attrib.

In these cases, we emit one variable per location, and so we must
flatten stuff.

14 files changed, 805 insertions, 80 deletions

diff --git a/reference/opt/shaders-msl/frag/array-component-io.frag b/reference/opt/shaders-msl/frag/array-component-io.frag
new file mode 100644
index 00000000..9b4c5b52
--- /dev/null
+++ b/reference/opt/shaders-msl/frag/array-component-io.frag
@@ -0,0 +1,99 @@
+#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 main0_out
+{
+    float4 m_location_0 [[color(0)]];
+    float4 m_location_1 [[color(1)]];
+    float4 m_location_2 [[color(2)]];
+};
+
+struct main0_in
+{
+    float InC_0 [[user(locn0_1), flat]];
+    float InA_0 [[user(locn1), flat]];
+    float InC_1 [[user(locn1_1), flat]];
+    float2 InB_0 [[user(locn1_2), flat]];
+    float InA_1 [[user(locn2), flat]];
+    float InC_2 [[user(locn2_1), flat]];
+    float2 InB_1 [[user(locn2_2), flat]];
+    float InD [[user(locn3_1), sample_perspective]];
+    float InE [[user(locn4_2), center_no_perspective]];
+    float InF [[user(locn5_3), centroid_perspective]];
+};
+
+fragment main0_out main0(main0_in in [[stage_in]])
+{
+    main0_out out = {};
+    spvUnsafeArray<float, 2> A = {};
+    spvUnsafeArray<float2, 2> B = {};
+    spvUnsafeArray<float, 3> C = {};
+    float D = {};
+    spvUnsafeArray<float, 2> InA = {};
+    spvUnsafeArray<float2, 2> InB = {};
+    spvUnsafeArray<float, 3> InC = {};
+    InA[0] = in.InA_0;
+    InA[1] = in.InA_1;
+    InB[0] = in.InB_0;
+    InB[1] = in.InB_1;
+    InC[0] = in.InC_0;
+    InC[1] = in.InC_1;
+    InC[2] = in.InC_2;
+    A = InA;
+    B = InB;
+    C = InC;
+    D = (in.InD + in.InE) + in.InF;
+    out.m_location_1.x = A[0];
+    out.m_location_2.x = A[1];
+    out.m_location_1.zw = B[0];
+    out.m_location_2.zw = B[1];
+    out.m_location_0.y = C[0];
+    out.m_location_1.y = C[1];
+    out.m_location_2.y = C[2];
+    out.m_location_0.w = D;
+    return out;
+}
+
diff --git a/reference/opt/shaders-msl/vert/array-component-io.for-tess.vert b/reference/opt/shaders-msl/vert/array-component-io.for-tess.vert
new file mode 100644
index 00000000..24958eb5
--- /dev/null
+++ b/reference/opt/shaders-msl/vert/array-component-io.for-tess.vert
@@ -0,0 +1,98 @@
+#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 main0_out
+{
+    float4 m_location_0;
+    float4 m_location_1;
+    float4 m_location_2;
+    float4 gl_Position;
+};
+
+struct main0_in
+{
+    float4 m_location_0 [[attribute(0)]];
+    float4 m_location_1 [[attribute(1)]];
+    float4 m_location_2 [[attribute(2)]];
+    float4 Pos [[attribute(4)]];
+};
+
+kernel void main0(main0_in in [[stage_in]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 spvStageInputSize [[grid_size]], device main0_out* spvOut [[buffer(28)]])
+{
+    spvUnsafeArray<float, 2> A = {};
+    spvUnsafeArray<float2, 2> B = {};
+    spvUnsafeArray<float, 3> C = {};
+    float D = {};
+    spvUnsafeArray<float, 2> InA = {};
+    spvUnsafeArray<float2, 2> InB = {};
+    spvUnsafeArray<float, 3> InC = {};
+    float InD = {};
+    device main0_out& out = spvOut[gl_GlobalInvocationID.y * spvStageInputSize.x + gl_GlobalInvocationID.x];
+    InA[0] = in.m_location_1.x;
+    InA[1] = in.m_location_2.x;
+    InB[0] = in.m_location_1.zw;
+    InB[1] = in.m_location_2.zw;
+    InC[0] = in.m_location_0.y;
+    InC[1] = in.m_location_1.y;
+    InC[2] = in.m_location_2.y;
+    InD = in.m_location_0.w;
+    if (any(gl_GlobalInvocationID >= spvStageInputSize))
+        return;
+    out.gl_Position = in.Pos;
+    A = InA;
+    B = InB;
+    C = InC;
+    D = InD;
+    out.m_location_1.x = A[0];
+    out.m_location_2.x = A[1];
+    out.m_location_1.zw = B[0];
+    out.m_location_2.zw = B[1];
+    out.m_location_0.y = C[0];
+    out.m_location_1.y = C[1];
+    out.m_location_2.y = C[2];
+    out.m_location_0.w = D;
+}
+
diff --git a/reference/opt/shaders-msl/vert/array-component-io.vert b/reference/opt/shaders-msl/vert/array-component-io.vert
new file mode 100644
index 00000000..352c9d2e
--- /dev/null
+++ b/reference/opt/shaders-msl/vert/array-component-io.vert
@@ -0,0 +1,100 @@
+#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 main0_out
+{
+    float C_0 [[user(locn0_1)]];
+    float D [[user(locn0_3)]];
+    float A_0 [[user(locn1)]];
+    float C_1 [[user(locn1_1)]];
+    float2 B_0 [[user(locn1_2)]];
+    float A_1 [[user(locn2)]];
+    float C_2 [[user(locn2_1)]];
+    float2 B_1 [[user(locn2_2)]];
+    float4 gl_Position [[position]];
+};
+
+struct main0_in
+{
+    float4 m_location_0 [[attribute(0)]];
+    float4 m_location_1 [[attribute(1)]];
+    float4 m_location_2 [[attribute(2)]];
+    float4 Pos [[attribute(4)]];
+};
+
+vertex main0_out main0(main0_in in [[stage_in]])
+{
+    main0_out out = {};
+    spvUnsafeArray<float, 2> A = {};
+    spvUnsafeArray<float2, 2> B = {};
+    spvUnsafeArray<float, 3> C = {};
+    spvUnsafeArray<float, 2> InA = {};
+    spvUnsafeArray<float2, 2> InB = {};
+    spvUnsafeArray<float, 3> InC = {};
+    float InD = {};
+    InA[0] = in.m_location_1.x;
+    InA[1] = in.m_location_2.x;
+    InB[0] = in.m_location_1.zw;
+    InB[1] = in.m_location_2.zw;
+    InC[0] = in.m_location_0.y;
+    InC[1] = in.m_location_1.y;
+    InC[2] = in.m_location_2.y;
+    InD = in.m_location_0.w;
+    out.gl_Position = in.Pos;
+    A = InA;
+    B = InB;
+    C = InC;
+    out.D = InD;
+    out.A_0 = A[0];
+    out.A_1 = A[1];
+    out.B_0 = B[0];
+    out.B_1 = B[1];
+    out.C_0 = C[0];
+    out.C_1 = C[1];
+    out.C_2 = C[2];
+    return out;
+}
+
diff --git a/reference/shaders-msl-no-opt/components/fragment-output-component.frag b/reference/shaders-msl-no-opt/components/fragment-output-component.frag
index 7e030aa1..45b05b9d 100644
--- a/reference/shaders-msl-no-opt/components/fragment-output-component.frag
+++ b/reference/shaders-msl-no-opt/components/fragment-output-component.frag
@@ -5,7 +5,7 @@ using namespace metal;
 
 struct main0_out
 {
-    float4 FragColor0 [[color(0)]];
+    float4 m_location_0 [[color(0)]];
 };
 
 fragment main0_out main0()
@@ -17,9 +17,9 @@ fragment main0_out main0()
     FragColor0 = 1.0;
     FragColor1 = float2(2.0, 3.0);
     FragColor3 = 4.0;
-    out.FragColor0.x = FragColor0;
-    out.FragColor0.yz = FragColor1;
-    out.FragColor0.w = FragColor3;
+    out.m_location_0.x = FragColor0;
+    out.m_location_0.yz = FragColor1;
+    out.m_location_0.w = FragColor3;
     return out;
 }
 
diff --git a/reference/shaders-msl-no-opt/components/fragment-output-component.pad-fragment.frag b/reference/shaders-msl-no-opt/components/fragment-output-component.pad-fragment.frag
index 13eb7d57..0e4bee12 100644
--- a/reference/shaders-msl-no-opt/components/fragment-output-component.pad-fragment.frag
+++ b/reference/shaders-msl-no-opt/components/fragment-output-component.pad-fragment.frag
@@ -5,7 +5,7 @@ using namespace metal;
 
 struct main0_out
 {
-    float4 FragColor0 [[color(0)]];
+    float3 m_location_0 [[color(0)]];
 };
 
 fragment main0_out main0()
@@ -15,8 +15,8 @@ fragment main0_out main0()
     float2 FragColor1 = {};
     FragColor0 = 1.0;
     FragColor1 = float2(2.0, 3.0);
-    out.FragColor0.x = FragColor0;
-    out.FragColor0.yz = FragColor1;
+    out.m_location_0.x = FragColor0;
+    out.m_location_0.yz = FragColor1;
     return out;
 }
 
diff --git a/reference/shaders-msl-no-opt/components/vertex-input-component.vert b/reference/shaders-msl-no-opt/components/vertex-input-component.vert
index 1aae280d..7a099f50 100644
--- a/reference/shaders-msl-no-opt/components/vertex-input-component.vert
+++ b/reference/shaders-msl-no-opt/components/vertex-input-component.vert
@@ -11,7 +11,7 @@ struct main0_out
 
 struct main0_in
 {
-    float4 Foo3 [[attribute(0)]];
+    float4 m_location_0 [[attribute(0)]];
 };
 
 vertex main0_out main0(main0_in in [[stage_in]])
@@ -19,8 +19,8 @@ vertex main0_out main0(main0_in in [[stage_in]])
     main0_out out = {};
     float3 Foo3 = {};
     float Foo1 = {};
-    Foo3 = in.Foo3.xyz;
-    Foo1 = in.Foo3.w;
+    Foo3 = in.m_location_0.xyz;
+    Foo1 = in.m_location_0.w;
     out.gl_Position = float4(Foo3, Foo1);
     out.Foo = Foo3 + float3(Foo1);
     return out;
diff --git a/reference/shaders-msl/frag/array-component-io.frag b/reference/shaders-msl/frag/array-component-io.frag
new file mode 100644
index 00000000..9b4c5b52
--- /dev/null
+++ b/reference/shaders-msl/frag/array-component-io.frag
@@ -0,0 +1,99 @@
+#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 main0_out
+{
+    float4 m_location_0 [[color(0)]];
+    float4 m_location_1 [[color(1)]];
+    float4 m_location_2 [[color(2)]];
+};
+
+struct main0_in
+{
+    float InC_0 [[user(locn0_1), flat]];
+    float InA_0 [[user(locn1), flat]];
+    float InC_1 [[user(locn1_1), flat]];
+    float2 InB_0 [[user(locn1_2), flat]];
+    float InA_1 [[user(locn2), flat]];
+    float InC_2 [[user(locn2_1), flat]];
+    float2 InB_1 [[user(locn2_2), flat]];
+    float InD [[user(locn3_1), sample_perspective]];
+    float InE [[user(locn4_2), center_no_perspective]];
+    float InF [[user(locn5_3), centroid_perspective]];
+};
+
+fragment main0_out main0(main0_in in [[stage_in]])
+{
+    main0_out out = {};
+    spvUnsafeArray<float, 2> A = {};
+    spvUnsafeArray<float2, 2> B = {};
+    spvUnsafeArray<float, 3> C = {};
+    float D = {};
+    spvUnsafeArray<float, 2> InA = {};
+    spvUnsafeArray<float2, 2> InB = {};
+    spvUnsafeArray<float, 3> InC = {};
+    InA[0] = in.InA_0;
+    InA[1] = in.InA_1;
+    InB[0] = in.InB_0;
+    InB[1] = in.InB_1;
+    InC[0] = in.InC_0;
+    InC[1] = in.InC_1;
+    InC[2] = in.InC_2;
+    A = InA;
+    B = InB;
+    C = InC;
+    D = (in.InD + in.InE) + in.InF;
+    out.m_location_1.x = A[0];
+    out.m_location_2.x = A[1];
+    out.m_location_1.zw = B[0];
+    out.m_location_2.zw = B[1];
+    out.m_location_0.y = C[0];
+    out.m_location_1.y = C[1];
+    out.m_location_2.y = C[2];
+    out.m_location_0.w = D;
+    return out;
+}
+
diff --git a/reference/shaders-msl/vert/array-component-io.for-tess.vert b/reference/shaders-msl/vert/array-component-io.for-tess.vert
new file mode 100644
index 00000000..24958eb5
--- /dev/null
+++ b/reference/shaders-msl/vert/array-component-io.for-tess.vert
@@ -0,0 +1,98 @@
+#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 main0_out
+{
+    float4 m_location_0;
+    float4 m_location_1;
+    float4 m_location_2;
+    float4 gl_Position;
+};
+
+struct main0_in
+{
+    float4 m_location_0 [[attribute(0)]];
+    float4 m_location_1 [[attribute(1)]];
+    float4 m_location_2 [[attribute(2)]];
+    float4 Pos [[attribute(4)]];
+};
+
+kernel void main0(main0_in in [[stage_in]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 spvStageInputSize [[grid_size]], device main0_out* spvOut [[buffer(28)]])
+{
+    spvUnsafeArray<float, 2> A = {};
+    spvUnsafeArray<float2, 2> B = {};
+    spvUnsafeArray<float, 3> C = {};
+    float D = {};
+    spvUnsafeArray<float, 2> InA = {};
+    spvUnsafeArray<float2, 2> InB = {};
+    spvUnsafeArray<float, 3> InC = {};
+    float InD = {};
+    device main0_out& out = spvOut[gl_GlobalInvocationID.y * spvStageInputSize.x + gl_GlobalInvocationID.x];
+    InA[0] = in.m_location_1.x;
+    InA[1] = in.m_location_2.x;
+    InB[0] = in.m_location_1.zw;
+    InB[1] = in.m_location_2.zw;
+    InC[0] = in.m_location_0.y;
+    InC[1] = in.m_location_1.y;
+    InC[2] = in.m_location_2.y;
+    InD = in.m_location_0.w;
+    if (any(gl_GlobalInvocationID >= spvStageInputSize))
+        return;
+    out.gl_Position = in.Pos;
+    A = InA;
+    B = InB;
+    C = InC;
+    D = InD;
+    out.m_location_1.x = A[0];
+    out.m_location_2.x = A[1];
+    out.m_location_1.zw = B[0];
+    out.m_location_2.zw = B[1];
+    out.m_location_0.y = C[0];
+    out.m_location_1.y = C[1];
+    out.m_location_2.y = C[2];
+    out.m_location_0.w = D;
+}
+
diff --git a/reference/shaders-msl/vert/array-component-io.vert b/reference/shaders-msl/vert/array-component-io.vert
new file mode 100644
index 00000000..352c9d2e
--- /dev/null
+++ b/reference/shaders-msl/vert/array-component-io.vert
@@ -0,0 +1,100 @@
+#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 main0_out
+{
+    float C_0 [[user(locn0_1)]];
+    float D [[user(locn0_3)]];
+    float A_0 [[user(locn1)]];
+    float C_1 [[user(locn1_1)]];
+    float2 B_0 [[user(locn1_2)]];
+    float A_1 [[user(locn2)]];
+    float C_2 [[user(locn2_1)]];
+    float2 B_1 [[user(locn2_2)]];
+    float4 gl_Position [[position]];
+};
+
+struct main0_in
+{
+    float4 m_location_0 [[attribute(0)]];
+    float4 m_location_1 [[attribute(1)]];
+    float4 m_location_2 [[attribute(2)]];
+    float4 Pos [[attribute(4)]];
+};
+
+vertex main0_out main0(main0_in in [[stage_in]])
+{
+    main0_out out = {};
+    spvUnsafeArray<float, 2> A = {};
+    spvUnsafeArray<float2, 2> B = {};
+    spvUnsafeArray<float, 3> C = {};
+    spvUnsafeArray<float, 2> InA = {};
+    spvUnsafeArray<float2, 2> InB = {};
+    spvUnsafeArray<float, 3> InC = {};
+    float InD = {};
+    InA[0] = in.m_location_1.x;
+    InA[1] = in.m_location_2.x;
+    InB[0] = in.m_location_1.zw;
+    InB[1] = in.m_location_2.zw;
+    InC[0] = in.m_location_0.y;
+    InC[1] = in.m_location_1.y;
+    InC[2] = in.m_location_2.y;
+    InD = in.m_location_0.w;
+    out.gl_Position = in.Pos;
+    A = InA;
+    B = InB;
+    C = InC;
+    out.D = InD;
+    out.A_0 = A[0];
+    out.A_1 = A[1];
+    out.B_0 = B[0];
+    out.B_1 = B[1];
+    out.C_0 = C[0];
+    out.C_1 = C[1];
+    out.C_2 = C[2];
+    return out;
+}
+
diff --git a/shaders-msl/frag/array-component-io.frag b/shaders-msl/frag/array-component-io.frag
new file mode 100644
index 00000000..8d88249e
--- /dev/null
+++ b/shaders-msl/frag/array-component-io.frag
@@ -0,0 +1,21 @@
+#version 450
+
+layout(location = 1, component = 0) out float A[2];
+layout(location = 1, component = 2) out vec2 B[2];
+layout(location = 0, component = 1) out float C[3];
+layout(location = 0, component = 3) out float D;
+
+layout(location = 1, component = 0) flat in float InA[2];
+layout(location = 1, component = 2) flat in vec2 InB[2];
+layout(location = 0, component = 1) flat in float InC[3];
+layout(location = 3, component = 1) sample in float InD;
+layout(location = 4, component = 2) noperspective in float InE;
+layout(location = 5, component = 3) centroid in float InF;
+
+void main()
+{
+	A = InA;
+	B = InB;
+	C = InC;
+	D = InD + InE + InF;
+}
diff --git a/shaders-msl/vert/array-component-io.for-tess.vert b/shaders-msl/vert/array-component-io.for-tess.vert
new file mode 100644
index 00000000..257ac848
--- /dev/null
+++ b/shaders-msl/vert/array-component-io.for-tess.vert
@@ -0,0 +1,21 @@
+#version 450
+
+layout(location = 1, component = 0) out float A[2];
+layout(location = 1, component = 2) out vec2 B[2];
+layout(location = 0, component = 1) out float C[3];
+layout(location = 0, component = 3) out float D;
+
+layout(location = 1, component = 0) in float InA[2];
+layout(location = 1, component = 2) in vec2 InB[2];
+layout(location = 0, component = 1) in float InC[3];
+layout(location = 0, component = 3) in float InD;
+layout(location = 4) in vec4 Pos;
+
+void main()
+{
+	gl_Position = Pos;
+	A = InA;
+	B = InB;
+	C = InC;
+	D = InD;
+}
diff --git a/shaders-msl/vert/array-component-io.vert b/shaders-msl/vert/array-component-io.vert
new file mode 100644
index 00000000..257ac848
--- /dev/null
+++ b/shaders-msl/vert/array-component-io.vert
@@ -0,0 +1,21 @@
+#version 450
+
+layout(location = 1, component = 0) out float A[2];
+layout(location = 1, component = 2) out vec2 B[2];
+layout(location = 0, component = 1) out float C[3];
+layout(location = 0, component = 3) out float D;
+
+layout(location = 1, component = 0) in float InA[2];
+layout(location = 1, component = 2) in vec2 InB[2];
+layout(location = 0, component = 1) in float InC[3];
+layout(location = 0, component = 3) in float InD;
+layout(location = 4) in vec4 Pos;
+
+void main()
+{
+	gl_Position = Pos;
+	A = InA;
+	B = InB;
+	C = InC;
+	D = InD;
+}
diff --git a/spirv_msl.cpp b/spirv_msl.cpp
index e991b85a..8a5a22f8 100644
--- a/spirv_msl.cpp
+++ b/spirv_msl.cpp
@@ -1991,6 +1991,92 @@ uint32_t CompilerMSL::build_msl_interpolant_type(uint32_t type_id, bool is_noper
 	return new_type_id;
 }
 
+bool CompilerMSL::add_component_variable_to_interface_block(spv::StorageClass storage, const std::string &ib_var_ref,
+                                                            SPIRVariable &var,
+                                                            const SPIRType &type,
+                                                            InterfaceBlockMeta &meta)
+{
+	// Deal with Component decorations.
+	const InterfaceBlockMeta::LocationMeta *location_meta = nullptr;
+	uint32_t location = ~0u;
+	if (has_decoration(var.self, DecorationLocation))
+	{
+		location = get_decoration(var.self, DecorationLocation);
+		auto location_meta_itr = meta.location_meta.find(location);
+		if (location_meta_itr != end(meta.location_meta))
+			location_meta = &location_meta_itr->second;
+	}
+
+	// Check if we need to pad fragment output to match a certain number of components.
+	if (location_meta)
+	{
+		bool pad_fragment_output = has_decoration(var.self, DecorationLocation) &&
+		                           msl_options.pad_fragment_output_components &&
+		                           get_entry_point().model == ExecutionModelFragment && storage == StorageClassOutput;
+
+		auto &entry_func = get<SPIRFunction>(ir.default_entry_point);
+		uint32_t start_component = get_decoration(var.self, DecorationComponent);
+		uint32_t type_components = type.vecsize;
+		uint32_t num_components = location_meta->num_components;
+
+		if (pad_fragment_output)
+		{
+			uint32_t locn = get_decoration(var.self, DecorationLocation);
+			num_components = std::max(num_components, get_target_components_for_fragment_location(locn));
+		}
+
+		// We have already declared an IO block member as m_location_N.
+		// Just emit an early-declared variable and fixup as needed.
+		// Arrays need to be unrolled here since each location might need a different number of components.
+		entry_func.add_local_variable(var.self);
+		vars_needing_early_declaration.push_back(var.self);
+
+		if (var.storage == StorageClassInput)
+		{
+			entry_func.fixup_hooks_in.push_back([=, &type, &var]() {
+				if (!type.array.empty())
+				{
+					uint32_t array_size = to_array_size_literal(type);
+					for (uint32_t loc_off = 0; loc_off < array_size; loc_off++)
+					{
+						statement(to_name(var.self), "[", loc_off, "]", " = ", ib_var_ref,
+						          ".m_location_", location + loc_off,
+						          vector_swizzle(type_components, start_component), ";");
+					}
+				}
+				else
+				{
+					statement(to_name(var.self), " = ", ib_var_ref, ".m_location_", location,
+					          vector_swizzle(type_components, start_component), ";");
+				}
+			});
+		}
+		else
+		{
+			entry_func.fixup_hooks_out.push_back([=, &type, &var]() {
+				if (!type.array.empty())
+				{
+					uint32_t array_size = to_array_size_literal(type);
+					for (uint32_t loc_off = 0; loc_off < array_size; loc_off++)
+					{
+						statement(ib_var_ref, ".m_location_", location + loc_off,
+						          vector_swizzle(type_components, start_component), " = ",
+						          to_name(var.self), "[", loc_off, "];");
+					}
+				}
+				else
+				{
+					statement(ib_var_ref, ".m_location_", location,
+					          vector_swizzle(type_components, start_component), " = ", to_name(var.self), ";");
+				}
+			});
+		}
+		return true;
+	}
+	else
+		return false;
+}
+
 void CompilerMSL::add_plain_variable_to_interface_block(StorageClass storage, const string &ib_var_ref,
                                                         SPIRType &ib_type, SPIRVariable &var, InterfaceBlockMeta &meta)
 {
@@ -2019,65 +2105,14 @@ void CompilerMSL::add_plain_variable_to_interface_block(StorageClass storage, co
 
 	auto &entry_func = get<SPIRFunction>(ir.default_entry_point);
 
-	// Deal with Component decorations.
-	InterfaceBlockMeta::LocationMeta *location_meta = nullptr;
-	if (has_decoration(var.self, DecorationLocation))
-	{
-		auto location_meta_itr = meta.location_meta.find(get_decoration(var.self, DecorationLocation));
-		if (location_meta_itr != end(meta.location_meta))
-			location_meta = &location_meta_itr->second;
-	}
+	if (add_component_variable_to_interface_block(storage, ib_var_ref, var, type, meta))
+		return;
 
 	bool pad_fragment_output = has_decoration(var.self, DecorationLocation) &&
 	                           msl_options.pad_fragment_output_components &&
 	                           get_entry_point().model == ExecutionModelFragment && storage == StorageClassOutput;
 
-	// Check if we need to pad fragment output to match a certain number of components.
-	if (location_meta)
-	{
-		start_component = get_decoration(var.self, DecorationComponent);
-		uint32_t num_components = location_meta->num_components;
-		if (pad_fragment_output)
-		{
-			uint32_t locn = get_decoration(var.self, DecorationLocation);
-			num_components = std::max(num_components, get_target_components_for_fragment_location(locn));
-		}
-
-		if (location_meta->ib_index != ~0u)
-		{
-			// We have already declared the variable. Just emit an early-declared variable and fixup as needed.
-			entry_func.add_local_variable(var.self);
-			vars_needing_early_declaration.push_back(var.self);
-
-			if (var.storage == StorageClassInput)
-			{
-				uint32_t ib_index = location_meta->ib_index;
-				entry_func.fixup_hooks_in.push_back([=, &var]() {
-					statement(to_name(var.self), " = ", ib_var_ref, ".", to_member_name(ib_type, ib_index),
-					          vector_swizzle(type_components, start_component), ";");
-				});
-			}
-			else
-			{
-				uint32_t ib_index = location_meta->ib_index;
-				entry_func.fixup_hooks_out.push_back([=, &var]() {
-					statement(ib_var_ref, ".", to_member_name(ib_type, ib_index),
-					          vector_swizzle(type_components, start_component), " = ", to_name(var.self), ";");
-				});
-			}
-			return;
-		}
-		else
-		{
-			location_meta->ib_index = uint32_t(ib_type.member_types.size());
-			type_id = build_extended_vector_type(type_id, num_components);
-			if (var.storage == StorageClassInput)
-				padded_input = true;
-			else
-				padded_output = true;
-		}
-	}
-	else if (pad_fragment_output)
+	if (pad_fragment_output)
 	{
 		uint32_t locn = get_decoration(var.self, DecorationLocation);
 		target_components = get_target_components_for_fragment_location(locn);
@@ -2169,11 +2204,8 @@ void CompilerMSL::add_plain_variable_to_interface_block(StorageClass storage, co
 		uint32_t locn = get_decoration(var.self, DecorationLocation);
 		if (storage == StorageClassInput)
 		{
-			type_id = ensure_correct_input_type(var.basetype, locn,
-			                                    location_meta ? location_meta->num_components : 0,
-			                                    meta.strip_array);
-			if (!location_meta)
-				var.basetype = type_id;
+			type_id = ensure_correct_input_type(var.basetype, locn, 0, meta.strip_array);
+			var.basetype = type_id;
 
 			type_id = get_pointee_type_id(type_id);
 			if (meta.strip_array && is_array(get<SPIRType>(type_id)))
@@ -2193,13 +2225,10 @@ void CompilerMSL::add_plain_variable_to_interface_block(StorageClass storage, co
 		mark_location_as_used_by_shader(locn, type, storage);
 	}
 
-	if (!location_meta)
+	if (get_decoration_bitset(var.self).get(DecorationComponent))
 	{
-		if (get_decoration_bitset(var.self).get(DecorationComponent))
-		{
-			uint32_t component = get_decoration(var.self, DecorationComponent);
-			set_member_decoration(ib_type.self, ib_mbr_idx, DecorationComponent, component);
-		}
+		uint32_t component = get_decoration(var.self, DecorationComponent);
+		set_member_decoration(ib_type.self, ib_mbr_idx, DecorationComponent, component);
 	}
 
 	if (get_decoration_bitset(var.self).get(DecorationIndex))
@@ -2229,10 +2258,7 @@ void CompilerMSL::add_plain_variable_to_interface_block(StorageClass storage, co
 			set_member_decoration(ib_type.self, ib_mbr_idx, DecorationSample);
 	}
 
-	// If we have location meta, there is no unique OrigID. We won't need it, since we flatten/unflatten
-	// the variable to stack anyways here.
-	if (!location_meta)
-		set_extended_member_decoration(ib_type.self, ib_mbr_idx, SPIRVCrossDecorationInterfaceOrigID, var.self);
+	set_extended_member_decoration(ib_type.self, ib_mbr_idx, SPIRVCrossDecorationInterfaceOrigID, var.self);
 }
 
 void CompilerMSL::add_composite_variable_to_interface_block(StorageClass storage, const string &ib_var_ref,
@@ -2243,6 +2269,9 @@ void CompilerMSL::add_composite_variable_to_interface_block(StorageClass storage
 	auto &var_type = meta.strip_array ? get_variable_element_type(var) : get_variable_data_type(var);
 	uint32_t elem_cnt = 0;
 
+	if (add_component_variable_to_interface_block(storage, ib_var_ref, var, var_type, meta))
+		return;
+
 	if (is_matrix(var_type))
 	{
 		if (is_array(var_type))
@@ -3339,6 +3368,13 @@ uint32_t CompilerMSL::add_interface_block(StorageClass storage, bool patch)
 						{
 							auto &location_meta = meta.location_meta[location + location_offset];
 							location_meta.num_components = std::max(location_meta.num_components, component + type.vecsize);
+
+							// For variables sharing location, decorations and base type must match.
+							location_meta.base_type_id = type.self;
+							location_meta.flat = has_decoration(var.self, DecorationFlat);
+							location_meta.noperspective = has_decoration(var.self, DecorationNoPerspective);
+							location_meta.centroid = has_decoration(var.self, DecorationCentroid);
+							location_meta.sample = has_decoration(var.self, DecorationSample);
 						}
 					}
 				}
@@ -3588,6 +3624,31 @@ uint32_t CompilerMSL::add_interface_block(StorageClass storage, bool patch)
 		}
 	}
 
+	// When multiple variables need to access same location,
+	// unroll locations one by one and we will flatten output or input as necessary.
+	for (auto &loc : meta.location_meta)
+	{
+		uint32_t location = loc.first;
+		auto &location_meta = loc.second;
+
+		uint32_t ib_mbr_idx = uint32_t(ib_type.member_types.size());
+		uint32_t type_id = build_extended_vector_type(location_meta.base_type_id, location_meta.num_components);
+		ib_type.member_types.push_back(type_id);
+
+		set_member_name(ib_type.self, ib_mbr_idx, join("m_location_", location));
+		set_member_decoration(ib_type.self, ib_mbr_idx, DecorationLocation, location);
+		mark_location_as_used_by_shader(location, get<SPIRType>(type_id), storage);
+
+		if (location_meta.flat)
+			set_member_decoration(ib_type.self, ib_mbr_idx, DecorationFlat);
+		if (location_meta.noperspective)
+			set_member_decoration(ib_type.self, ib_mbr_idx, DecorationNoPerspective);
+		if (location_meta.centroid)
+			set_member_decoration(ib_type.self, ib_mbr_idx, DecorationCentroid);
+		if (location_meta.sample)
+			set_member_decoration(ib_type.self, ib_mbr_idx, DecorationSample);
+	}
+
 	// Sort the members of the structure by their locations.
 	MemberSorter member_sorter(ib_type, ir.meta[ib_type_id], MemberSorter::LocationThenBuiltInType);
 	member_sorter.sort();
diff --git a/spirv_msl.hpp b/spirv_msl.hpp
index 95e0c966..61f54be1 100644
--- a/spirv_msl.hpp
+++ b/spirv_msl.hpp
@@ -799,8 +799,12 @@ protected:
 	{
 		struct LocationMeta
 		{
+			uint32_t base_type_id = 0;
 			uint32_t num_components = 0;
-			uint32_t ib_index = ~0u;
+			bool flat = false;
+			bool noperspective = false;
+			bool centroid = false;
+			bool sample = false;
 		};
 		std::unordered_map<uint32_t, LocationMeta> location_meta;
 		bool strip_array = false;
@@ -815,6 +819,9 @@ protected:
 	                                               SPIRType &ib_type, SPIRVariable &var, InterfaceBlockMeta &meta);
 	void add_plain_variable_to_interface_block(spv::StorageClass storage, const std::string &ib_var_ref,
 	                                           SPIRType &ib_type, SPIRVariable &var, InterfaceBlockMeta &meta);
+	bool add_component_variable_to_interface_block(spv::StorageClass storage, const std::string &ib_var_ref,
+	                                               SPIRVariable &var, const SPIRType &type,
+	                                               InterfaceBlockMeta &meta);
 	void add_plain_member_variable_to_interface_block(spv::StorageClass storage, const std::string &ib_var_ref,
 	                                                  SPIRType &ib_type, SPIRVariable &var, uint32_t index,
 	                                                  InterfaceBlockMeta &meta);