MSL: Handle descriptor aliasing of raw buffer descriptors.

It is allowed to redeclare descriptors with different types in Vulkan.
MSL in general does not allow this, but for raw buffers, we can cast the
reference type at the very least.

For typed resources we are kinda hosed. Without descriptor indexing's
PARTIALLY_BOUND_BIT, descriptors must be valid if they are statically
accessed, so it would not be valid to access differently typed aliases
unless that flag is used. There might be a way to reinterpret cast
descriptors, but that seems very sketchy.

Implements support for:

- Single discrete descriptor
- Single argument buffer descriptor
- Array of argument buffer descriptors

Other cases are unimplemented for now since they are extremely painful
to unroll.

9 files changed, 878 insertions, 23 deletions

diff --git a/reference/opt/shaders-msl/comp/raw-buffer-descriptor-aliasing.argument.discrete.device-argument-buffer.msl2.comp b/reference/opt/shaders-msl/comp/raw-buffer-descriptor-aliasing.argument.discrete.device-argument-buffer.msl2.comp
new file mode 100644
index 00000000..a3c1a5b3
--- /dev/null
+++ b/reference/opt/shaders-msl/comp/raw-buffer-descriptor-aliasing.argument.discrete.device-argument-buffer.msl2.comp
@@ -0,0 +1,111 @@
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+struct SSBO_A
+{
+    float data[1];
+};
+
+struct UBO_C
+{
+    float4 data[1024];
+};
+
+struct SSBO_B
+{
+    uint2 data[1];
+};
+
+struct UBO_D
+{
+    uint4 data[1024];
+};
+
+struct SSBO_BRO
+{
+    uint2 data[1];
+};
+
+struct SSBO_As
+{
+    float data[1];
+};
+
+struct UBO_Cs
+{
+    float4 data[1024];
+};
+
+struct SSBO_Bs
+{
+    uint2 data[1024];
+};
+
+struct UBO_Ds
+{
+    uint4 data[1024];
+};
+
+struct SSBO_BsRO
+{
+    uint2 data[1024];
+};
+
+struct SSBO_E
+{
+    float data[1];
+};
+
+struct UBO_G
+{
+    float4 data[1024];
+};
+
+struct SSBO_F
+{
+    uint2 data[1];
+};
+
+struct UBO_H
+{
+    uint4 data[1024];
+};
+
+struct SSBO_I
+{
+    uint2 data[1];
+};
+
+constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(64u, 1u, 1u);
+
+struct spvDescriptorSetBuffer0
+{
+    device SSBO_A* ssbo_a [[id(0)]];
+    constant UBO_C* ubo_c [[id(1)]];
+    device SSBO_As* ssbo_as [[id(2)]][4];
+    constant UBO_Cs* ubo_cs [[id(6)]][4];
+};
+
+kernel void main0(const device spvDescriptorSetBuffer0& spvDescriptorSet0 [[buffer(0)]], device void* spvBufferAliasSet2Binding0 [[buffer(1)]], constant void* spvBufferAliasSet2Binding1 [[buffer(2)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]])
+{
+    device auto& ssbo_e = *(device SSBO_E*)spvBufferAliasSet2Binding0;
+    constant auto& ubo_g = *(constant UBO_G*)spvBufferAliasSet2Binding1;
+    device auto& ssbo_f = *(device SSBO_F*)spvBufferAliasSet2Binding0;
+    constant auto& ubo_h = *(constant UBO_H*)spvBufferAliasSet2Binding1;
+    const device auto& ssbo_i = *(const device SSBO_I*)spvBufferAliasSet2Binding0;
+    device auto& ssbo_b = (device SSBO_B&)(*spvDescriptorSet0.ssbo_a);
+    constant auto& ubo_d = (constant UBO_D&)(*spvDescriptorSet0.ubo_c);
+    const device auto& ssbo_b_readonly = (const device SSBO_BRO&)(*spvDescriptorSet0.ssbo_a);
+    const device auto& ssbo_bs = (device SSBO_Bs* const device (&)[4])spvDescriptorSet0.ssbo_as;
+    const device auto& ubo_ds = (constant UBO_Ds* const device (&)[4])spvDescriptorSet0.ubo_cs;
+    const device auto& ssbo_bs_readonly = (const device SSBO_BsRO* const device (&)[4])spvDescriptorSet0.ssbo_as;
+    (*spvDescriptorSet0.ssbo_a).data[gl_GlobalInvocationID.x] = (*spvDescriptorSet0.ubo_c).data[gl_WorkGroupID.x].x;
+    ssbo_b.data[gl_GlobalInvocationID.x] = ubo_d.data[gl_WorkGroupID.y].xy + ssbo_b_readonly.data[gl_GlobalInvocationID.x];
+    spvDescriptorSet0.ssbo_as[gl_WorkGroupID.x]->data[gl_GlobalInvocationID.x] = spvDescriptorSet0.ubo_cs[gl_WorkGroupID.x]->data[0].x;
+    ssbo_bs[gl_WorkGroupID.x]->data[gl_GlobalInvocationID.x] = ubo_ds[gl_WorkGroupID.x]->data[0].xy + ssbo_bs_readonly[gl_WorkGroupID.x]->data[gl_GlobalInvocationID.x];
+    ssbo_e.data[gl_GlobalInvocationID.x] = ubo_g.data[gl_WorkGroupID.x].x;
+    ssbo_f.data[gl_GlobalInvocationID.x] = ubo_h.data[gl_WorkGroupID.y].xy + ssbo_i.data[gl_GlobalInvocationID.x];
+}
+
diff --git a/reference/opt/shaders-msl/comp/raw-buffer-descriptor-aliasing.argument.discrete.msl2.comp b/reference/opt/shaders-msl/comp/raw-buffer-descriptor-aliasing.argument.discrete.msl2.comp
new file mode 100644
index 00000000..bc0aa461
--- /dev/null
+++ b/reference/opt/shaders-msl/comp/raw-buffer-descriptor-aliasing.argument.discrete.msl2.comp
@@ -0,0 +1,111 @@
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+struct SSBO_A
+{
+    float data[1];
+};
+
+struct UBO_C
+{
+    float4 data[1024];
+};
+
+struct SSBO_B
+{
+    uint2 data[1];
+};
+
+struct UBO_D
+{
+    uint4 data[1024];
+};
+
+struct SSBO_BRO
+{
+    uint2 data[1];
+};
+
+struct SSBO_As
+{
+    float data[1];
+};
+
+struct UBO_Cs
+{
+    float4 data[1024];
+};
+
+struct SSBO_Bs
+{
+    uint2 data[1024];
+};
+
+struct UBO_Ds
+{
+    uint4 data[1024];
+};
+
+struct SSBO_BsRO
+{
+    uint2 data[1024];
+};
+
+struct SSBO_E
+{
+    float data[1];
+};
+
+struct UBO_G
+{
+    float4 data[1024];
+};
+
+struct SSBO_F
+{
+    uint2 data[1];
+};
+
+struct UBO_H
+{
+    uint4 data[1024];
+};
+
+struct SSBO_I
+{
+    uint2 data[1];
+};
+
+constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(64u, 1u, 1u);
+
+struct spvDescriptorSetBuffer0
+{
+    device SSBO_A* ssbo_a [[id(0)]];
+    constant UBO_C* ubo_c [[id(1)]];
+    device SSBO_As* ssbo_as [[id(2)]][4];
+    constant UBO_Cs* ubo_cs [[id(6)]][4];
+};
+
+kernel void main0(constant spvDescriptorSetBuffer0& spvDescriptorSet0 [[buffer(0)]], device void* spvBufferAliasSet2Binding0 [[buffer(1)]], constant void* spvBufferAliasSet2Binding1 [[buffer(2)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]])
+{
+    device auto& ssbo_e = *(device SSBO_E*)spvBufferAliasSet2Binding0;
+    constant auto& ubo_g = *(constant UBO_G*)spvBufferAliasSet2Binding1;
+    device auto& ssbo_f = *(device SSBO_F*)spvBufferAliasSet2Binding0;
+    constant auto& ubo_h = *(constant UBO_H*)spvBufferAliasSet2Binding1;
+    const device auto& ssbo_i = *(const device SSBO_I*)spvBufferAliasSet2Binding0;
+    device auto& ssbo_b = (device SSBO_B&)(*spvDescriptorSet0.ssbo_a);
+    constant auto& ubo_d = (constant UBO_D&)(*spvDescriptorSet0.ubo_c);
+    const device auto& ssbo_b_readonly = (const device SSBO_BRO&)(*spvDescriptorSet0.ssbo_a);
+    constant auto& ssbo_bs = (device SSBO_Bs* constant (&)[4])spvDescriptorSet0.ssbo_as;
+    constant auto& ubo_ds = (constant UBO_Ds* constant (&)[4])spvDescriptorSet0.ubo_cs;
+    constant auto& ssbo_bs_readonly = (const device SSBO_BsRO* constant (&)[4])spvDescriptorSet0.ssbo_as;
+    (*spvDescriptorSet0.ssbo_a).data[gl_GlobalInvocationID.x] = (*spvDescriptorSet0.ubo_c).data[gl_WorkGroupID.x].x;
+    ssbo_b.data[gl_GlobalInvocationID.x] = ubo_d.data[gl_WorkGroupID.y].xy + ssbo_b_readonly.data[gl_GlobalInvocationID.x];
+    spvDescriptorSet0.ssbo_as[gl_WorkGroupID.x]->data[gl_GlobalInvocationID.x] = spvDescriptorSet0.ubo_cs[gl_WorkGroupID.x]->data[0].x;
+    ssbo_bs[gl_WorkGroupID.x]->data[gl_GlobalInvocationID.x] = ubo_ds[gl_WorkGroupID.x]->data[0].xy + ssbo_bs_readonly[gl_WorkGroupID.x]->data[gl_GlobalInvocationID.x];
+    ssbo_e.data[gl_GlobalInvocationID.x] = ubo_g.data[gl_WorkGroupID.x].x;
+    ssbo_f.data[gl_GlobalInvocationID.x] = ubo_h.data[gl_WorkGroupID.y].xy + ssbo_i.data[gl_GlobalInvocationID.x];
+}
+
diff --git a/reference/shaders-msl/comp/raw-buffer-descriptor-aliasing.argument.discrete.device-argument-buffer.msl2.comp b/reference/shaders-msl/comp/raw-buffer-descriptor-aliasing.argument.discrete.device-argument-buffer.msl2.comp
new file mode 100644
index 00000000..c5dc95e8
--- /dev/null
+++ b/reference/shaders-msl/comp/raw-buffer-descriptor-aliasing.argument.discrete.device-argument-buffer.msl2.comp
@@ -0,0 +1,137 @@
+#pragma clang diagnostic ignored "-Wmissing-prototypes"
+
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+struct SSBO_A
+{
+    float data[1];
+};
+
+struct UBO_C
+{
+    float4 data[1024];
+};
+
+struct SSBO_B
+{
+    uint2 data[1];
+};
+
+struct UBO_D
+{
+    uint4 data[1024];
+};
+
+struct SSBO_BRO
+{
+    uint2 data[1];
+};
+
+struct SSBO_As
+{
+    float data[1];
+};
+
+struct UBO_Cs
+{
+    float4 data[1024];
+};
+
+struct SSBO_Bs
+{
+    uint2 data[1024];
+};
+
+struct UBO_Ds
+{
+    uint4 data[1024];
+};
+
+struct SSBO_BsRO
+{
+    uint2 data[1024];
+};
+
+struct SSBO_E
+{
+    float data[1];
+};
+
+struct UBO_G
+{
+    float4 data[1024];
+};
+
+struct SSBO_F
+{
+    uint2 data[1];
+};
+
+struct UBO_H
+{
+    uint4 data[1024];
+};
+
+struct SSBO_I
+{
+    uint2 data[1];
+};
+
+constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(64u, 1u, 1u);
+
+struct spvDescriptorSetBuffer0
+{
+    device SSBO_A* ssbo_a [[id(0)]];
+    constant UBO_C* ubo_c [[id(1)]];
+    device SSBO_As* ssbo_as [[id(2)]][4];
+    constant UBO_Cs* ubo_cs [[id(6)]][4];
+};
+
+static inline __attribute__((always_inline))
+void func0(device SSBO_A& ssbo_a, thread uint3& gl_GlobalInvocationID, constant UBO_C& ubo_c, thread uint3& gl_WorkGroupID, device SSBO_B& ssbo_b, constant UBO_D& ubo_d, const device SSBO_BRO& ssbo_b_readonly)
+{
+    ssbo_a.data[gl_GlobalInvocationID.x] = ubo_c.data[gl_WorkGroupID.x].x;
+    ssbo_b.data[gl_GlobalInvocationID.x] = ubo_d.data[gl_WorkGroupID.y].xy + ssbo_b_readonly.data[gl_GlobalInvocationID.x];
+}
+
+static inline __attribute__((always_inline))
+void func1(thread uint3& gl_GlobalInvocationID, thread uint3& gl_WorkGroupID, device SSBO_As* const device (&ssbo_as)[4], constant UBO_Cs* const device (&ubo_cs)[4])
+{
+    ssbo_as[gl_WorkGroupID.x]->data[gl_GlobalInvocationID.x] = ubo_cs[gl_WorkGroupID.x]->data[0].x;
+}
+
+static inline __attribute__((always_inline))
+void func2(thread uint3& gl_GlobalInvocationID, thread uint3& gl_WorkGroupID, device SSBO_Bs* const device (&ssbo_bs)[4], constant UBO_Ds* const device (&ubo_ds)[4], const device SSBO_BsRO* const device (&ssbo_bs_readonly)[4])
+{
+    ssbo_bs[gl_WorkGroupID.x]->data[gl_GlobalInvocationID.x] = ubo_ds[gl_WorkGroupID.x]->data[0].xy + ssbo_bs_readonly[gl_WorkGroupID.x]->data[gl_GlobalInvocationID.x];
+}
+
+static inline __attribute__((always_inline))
+void func3(thread uint3& gl_GlobalInvocationID, thread uint3& gl_WorkGroupID, device SSBO_E& ssbo_e, constant UBO_G& ubo_g, device SSBO_F& ssbo_f, constant UBO_H& ubo_h, const device SSBO_I& ssbo_i)
+{
+    ssbo_e.data[gl_GlobalInvocationID.x] = ubo_g.data[gl_WorkGroupID.x].x;
+    ssbo_f.data[gl_GlobalInvocationID.x] = ubo_h.data[gl_WorkGroupID.y].xy + ssbo_i.data[gl_GlobalInvocationID.x];
+}
+
+kernel void main0(const device spvDescriptorSetBuffer0& spvDescriptorSet0 [[buffer(0)]], device void* spvBufferAliasSet2Binding0 [[buffer(1)]], constant void* spvBufferAliasSet2Binding1 [[buffer(2)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]])
+{
+    device auto& ssbo_e = *(device SSBO_E*)spvBufferAliasSet2Binding0;
+    constant auto& ubo_g = *(constant UBO_G*)spvBufferAliasSet2Binding1;
+    device auto& ssbo_f = *(device SSBO_F*)spvBufferAliasSet2Binding0;
+    constant auto& ubo_h = *(constant UBO_H*)spvBufferAliasSet2Binding1;
+    const device auto& ssbo_i = *(const device SSBO_I*)spvBufferAliasSet2Binding0;
+    device auto& ssbo_b = (device SSBO_B&)(*spvDescriptorSet0.ssbo_a);
+    constant auto& ubo_d = (constant UBO_D&)(*spvDescriptorSet0.ubo_c);
+    const device auto& ssbo_b_readonly = (const device SSBO_BRO&)(*spvDescriptorSet0.ssbo_a);
+    const device auto& ssbo_bs = (device SSBO_Bs* const device (&)[4])spvDescriptorSet0.ssbo_as;
+    const device auto& ubo_ds = (constant UBO_Ds* const device (&)[4])spvDescriptorSet0.ubo_cs;
+    const device auto& ssbo_bs_readonly = (const device SSBO_BsRO* const device (&)[4])spvDescriptorSet0.ssbo_as;
+    func0((*spvDescriptorSet0.ssbo_a), gl_GlobalInvocationID, (*spvDescriptorSet0.ubo_c), gl_WorkGroupID, ssbo_b, ubo_d, ssbo_b_readonly);
+    func1(gl_GlobalInvocationID, gl_WorkGroupID, spvDescriptorSet0.ssbo_as, spvDescriptorSet0.ubo_cs);
+    func2(gl_GlobalInvocationID, gl_WorkGroupID, ssbo_bs, ubo_ds, ssbo_bs_readonly);
+    func3(gl_GlobalInvocationID, gl_WorkGroupID, ssbo_e, ubo_g, ssbo_f, ubo_h, ssbo_i);
+}
+
diff --git a/reference/shaders-msl/comp/raw-buffer-descriptor-aliasing.argument.discrete.msl2.comp b/reference/shaders-msl/comp/raw-buffer-descriptor-aliasing.argument.discrete.msl2.comp
new file mode 100644
index 00000000..bdc5bc1c
--- /dev/null
+++ b/reference/shaders-msl/comp/raw-buffer-descriptor-aliasing.argument.discrete.msl2.comp
@@ -0,0 +1,137 @@
+#pragma clang diagnostic ignored "-Wmissing-prototypes"
+
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+struct SSBO_A
+{
+    float data[1];
+};
+
+struct UBO_C
+{
+    float4 data[1024];
+};
+
+struct SSBO_B
+{
+    uint2 data[1];
+};
+
+struct UBO_D
+{
+    uint4 data[1024];
+};
+
+struct SSBO_BRO
+{
+    uint2 data[1];
+};
+
+struct SSBO_As
+{
+    float data[1];
+};
+
+struct UBO_Cs
+{
+    float4 data[1024];
+};
+
+struct SSBO_Bs
+{
+    uint2 data[1024];
+};
+
+struct UBO_Ds
+{
+    uint4 data[1024];
+};
+
+struct SSBO_BsRO
+{
+    uint2 data[1024];
+};
+
+struct SSBO_E
+{
+    float data[1];
+};
+
+struct UBO_G
+{
+    float4 data[1024];
+};
+
+struct SSBO_F
+{
+    uint2 data[1];
+};
+
+struct UBO_H
+{
+    uint4 data[1024];
+};
+
+struct SSBO_I
+{
+    uint2 data[1];
+};
+
+constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(64u, 1u, 1u);
+
+struct spvDescriptorSetBuffer0
+{
+    device SSBO_A* ssbo_a [[id(0)]];
+    constant UBO_C* ubo_c [[id(1)]];
+    device SSBO_As* ssbo_as [[id(2)]][4];
+    constant UBO_Cs* ubo_cs [[id(6)]][4];
+};
+
+static inline __attribute__((always_inline))
+void func0(device SSBO_A& ssbo_a, thread uint3& gl_GlobalInvocationID, constant UBO_C& ubo_c, thread uint3& gl_WorkGroupID, device SSBO_B& ssbo_b, constant UBO_D& ubo_d, const device SSBO_BRO& ssbo_b_readonly)
+{
+    ssbo_a.data[gl_GlobalInvocationID.x] = ubo_c.data[gl_WorkGroupID.x].x;
+    ssbo_b.data[gl_GlobalInvocationID.x] = ubo_d.data[gl_WorkGroupID.y].xy + ssbo_b_readonly.data[gl_GlobalInvocationID.x];
+}
+
+static inline __attribute__((always_inline))
+void func1(thread uint3& gl_GlobalInvocationID, thread uint3& gl_WorkGroupID, device SSBO_As* constant (&ssbo_as)[4], constant UBO_Cs* constant (&ubo_cs)[4])
+{
+    ssbo_as[gl_WorkGroupID.x]->data[gl_GlobalInvocationID.x] = ubo_cs[gl_WorkGroupID.x]->data[0].x;
+}
+
+static inline __attribute__((always_inline))
+void func2(thread uint3& gl_GlobalInvocationID, thread uint3& gl_WorkGroupID, device SSBO_Bs* constant (&ssbo_bs)[4], constant UBO_Ds* constant (&ubo_ds)[4], const device SSBO_BsRO* constant (&ssbo_bs_readonly)[4])
+{
+    ssbo_bs[gl_WorkGroupID.x]->data[gl_GlobalInvocationID.x] = ubo_ds[gl_WorkGroupID.x]->data[0].xy + ssbo_bs_readonly[gl_WorkGroupID.x]->data[gl_GlobalInvocationID.x];
+}
+
+static inline __attribute__((always_inline))
+void func3(thread uint3& gl_GlobalInvocationID, thread uint3& gl_WorkGroupID, device SSBO_E& ssbo_e, constant UBO_G& ubo_g, device SSBO_F& ssbo_f, constant UBO_H& ubo_h, const device SSBO_I& ssbo_i)
+{
+    ssbo_e.data[gl_GlobalInvocationID.x] = ubo_g.data[gl_WorkGroupID.x].x;
+    ssbo_f.data[gl_GlobalInvocationID.x] = ubo_h.data[gl_WorkGroupID.y].xy + ssbo_i.data[gl_GlobalInvocationID.x];
+}
+
+kernel void main0(constant spvDescriptorSetBuffer0& spvDescriptorSet0 [[buffer(0)]], device void* spvBufferAliasSet2Binding0 [[buffer(1)]], constant void* spvBufferAliasSet2Binding1 [[buffer(2)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]])
+{
+    device auto& ssbo_e = *(device SSBO_E*)spvBufferAliasSet2Binding0;
+    constant auto& ubo_g = *(constant UBO_G*)spvBufferAliasSet2Binding1;
+    device auto& ssbo_f = *(device SSBO_F*)spvBufferAliasSet2Binding0;
+    constant auto& ubo_h = *(constant UBO_H*)spvBufferAliasSet2Binding1;
+    const device auto& ssbo_i = *(const device SSBO_I*)spvBufferAliasSet2Binding0;
+    device auto& ssbo_b = (device SSBO_B&)(*spvDescriptorSet0.ssbo_a);
+    constant auto& ubo_d = (constant UBO_D&)(*spvDescriptorSet0.ubo_c);
+    const device auto& ssbo_b_readonly = (const device SSBO_BRO&)(*spvDescriptorSet0.ssbo_a);
+    constant auto& ssbo_bs = (device SSBO_Bs* constant (&)[4])spvDescriptorSet0.ssbo_as;
+    constant auto& ubo_ds = (constant UBO_Ds* constant (&)[4])spvDescriptorSet0.ubo_cs;
+    constant auto& ssbo_bs_readonly = (const device SSBO_BsRO* constant (&)[4])spvDescriptorSet0.ssbo_as;
+    func0((*spvDescriptorSet0.ssbo_a), gl_GlobalInvocationID, (*spvDescriptorSet0.ubo_c), gl_WorkGroupID, ssbo_b, ubo_d, ssbo_b_readonly);
+    func1(gl_GlobalInvocationID, gl_WorkGroupID, spvDescriptorSet0.ssbo_as, spvDescriptorSet0.ubo_cs);
+    func2(gl_GlobalInvocationID, gl_WorkGroupID, ssbo_bs, ubo_ds, ssbo_bs_readonly);
+    func3(gl_GlobalInvocationID, gl_WorkGroupID, ssbo_e, ubo_g, ssbo_f, ubo_h, ssbo_i);
+}
+
diff --git a/shaders-msl-no-opt/asm/comp/block-like-array-type-construct-2.asm.comp b/shaders-msl-no-opt/asm/comp/block-like-array-type-construct-2.asm.comp
index aff98231..37ff035f 100644
--- a/shaders-msl-no-opt/asm/comp/block-like-array-type-construct-2.asm.comp
+++ b/shaders-msl-no-opt/asm/comp/block-like-array-type-construct-2.asm.comp
@@ -24,7 +24,7 @@
                OpDecorate %CommonConstants DescriptorSet 0
                OpDecorate %CommonConstants Binding 0
                OpDecorate %g_data DescriptorSet 0
