1 files changed, 111 insertions, 0 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];
+}
+