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 +#include + +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 +#include + +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 +#include + +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 +#include + +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(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(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(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 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 buffers_requiring_array_length; - SmallVector buffer_arrays; + SmallVector buffer_arrays_discrete; + SmallVector> buffer_aliases_argument; + SmallVector buffer_aliases_discrete; std::unordered_set atomic_image_vars; // Emulate texture2D atomic operations std::unordered_set pull_model_inputs;