-               OpDecorate %g_data Binding 0
+               OpDecorate %g_data Binding 1
                OpMemberDecorate %type_CommonConstants 0 Offset 0
                OpMemberDecorate %type_CommonConstants 1 Offset 4
                OpDecorate %type_CommonConstants Block
diff --git a/shaders-msl/comp/raw-buffer-descriptor-aliasing.argument.discrete.device-argument-buffer.msl2.comp b/shaders-msl/comp/raw-buffer-descriptor-aliasing.argument.discrete.device-argument-buffer.msl2.comp
new file mode 100644
index 00000000..25ec7840
--- /dev/null
+++ b/shaders-msl/comp/raw-buffer-descriptor-aliasing.argument.discrete.device-argument-buffer.msl2.comp
@@ -0,0 +1,109 @@
+#version 450
+layout(local_size_x = 64) in;
+
+layout(set = 0, binding = 0) buffer SSBO_A
+{
+	float data[];
+} ssbo_a;
+
+layout(set = 0, binding = 0) buffer SSBO_B
+{
+	uvec2 data[];
+} ssbo_b;
+
+layout(set = 0, binding = 0) readonly buffer SSBO_BRO
+{
+	uvec2 data[];
+} ssbo_b_readonly;
+
+layout(set = 0, binding = 1) uniform UBO_C
+{
+	float data[1024];
+} ubo_c;
+
+layout(set = 0, binding = 1) uniform UBO_D
+{
+	uvec2 data[1024];
+} ubo_d;
+
+layout(set = 0, binding = 2) buffer SSBO_As 
+{
+	float data[];
+} ssbo_as[4];
+
+layout(set = 0, binding = 2) buffer SSBO_Bs
+{
+	uvec2 data[1024];
+} ssbo_bs[4];
+
+layout(set = 0, binding = 2) readonly buffer SSBO_BsRO
+{
+	uvec2 data[1024];
+} ssbo_bs_readonly[4];
+
+layout(set = 0, binding = 3) uniform UBO_Cs
+{
+	float data[1024];
+} ubo_cs[4];
+
+layout(set = 0, binding = 3) uniform UBO_Ds
+{
+	uvec2 data[1024];
+} ubo_ds[4];
+
+layout(set = 2, binding = 0) buffer SSBO_E
+{
+	float data[];
+} ssbo_e;
+
+layout(set = 2, binding = 0) buffer SSBO_F
+{
+	uvec2 data[];
+} ssbo_f;
+
+layout(set = 2, binding = 1) uniform UBO_G
+{
+	float data[1024];
+} ubo_g;
+
+layout(set = 2, binding = 1) uniform UBO_H
+{
+	uvec2 data[1024];
+} ubo_h;
+
+layout(set = 2, binding = 0) readonly buffer SSBO_I
+{
+	uvec2 data[];
+} ssbo_i;
+
+void func0()
+{
+	ssbo_a.data[gl_GlobalInvocationID.x] = ubo_c.data[gl_WorkGroupID.x];
+	ssbo_b.data[gl_GlobalInvocationID.x] =
+		ubo_d.data[gl_WorkGroupID.y] + ssbo_b_readonly.data[gl_GlobalInvocationID.x];
+}
+
+void func1()
+{
+	ssbo_as[gl_WorkGroupID.x].data[gl_GlobalInvocationID.x] = ubo_cs[gl_WorkGroupID.x].data[0];
+}
+
+void func2()
+{
+	ssbo_bs[gl_WorkGroupID.x].data[gl_GlobalInvocationID.x] =
+		ubo_ds[gl_WorkGroupID.x].data[0] + ssbo_bs_readonly[gl_WorkGroupID.x].data[gl_GlobalInvocationID.x];
+}
+
+void func3()
+{
+	ssbo_e.data[gl_GlobalInvocationID.x] = ubo_g.data[gl_WorkGroupID.x];
+	ssbo_f.data[gl_GlobalInvocationID.x] = ubo_h.data[gl_WorkGroupID.y] + ssbo_i.data[gl_GlobalInvocationID.x];
+}
+
+void main()
+{
+	func0();
+	func1();
+	func2();
+	func3();
+}
diff --git a/shaders-msl/comp/raw-buffer-descriptor-aliasing.argument.discrete.msl2.comp b/shaders-msl/comp/raw-buffer-descriptor-aliasing.argument.discrete.msl2.comp
new file mode 100644
index 00000000..25ec7840
--- /dev/null
+++ b/shaders-msl/comp/raw-buffer-descriptor-aliasing.argument.discrete.msl2.comp
@@ -0,0 +1,109 @@
+#version 450
+layout(local_size_x = 64) in;
+
+layout(set = 0, binding = 0) buffer SSBO_A
+{
+	float data[];
+} ssbo_a;
+
+layout(set = 0, binding = 0) buffer SSBO_B
+{
+	uvec2 data[];
+} ssbo_b;
+
+layout(set = 0, binding = 0) readonly buffer SSBO_BRO
+{
+	uvec2 data[];
+} ssbo_b_readonly;
+
+layout(set = 0, binding = 1) uniform UBO_C
+{
+	float data[1024];
+} ubo_c;
+
+layout(set = 0, binding = 1) uniform UBO_D
+{
+	uvec2 data[1024];
+} ubo_d;
+
+layout(set = 0, binding = 2) buffer SSBO_As 
+{
+	float data[];
+} ssbo_as[4];
+
+layout(set = 0, binding = 2) buffer SSBO_Bs
+{
+	uvec2 data[1024];
+} ssbo_bs[4];
+
+layout(set = 0, binding = 2) readonly buffer SSBO_BsRO
+{
+	uvec2 data[1024];
+} ssbo_bs_readonly[4];
+
+layout(set = 0, binding = 3) uniform UBO_Cs
+{
+	float data[1024];
+} ubo_cs[4];
+
+layout(set = 0, binding = 3) uniform UBO_Ds
+{
+	uvec2 data[1024];
+} ubo_ds[4];
+
+layout(set = 2, binding = 0) buffer SSBO_E
+{
+	float data[];
+} ssbo_e;
+
+layout(set = 2, binding = 0) buffer SSBO_F
+{
+	uvec2 data[];
+} ssbo_f;
+
+layout(set = 2, binding = 1) uniform UBO_G
+{
+	float data[1024];
+} ubo_g;
+
+layout(set = 2, binding = 1) uniform UBO_H
+{
+	uvec2 data[1024];
+} ubo_h;
+
+layout(set = 2, binding = 0) readonly buffer SSBO_I
+{
+	uvec2 data[];
+} ssbo_i;
+
+void func0()
+{
+	ssbo_a.data[gl_GlobalInvocationID.x] = ubo_c.data[gl_WorkGroupID.x];
+	ssbo_b.data[gl_GlobalInvocationID.x] =
+		ubo_d.data[gl_WorkGroupID.y] + ssbo_b_readonly.data[gl_GlobalInvocationID.x];
+}
+
+void func1()
+{
+	ssbo_as[gl_WorkGroupID.x].data[gl_GlobalInvocationID.x] = ubo_cs[gl_WorkGroupID.x].data[0];
+}
+
+void func2()
+{
+	ssbo_bs[gl_WorkGroupID.x].data[gl_GlobalInvocationID.x] =
+		ubo_ds[gl_WorkGroupID.x].data[0] + ssbo_bs_readonly[gl_WorkGroupID.x].data[gl_GlobalInvocationID.x];
+}
+
+void func3()
+{
+	ssbo_e.data[gl_GlobalInvocationID.x] = ubo_g.data[gl_WorkGroupID.x];
+	ssbo_f.data[gl_GlobalInvocationID.x] = ubo_h.data[gl_WorkGroupID.y] + ssbo_i.data[gl_GlobalInvocationID.x];
+}
+
+void main()
+{
+	func0();
+	func1();
+	func2();
+	func3();
+}
diff --git a/spirv_msl.cpp b/spirv_msl.cpp
index 9610bfa8..d01ffd4b 100644
--- a/spirv_msl.cpp
+++ b/spirv_msl.cpp
@@ -1314,7 +1314,7 @@ void CompilerMSL::emit_entry_point_declarations()
 	}
 
 	// Emit buffer arrays here.
