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

MSL: Reassign bindings so they work with decoration binding in test.

Browse Source
This commit is contained in:
Hans-Kristian Arntzen 2024-04-03 13:45:18 +02:00
parent ee77265ae5
commit 8219ccf706
3 changed files with 36 additions and 36 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

26
reference/opt/shaders-msl/comp/raw-buffer-descriptor-aliasing.argument.discrete.device-argument-buffer.msl2.decoration-binding.comp
View File

@ -88,31 +88,31 @@ constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(64u, 1u, 1u);
struct spvDescriptorSetBuffer0
{
device SSBO_A* ssbo_a [[id(0)]];
// Overlapping binding: device SSBO_B* ssbo_b [[id(0)]];
// Overlapping binding: const device SSBO_BRO* ssbo_b_readonly [[id(0)]];
constant UBO_C* ubo_c [[id(1)]];
// Overlapping binding: constant UBO_D* ubo_d [[id(1)]];
device SSBO_As* ssbo_as [[id(2)]][4];
// Overlapping binding: device SSBO_Bs* ssbo_bs [[id(2)]][4];
// Overlapping binding: const device SSBO_BsRO* ssbo_bs_readonly [[id(2)]][4];
constant UBO_Cs* ubo_cs [[id(3)]][4];
// Overlapping binding: constant UBO_Ds* ubo_ds [[id(3)]][4];
constant UBO_Cs* ubo_cs [[id(6)]][4];
// Overlapping binding: constant UBO_Ds* ubo_ds [[id(6)]][4];
device SSBO_A* ssbo_a [[id(10)]];
// Overlapping binding: device SSBO_B* ssbo_b [[id(10)]];
// Overlapping binding: const device SSBO_BRO* ssbo_b_readonly [[id(10)]];
};
kernel void main0(const device spvDescriptorSetBuffer0& spvDescriptorSet0 [[buffer(0)]], device void* spvBufferAliasSet2Binding0 [[buffer(0)]], constant Registers& _42 [[buffer(1)]], constant void* spvBufferAliasSet2Binding1 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]])
kernel void main0(const device spvDescriptorSetBuffer0& spvDescriptorSet0 [[buffer(0)]], constant Registers& _42 [[buffer(1)]], device void* spvBufferAliasSet2Binding11 [[buffer(11)]], constant void* spvBufferAliasSet2Binding12 [[buffer(12)]], 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 = *reinterpret_cast<device SSBO_B* const device &>(spvDescriptorSet0.ssbo_a);
const device auto &ssbo_b_readonly = *reinterpret_cast<const device SSBO_BRO* const device &>(spvDescriptorSet0.ssbo_a);
device auto& ssbo_e = *(device SSBO_E*)spvBufferAliasSet2Binding11;
constant auto& ubo_g = *(constant UBO_G*)spvBufferAliasSet2Binding12;
device auto& ssbo_f = *(device SSBO_F*)spvBufferAliasSet2Binding11;
constant auto& ubo_h = *(constant UBO_H*)spvBufferAliasSet2Binding12;
const device auto& ssbo_i = *(const device SSBO_I*)spvBufferAliasSet2Binding11;
constant auto &ubo_d = *reinterpret_cast<constant UBO_D* const device &>(spvDescriptorSet0.ubo_c);
const device auto &ssbo_bs = reinterpret_cast<device SSBO_Bs* const device (&)[4]>(spvDescriptorSet0.ssbo_as);
const device auto &ssbo_bs_readonly = reinterpret_cast<const device SSBO_BsRO* const device (&)[4]>(spvDescriptorSet0.ssbo_as);
const device auto &ubo_ds = reinterpret_cast<constant UBO_Ds* const device (&)[4]>(spvDescriptorSet0.ubo_cs);
device auto &ssbo_b = *reinterpret_cast<device SSBO_B* const device &>(spvDescriptorSet0.ssbo_a);
const device auto &ssbo_b_readonly = *reinterpret_cast<const device SSBO_BRO* const device &>(spvDescriptorSet0.ssbo_a);
(*spvDescriptorSet0.ssbo_a).data[gl_GlobalInvocationID.x] = (*spvDescriptorSet0.ubo_c).data[gl_WorkGroupID.x].x + _42.reg;
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;

26
reference/shaders-msl/comp/raw-buffer-descriptor-aliasing.argument.discrete.device-argument-buffer.msl2.decoration-binding.comp
View File

@ -90,16 +90,16 @@ constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(64u, 1u, 1u);
struct spvDescriptorSetBuffer0
{
device SSBO_A* ssbo_a [[id(0)]];
// Overlapping binding: device SSBO_B* ssbo_b [[id(0)]];
// Overlapping binding: const device SSBO_BRO* ssbo_b_readonly [[id(0)]];
constant UBO_C* ubo_c [[id(1)]];
// Overlapping binding: constant UBO_D* ubo_d [[id(1)]];
device SSBO_As* ssbo_as [[id(2)]][4];
// Overlapping binding: device SSBO_Bs* ssbo_bs [[id(2)]][4];
// Overlapping binding: const device SSBO_BsRO* ssbo_bs_readonly [[id(2)]][4];
constant UBO_Cs* ubo_cs [[id(3)]][4];
// Overlapping binding: constant UBO_Ds* ubo_ds [[id(3)]][4];
constant UBO_Cs* ubo_cs [[id(6)]][4];
// Overlapping binding: constant UBO_Ds* ubo_ds [[id(6)]][4];
device SSBO_A* ssbo_a [[id(10)]];
// Overlapping binding: device SSBO_B* ssbo_b [[id(10)]];
// Overlapping binding: const device SSBO_BRO* ssbo_b_readonly [[id(10)]];
};
static inline __attribute__((always_inline))
@ -128,19 +128,19 @@ void func3(thread uint3& gl_GlobalInvocationID, thread uint3& gl_WorkGroupID, de
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(0)]], constant Registers& _42 [[buffer(1)]], constant void* spvBufferAliasSet2Binding1 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]])
kernel void main0(const device spvDescriptorSetBuffer0& spvDescriptorSet0 [[buffer(0)]], constant Registers& _42 [[buffer(1)]], device void* spvBufferAliasSet2Binding11 [[buffer(11)]], constant void* spvBufferAliasSet2Binding12 [[buffer(12)]], 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 = *reinterpret_cast<device SSBO_B* const device &>(spvDescriptorSet0.ssbo_a);
const device auto &ssbo_b_readonly = *reinterpret_cast<const device SSBO_BRO* const device &>(spvDescriptorSet0.ssbo_a);
device auto& ssbo_e = *(device SSBO_E*)spvBufferAliasSet2Binding11;
constant auto& ubo_g = *(constant UBO_G*)spvBufferAliasSet2Binding12;
device auto& ssbo_f = *(device SSBO_F*)spvBufferAliasSet2Binding11;
constant auto& ubo_h = *(constant UBO_H*)spvBufferAliasSet2Binding12;
const device auto& ssbo_i = *(const device SSBO_I*)spvBufferAliasSet2Binding11;
constant auto &ubo_d = *reinterpret_cast<constant UBO_D* const device &>(spvDescriptorSet0.ubo_c);
const device auto &ssbo_bs = reinterpret_cast<device SSBO_Bs* const device (&)[4]>(spvDescriptorSet0.ssbo_as);
const device auto &ssbo_bs_readonly = reinterpret_cast<const device SSBO_BsRO* const device (&)[4]>(spvDescriptorSet0.ssbo_as);
const device auto &ubo_ds = reinterpret_cast<constant UBO_Ds* const device (&)[4]>(spvDescriptorSet0.ubo_cs);
device auto &ssbo_b = *reinterpret_cast<device SSBO_B* const device &>(spvDescriptorSet0.ssbo_a);
const device auto &ssbo_b_readonly = *reinterpret_cast<const device SSBO_BRO* const device &>(spvDescriptorSet0.ssbo_a);
func0((*spvDescriptorSet0.ssbo_a), gl_GlobalInvocationID, (*spvDescriptorSet0.ubo_c), gl_WorkGroupID, _42, 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);

20
shaders-msl/comp/raw-buffer-descriptor-aliasing.argument.discrete.device-argument-buffer.msl2.decoration-binding.comp
View File

@ -1,17 +1,17 @@
#version 450
layout(local_size_x = 64) in;
layout(set = 0, binding = 0) buffer SSBO_A
layout(set = 0, binding = 10) buffer SSBO_A
{
float data[];
} ssbo_a;
layout(set = 0, binding = 0) buffer SSBO_B
layout(set = 0, binding = 10) buffer SSBO_B
{
uvec2 data[];
} ssbo_b;
layout(set = 0, binding = 0) readonly buffer SSBO_BRO
layout(set = 0, binding = 10) readonly buffer SSBO_BRO
{
uvec2 data[];
} ssbo_b_readonly;
@ -41,37 +41,37 @@ layout(set = 0, binding = 2) readonly buffer SSBO_BsRO
uvec2 data[1024];
} ssbo_bs_readonly[4];
layout(set = 0, binding = 3) uniform UBO_Cs
layout(set = 0, binding = 6) uniform UBO_Cs
{
float data[1024];
} ubo_cs[4];
layout(set = 0, binding = 3) uniform UBO_Ds
layout(set = 0, binding = 6) uniform UBO_Ds
{
uvec2 data[1024];
} ubo_ds[4];
layout(set = 2, binding = 0) buffer SSBO_E
layout(set = 2, binding = 11) buffer SSBO_E
{
float data[];
} ssbo_e;
layout(set = 2, binding = 0) buffer SSBO_F
layout(set = 2, binding = 11) buffer SSBO_F
{
uvec2 data[];
} ssbo_f;
layout(set = 2, binding = 1) uniform UBO_G
layout(set = 2, binding = 12) uniform UBO_G
{
float data[1024];
} ubo_g;
layout(set = 2, binding = 1) uniform UBO_H
layout(set = 2, binding = 12) uniform UBO_H
{
uvec2 data[1024];
} ubo_h;
layout(set = 2, binding = 0) readonly buffer SSBO_I
layout(set = 2, binding = 11) readonly buffer SSBO_I
{
uvec2 data[];
} ssbo_i;