From 8219ccf706b368bf7e79b522eb6c9f48a5abed32 Mon Sep 17 00:00:00 2001 From: Hans-Kristian Arntzen Date: Wed, 3 Apr 2024 13:45:18 +0200 Subject: [PATCH] MSL: Reassign bindings so they work with decoration binding in test. --- ...gument-buffer.msl2.decoration-binding.comp | 26 +++++++++---------- ...gument-buffer.msl2.decoration-binding.comp | 26 +++++++++---------- ...gument-buffer.msl2.decoration-binding.comp | 20 +++++++------- 3 files changed, 36 insertions(+), 36 deletions(-) diff --git a/reference/opt/shaders-msl/comp/raw-buffer-descriptor-aliasing.argument.discrete.device-argument-buffer.msl2.decoration-binding.comp b/reference/opt/shaders-msl/comp/raw-buffer-descriptor-aliasing.argument.discrete.device-argument-buffer.msl2.decoration-binding.comp index 4cce144c..11140775 100644 --- a/reference/opt/shaders-msl/comp/raw-buffer-descriptor-aliasing.argument.discrete.device-argument-buffer.msl2.decoration-binding.comp +++ b/reference/opt/shaders-msl/comp/raw-buffer-descriptor-aliasing.argument.discrete.device-argument-buffer.msl2.decoration-binding.comp @@ -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(spvDescriptorSet0.ssbo_a); - const device auto &ssbo_b_readonly = *reinterpret_cast(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(spvDescriptorSet0.ubo_c); const device auto &ssbo_bs = reinterpret_cast(spvDescriptorSet0.ssbo_as); const device auto &ssbo_bs_readonly = reinterpret_cast(spvDescriptorSet0.ssbo_as); const device auto &ubo_ds = reinterpret_cast(spvDescriptorSet0.ubo_cs); + device auto &ssbo_b = *reinterpret_cast(spvDescriptorSet0.ssbo_a); + const device auto &ssbo_b_readonly = *reinterpret_cast(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; diff --git a/reference/shaders-msl/comp/raw-buffer-descriptor-aliasing.argument.discrete.device-argument-buffer.msl2.decoration-binding.comp b/reference/shaders-msl/comp/raw-buffer-descriptor-aliasing.argument.discrete.device-argument-buffer.msl2.decoration-binding.comp index 94b45470..a6897467 100644 --- a/reference/shaders-msl/comp/raw-buffer-descriptor-aliasing.argument.discrete.device-argument-buffer.msl2.decoration-binding.comp +++ b/reference/shaders-msl/comp/raw-buffer-descriptor-aliasing.argument.discrete.device-argument-buffer.msl2.decoration-binding.comp @@ -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(spvDescriptorSet0.ssbo_a); - const device auto &ssbo_b_readonly = *reinterpret_cast(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(spvDescriptorSet0.ubo_c); const device auto &ssbo_bs = reinterpret_cast(spvDescriptorSet0.ssbo_as); const device auto &ssbo_bs_readonly = reinterpret_cast(spvDescriptorSet0.ssbo_as); const device auto &ubo_ds = reinterpret_cast(spvDescriptorSet0.ubo_cs); + device auto &ssbo_b = *reinterpret_cast(spvDescriptorSet0.ssbo_a); + const device auto &ssbo_b_readonly = *reinterpret_cast(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); diff --git a/shaders-msl/comp/raw-buffer-descriptor-aliasing.argument.discrete.device-argument-buffer.msl2.decoration-binding.comp b/shaders-msl/comp/raw-buffer-descriptor-aliasing.argument.discrete.device-argument-buffer.msl2.decoration-binding.comp index eea6a3df..b9185862 100644 --- a/shaders-msl/comp/raw-buffer-descriptor-aliasing.argument.discrete.device-argument-buffer.msl2.decoration-binding.comp +++ b/shaders-msl/comp/raw-buffer-descriptor-aliasing.argument.discrete.device-argument-buffer.msl2.decoration-binding.comp @@ -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;