-	for (uint32_t array_id : buffer_arrays)
+	for (uint32_t array_id : buffer_arrays_discrete)
 	{
 		const auto &var = get<SPIRVariable>(array_id);
 		const auto &type = get_variable_data_type(var);
@@ -1328,8 +1328,57 @@ void CompilerMSL::emit_entry_point_declarations()
 		end_scope_decl();
 		statement_no_indent("");
 	}
-	// For some reason, without this, we end up emitting the arrays twice.
-	buffer_arrays.clear();
+	// Discrete descriptors are processed in entry point emission every compiler iteration.
+	buffer_arrays_discrete.clear();
+
+	// Emit buffer aliases here.
+	for (auto &var_id : buffer_aliases_discrete)
+	{
+		const auto &var = get<SPIRVariable>(var_id);
+		const auto &type = get_variable_data_type(var);
+		auto addr_space = get_argument_address_space(var);
+		auto name = to_name(var_id);
+
+		uint32_t desc_set = get_decoration(var_id, DecorationDescriptorSet);
+		uint32_t desc_binding = get_decoration(var_id, DecorationBinding);
+		auto alias_name = join("spvBufferAliasSet", desc_set, "Binding", desc_binding);
+
+		statement(addr_space, " auto& ", to_restrict(var_id),
+		          name,
+		          " = *(", addr_space, " ", type_to_glsl(type), "*)", alias_name, ";");
+	}
+	// Discrete descriptors are processed in entry point emission every compiler iteration.
+	buffer_aliases_discrete.clear();
+
+	for (auto &var_pair : buffer_aliases_argument)
+	{
+		uint32_t var_id = var_pair.first;
+		uint32_t alias_id = var_pair.second;
+
+		const auto &var = get<SPIRVariable>(var_id);
+		const auto &type = get_variable_data_type(var);
+		auto addr_space = get_argument_address_space(var);
+
+		if (type.array.empty())
+		{
+			statement(addr_space, " auto& ", to_restrict(var_id), to_name(var_id), " = (", addr_space, " ",
+			          type_to_glsl(type), "&)", ir.meta[alias_id].decoration.qualified_alias, ";");
+		}
+		else
+		{
+			const char *desc_addr_space = descriptor_address_space(var_id, var.storage, "thread");
+
+			// Esoteric type cast. Reference to array of pointers.
+			// Auto here defers to UBO or SSBO. The address space of the reference needs to refer to the
+			// address space of the argument buffer itself, which is usually constant, but can be const device for
+			// large argument buffers.
+			is_using_builtin_array = true;
+			statement(desc_addr_space, " auto& ", to_restrict(var_id), to_name(var_id), " = (", addr_space, " ",
+			          type_to_glsl(type), "* ", desc_addr_space, " (&)",
+			          type_to_array_glsl(type), ")", ir.meta[alias_id].decoration.qualified_alias, ";");
+			is_using_builtin_array = false;
+		}
+	}
 
 	// Emit disabled fragment outputs.
 	std::sort(disabled_frag_outputs.begin(), disabled_frag_outputs.end());
