AuroraMiddleware/SPIRV-Cross
1
0
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity

MSL: Handle descriptor aliasing of raw buffer descriptors.

Browse Source 
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.
This commit is contained in:
Hans-Kristian Arntzen 2022-09-20 11:57:19 +02:00
parent 1ad6006130
commit 24dc49e692
9 changed files with 878 additions and 23 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

111
reference/opt/shaders-msl/comp/raw-buffer-descriptor-aliasing.argument.discrete.device-argument-buffer.msl2.comp Normal file
View File

@ -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];
}

111
reference/opt/shaders-msl/comp/raw-buffer-descriptor-aliasing.argument.discrete.msl2.comp Normal file
View File

@ -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];
}

137
reference/shaders-msl/comp/raw-buffer-descriptor-aliasing.argument.discrete.device-argument-buffer.msl2.comp Normal file
View File

@ -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);
}

137
reference/shaders-msl/comp/raw-buffer-descriptor-aliasing.argument.discrete.msl2.comp Normal file
View File

@ -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);
}

2
shaders-msl-no-opt/asm/comp/block-like-array-type-construct-2.asm.comp
View File

@ -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

109
shaders-msl/comp/raw-buffer-descriptor-aliasing.argument.discrete.device-argument-buffer.msl2.comp Normal file
View File

@ -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();
}

109
shaders-msl/comp/raw-buffer-descriptor-aliasing.argument.discrete.msl2.comp Normal file
View File

@ -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();
}

181
spirv_msl.cpp
View File

@ -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));

4
spirv_msl.hpp
View File

@ -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;