@@ -12392,6 +12441,7 @@ void CompilerMSL::entry_point_args_discrete_descriptors(string &ep_args)
 	struct Resource
 	{
 		SPIRVariable *var;
+		SPIRVariable *descriptor_alias;
 		string name;
 		SPIRType::BaseType basetype;
 		uint32_t index;
@@ -12415,6 +12465,31 @@ void CompilerMSL::entry_point_args_discrete_descriptors(string &ep_args)
 					return;
 			}
 
+			// Handle descriptor aliasing. We can handle aliasing of buffers by casting pointers,
+			// but not for typed resources.
+			SPIRVariable *descriptor_alias = nullptr;
+			for (auto &resource : resources)
+			{
+				if (get_decoration(resource.var->self, DecorationDescriptorSet) == get_decoration(var_id, DecorationDescriptorSet) &&
+				    get_decoration(resource.var->self, DecorationBinding) == get_decoration(var_id, DecorationBinding) &&
+				    resource.basetype == SPIRType::Struct &&
+				    type.basetype == SPIRType::Struct)
+				{
+					// Possible, but horrible to implement, ignore for now.
+					if (!type.array.empty())
+						SPIRV_CROSS_THROW("Aliasing arrayed discrete descriptors is currently not supported.");
+
+					descriptor_alias = resource.var;
+					// Self-reference marks that we should declare the resource,
+					// and it's being used as an alias (so we can emit void* instead).
+					resource.descriptor_alias = resource.var;
+					// Need to promote interlocked usage so that the primary declaration is correct.
+					if (interlocked_resources.count(var_id))
+						interlocked_resources.insert(resource.var->self);
+					break;
+				}
+			}
+
 			const MSLConstexprSampler *constexpr_sampler = nullptr;
 			if (type.basetype == SPIRType::SampledImage || type.basetype == SPIRType::Sampler)
 			{
@@ -12442,12 +12517,12 @@ void CompilerMSL::entry_point_args_discrete_descriptors(string &ep_args)
 					plane_count = constexpr_sampler->planes;
 
 				for (uint32_t i = 0; i < plane_count; i++)
-					resources.push_back({ &var, to_name(var_id), SPIRType::Image,
+					resources.push_back({ &var, descriptor_alias, to_name(var_id), SPIRType::Image,
 					                      get_metal_resource_index(var, SPIRType::Image, i), i, secondary_index });
 
 				if (type.image.dim != DimBuffer && !constexpr_sampler)
 				{
-					resources.push_back({ &var, to_sampler_expression(var_id), SPIRType::Sampler,
+					resources.push_back({ &var, descriptor_alias, to_sampler_expression(var_id), SPIRType::Sampler,
 					                      get_metal_resource_index(var, SPIRType::Sampler), 0, 0 });
 				}
 			}
@@ -12455,13 +12530,19 @@ void CompilerMSL::entry_point_args_discrete_descriptors(string &ep_args)
 			{
 				// constexpr samplers are not declared as resources.
 				add_resource_name(var_id);
-				resources.push_back({ &var, to_name(var_id), type.basetype,
-				                      get_metal_resource_index(var, type.basetype), 0, secondary_index });
+
+				// Don't allocate resource indices for aliases.
+				uint32_t resource_index = ~0u;
+				if (!descriptor_alias)
+					resource_index = get_metal_resource_index(var, type.basetype);
+
+				resources.push_back({ &var, descriptor_alias, to_name(var_id), type.basetype,
+				                      resource_index, 0, secondary_index });
 			}
 		}
 	});
 
-	sort(resources.begin(), resources.end(), [](const Resource &lhs, const Resource &rhs) {
+	stable_sort(resources.begin(), resources.end(), [](const Resource &lhs, const Resource &rhs) {
 		return tie(lhs.basetype, lhs.index) < tie(rhs.basetype, rhs.index);
 	});
 
@@ -12479,7 +12560,29 @@ void CompilerMSL::entry_point_args_discrete_descriptors(string &ep_args)
 			auto &m = ir.meta[type.self];
 			if (m.members.size() == 0)
 				break;
-			if (!type.array.empty())
+
+			if (r.descriptor_alias)
+			{
+				if (r.var == r.descriptor_alias)
+				{
+					auto primary_name = join("spvBufferAliasSet",
+					                         get_decoration(var_id, DecorationDescriptorSet),
+					                         "Binding",
+					                         get_decoration(var_id, DecorationBinding));
+
+					// Declare the primary alias as void*
+					if (!ep_args.empty())
+						ep_args += ", ";
+					ep_args += get_argument_address_space(var) + " void* " + primary_name;
+					ep_args += " [[buffer(" + convert_to_string(r.index) + ")";
+					if (interlocked_resources.count(var_id))
+						ep_args += ", raster_order_group(0)";
+					ep_args += "]]";
+				}
+
+				buffer_aliases_discrete.push_back(r.var->self);
+			}
+			else if (!type.array.empty())
 			{
 				if (type.array.size() > 1)
 					SPIRV_CROSS_THROW("Arrays of arrays of buffers are not supported.");
@@ -12494,7 +12597,7 @@ void CompilerMSL::entry_point_args_discrete_descriptors(string &ep_args)
 
 				// Allow Metal to use the array<T> template to make arrays a value type
 				is_using_builtin_array = true;
-				buffer_arrays.push_back(var_id);
+				buffer_arrays_discrete.push_back(var_id);
 				for (uint32_t i = 0; i < array_size; ++i)
 				{
 					if (!ep_args.empty())
@@ -16471,6 +16574,7 @@ void CompilerMSL::analyze_argument_buffers()
 	struct Resource
 	{
 		SPIRVariable *var;
+		SPIRVariable *descriptor_alias;
 		string name;
 		SPIRType::BaseType basetype;
 		uint32_t index;
@@ -16510,6 +16614,27 @@ void CompilerMSL::analyze_argument_buffers()
 				}
 			}
 
+			// Handle descriptor aliasing as well as we can.
+			// We can handle aliasing of buffers by casting pointers, but not for typed resources.
+			// Inline UBOs cannot be handled since it's not a pointer, but inline data.
+			SPIRVariable *descriptor_alias = nullptr;
+			for (auto &resource : resources_in_set[desc_set])
+			{
+				if (get_decoration(resource.var->self, DecorationBinding) == get_decoration(var_id, DecorationBinding) &&
+				    resource.basetype == SPIRType::Struct &&
+				    type.basetype == SPIRType::Struct)
+				{
+					descriptor_alias = resource.var;
+					// Self-reference marks that we should declare the resource,
+					// and it's being used as an alias (so we can emit void* instead).
+					resource.descriptor_alias = resource.var;
+					// Need to promote interlocked usage so that the primary declaration is correct.
+					if (interlocked_resources.count(var_id))
+						interlocked_resources.insert(resource.var->self);
+					break;
+				}
+			}
+
 			uint32_t binding = get_decoration(var_id, DecorationBinding);
 			if (type.basetype == SPIRType::SampledImage)
 			{
@@ -16523,14 +16648,14 @@ void CompilerMSL::analyze_argument_buffers()
 				{
 					uint32_t image_resource_index = get_metal_resource_index(var, SPIRType::Image, i);
 					resources_in_set[desc_set].push_back(
-					    { &var, to_name(var_id), SPIRType::Image, image_resource_index, i });
+					    { &var, descriptor_alias, to_name(var_id), SPIRType::Image, image_resource_index, i });
 				}
 
 				if (type.image.dim != DimBuffer && !constexpr_sampler)
 				{
 					uint32_t sampler_resource_index = get_metal_resource_index(var, SPIRType::Sampler);
 					resources_in_set[desc_set].push_back(
-					    { &var, to_sampler_expression(var_id), SPIRType::Sampler, sampler_resource_index, 0 });
+					    { &var, descriptor_alias, to_sampler_expression(var_id), SPIRType::Sampler, sampler_resource_index, 0 });
 				}
 			}
 			else if (inline_uniform_blocks.count(SetBindingPair{ desc_set, binding }))
@@ -16542,15 +16667,20 @@ void CompilerMSL::analyze_argument_buffers()
 				// constexpr samplers are not declared as resources.
 				// Inline uniform blocks are always emitted at the end.
 				add_resource_name(var_id);
+
+				uint32_t resource_index = ~0u;
+				if (!descriptor_alias)
+					resource_index = get_metal_resource_index(var, type.basetype);
+
 				resources_in_set[desc_set].push_back(
-					{ &var, to_name(var_id), type.basetype, get_metal_resource_index(var, type.basetype), 0 });
+					{ &var, descriptor_alias, to_name(var_id), type.basetype, resource_index, 0 });
 
 				// Emulate texture2D atomic operations
 				if (atomic_image_vars.count(var.self))
 				{
 					uint32_t buffer_resource_index = get_metal_resource_index(var, SPIRType::AtomicCounter, 0);
 					resources_in_set[desc_set].push_back(
-						{ &var, to_name(var_id) + "_atomic", SPIRType::Struct, buffer_resource_index, 0 });
+						{ &var, descriptor_alias, to_name(var_id) + "_atomic", SPIRType::Struct, buffer_resource_index, 0 });
 				}
 			}
 
@@ -16597,7 +16727,7 @@ void CompilerMSL::analyze_argument_buffers()
 				set_decoration(var_id, DecorationDescriptorSet, desc_set);
 				set_decoration(var_id, DecorationBinding, kSwizzleBufferBinding);
 				resources_in_set[desc_set].push_back(
-				    { &var, to_name(var_id), SPIRType::UInt, get_metal_resource_index(var, SPIRType::UInt), 0 });
+				    { &var, nullptr, to_name(var_id), SPIRType::UInt, get_metal_resource_index(var, SPIRType::UInt), 0 });
 			}
 
 			if (set_needs_buffer_sizes[desc_set])
@@ -16608,7 +16738,7 @@ void CompilerMSL::analyze_argument_buffers()
 				set_decoration(var_id, DecorationDescriptorSet, desc_set);
 				set_decoration(var_id, DecorationBinding, kBufferSizeBufferBinding);
 				resources_in_set[desc_set].push_back(
-				    { &var, to_name(var_id), SPIRType::UInt, get_metal_resource_index(var, SPIRType::UInt), 0 });
+				    { &var, nullptr, to_name(var_id), SPIRType::UInt, get_metal_resource_index(var, SPIRType::UInt), 0 });
 			}
 		}
 	}
@@ -16620,7 +16750,7 @@ void CompilerMSL::analyze_argument_buffers()
 		uint32_t desc_set = get_decoration(var_id, DecorationDescriptorSet);
 		add_resource_name(var_id);
 		resources_in_set[desc_set].push_back(
-		    { &var, to_name(var_id), SPIRType::Struct, get_metal_resource_index(var, SPIRType::Struct), 0 });
+		    { &var, nullptr, to_name(var_id), SPIRType::Struct, get_metal_resource_index(var, SPIRType::Struct), 0 });
 	}
 
 	for (uint32_t desc_set = 0; desc_set < kMaxArgumentBuffers; desc_set++)
@@ -16664,7 +16794,7 @@ void CompilerMSL::analyze_argument_buffers()
 		set_name(buffer_variable_id, join("spvDescriptorSet", desc_set));
 
 		// Ids must be emitted in ID order.
-		sort(begin(resources), end(resources), [&](const Resource &lhs, const Resource &rhs) -> bool {
+		stable_sort(begin(resources), end(resources), [&](const Resource &lhs, const Resource &rhs) -> bool {
 			return tie(lhs.index, lhs.basetype) < tie(rhs.index, rhs.basetype);
 		});
 
@@ -16769,12 +16899,18 @@ void CompilerMSL::analyze_argument_buffers()
 				}
 				else if (buffers_requiring_dynamic_offset.count(pair))
 				{
+					if (resource.descriptor_alias)
+						SPIRV_CROSS_THROW("Descriptor aliasing is currently not supported with dynamic offsets.");
+
 					// Don't set the qualified name here; we'll define a variable holding the corrected buffer address later.
 					buffer_type.member_types.push_back(var.basetype);
 					buffers_requiring_dynamic_offset[pair].second = var.self;
 				}
 				else if (inline_uniform_blocks.count(pair))
 				{
+					if (resource.descriptor_alias)
+						SPIRV_CROSS_THROW("Descriptor aliasing is currently not supported with inline UBOs.");
+
 					// Put the buffer block itself into the argument buffer.
 					buffer_type.member_types.push_back(get_variable_data_type_id(var));
 					set_qualified_name(var.self, join(to_name(buffer_variable_id), ".", mbr_name));
@@ -16806,9 +16942,12 @@ void CompilerMSL::analyze_argument_buffers()
 				}
 				else
 				{
-					// Resources will be declared as pointers not references, so automatically dereference as appropriate.
-					buffer_type.member_types.push_back(var.basetype);
-					if (type.array.empty())
+					if (!resource.descriptor_alias || resource.descriptor_alias == resource.var)
+						buffer_type.member_types.push_back(var.basetype);
+
+					if (resource.descriptor_alias && resource.descriptor_alias != resource.var)
+						buffer_aliases_argument.push_back({ var.self, resource.descriptor_alias->self });
+					else if (type.array.empty())
 						set_qualified_name(var.self, join("(*", to_name(buffer_variable_id), ".", mbr_name, ")"));
 					else
 						set_qualified_name(var.self, join(to_name(buffer_variable_id), ".", mbr_name));
diff --git a/spirv_msl.hpp b/spirv_msl.hpp
index 0e5c76db..1a7ee5c0 100644
--- a/spirv_msl.hpp
+++ b/spirv_msl.hpp
@@ -1107,7 +1107,9 @@ protected:
 	const MSLConstexprSampler *find_constexpr_sampler(uint32_t id) const;
 
 	std::unordered_set<uint32_t> buffers_requiring_array_length;
-	SmallVector<uint32_t> buffer_arrays;
+	SmallVector<uint32_t> buffer_arrays_discrete;
+	SmallVector<std::pair<uint32_t, uint32_t>> buffer_aliases_argument;
+	SmallVector<uint32_t> buffer_aliases_discrete;
 	std::unordered_set<uint32_t> atomic_image_vars; // Emulate texture2D atomic operations
 	std::unordered_set<uint32_t> pull_model_inputs;