From a915e0bd4be6f012a17b2abbc56bf354094f0958 Mon Sep 17 00:00:00 2001 From: Hans-Kristian Arntzen Date: Wed, 21 Sep 2022 11:01:42 +0200 Subject: MSL: Do not attempt to alias push constants. --- ...ument.discrete.device-argument-buffer.msl2.comp | 9 ++- ...descriptor-aliasing.argument.discrete.msl2.comp | 9 ++- ...ument.discrete.device-argument-buffer.msl2.comp | 13 ++-- ...descriptor-aliasing.argument.discrete.msl2.comp | 13 ++-- ...ument.discrete.device-argument-buffer.msl2.comp | 7 ++- ...descriptor-aliasing.argument.discrete.msl2.comp | 7 ++- spirv_msl.cpp | 69 +++++++++++++--------- 7 files changed, 84 insertions(+), 43 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 index a3c1a5b3..c1191866 100644 --- 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 @@ -13,6 +13,11 @@ struct UBO_C float4 data[1024]; }; +struct Registers +{ + float reg; +}; + struct SSBO_B { uint2 data[1]; @@ -88,7 +93,7 @@ struct spvDescriptorSetBuffer0 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]]) +kernel void main0(const device spvDescriptorSetBuffer0& spvDescriptorSet0 [[buffer(0)]], constant Registers& _42 [[buffer(1)]], device void* spvBufferAliasSet2Binding0 [[buffer(2)]], constant void* spvBufferAliasSet2Binding1 [[buffer(3)]], 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; @@ -101,7 +106,7 @@ kernel void main0(const device spvDescriptorSetBuffer0& spvDescriptorSet0 [[buff 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; + (*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; 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]; 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 index bc0aa461..9cef6b20 100644 --- 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 @@ -13,6 +13,11 @@ struct UBO_C float4 data[1024]; }; +struct Registers +{ + float reg; +}; + struct SSBO_B { uint2 data[1]; @@ -88,7 +93,7 @@ struct spvDescriptorSetBuffer0 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]]) +kernel void main0(constant spvDescriptorSetBuffer0& spvDescriptorSet0 [[buffer(0)]], constant Registers& _42 [[buffer(1)]], device void* spvBufferAliasSet2Binding0 [[buffer(2)]], constant void* spvBufferAliasSet2Binding1 [[buffer(3)]], 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; @@ -101,7 +106,7 @@ kernel void main0(constant spvDescriptorSetBuffer0& spvDescriptorSet0 [[buffer(0 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; + (*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; 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]; 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 index c5dc95e8..14723cbe 100644 --- 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 @@ -15,6 +15,11 @@ struct UBO_C float4 data[1024]; }; +struct Registers +{ + float reg; +}; + struct SSBO_B { uint2 data[1]; @@ -91,9 +96,9 @@ struct spvDescriptorSetBuffer0 }; 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) +void func0(device SSBO_A& ssbo_a, thread uint3& gl_GlobalInvocationID, constant UBO_C& ubo_c, thread uint3& gl_WorkGroupID, constant Registers& v_42, 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_a.data[gl_GlobalInvocationID.x] = ubo_c.data[gl_WorkGroupID.x].x + v_42.reg; ssbo_b.data[gl_GlobalInvocationID.x] = ubo_d.data[gl_WorkGroupID.y].xy + ssbo_b_readonly.data[gl_GlobalInvocationID.x]; } @@ -116,7 +121,7 @@ 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(1)]], constant void* spvBufferAliasSet2Binding1 [[buffer(2)]], 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& v_42 [[buffer(1)]], device void* spvBufferAliasSet2Binding0 [[buffer(2)]], constant void* spvBufferAliasSet2Binding1 [[buffer(3)]], 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; @@ -129,7 +134,7 @@ kernel void main0(const device spvDescriptorSetBuffer0& spvDescriptorSet0 [[buff 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); + func0((*spvDescriptorSet0.ssbo_a), gl_GlobalInvocationID, (*spvDescriptorSet0.ubo_c), gl_WorkGroupID, v_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); 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 index bdc5bc1c..587f1ee8 100644 --- 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 @@ -15,6 +15,11 @@ struct UBO_C float4 data[1024]; }; +struct Registers +{ + float reg; +}; + struct SSBO_B { uint2 data[1]; @@ -91,9 +96,9 @@ struct spvDescriptorSetBuffer0 }; 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) +void func0(device SSBO_A& ssbo_a, thread uint3& gl_GlobalInvocationID, constant UBO_C& ubo_c, thread uint3& gl_WorkGroupID, constant Registers& v_42, 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_a.data[gl_GlobalInvocationID.x] = ubo_c.data[gl_WorkGroupID.x].x + v_42.reg; ssbo_b.data[gl_GlobalInvocationID.x] = ubo_d.data[gl_WorkGroupID.y].xy + ssbo_b_readonly.data[gl_GlobalInvocationID.x]; } @@ -116,7 +121,7 @@ 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(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]]) +kernel void main0(constant spvDescriptorSetBuffer0& spvDescriptorSet0 [[buffer(0)]], constant Registers& v_42 [[buffer(1)]], device void* spvBufferAliasSet2Binding0 [[buffer(2)]], constant void* spvBufferAliasSet2Binding1 [[buffer(3)]], 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; @@ -129,7 +134,7 @@ kernel void main0(constant spvDescriptorSetBuffer0& spvDescriptorSet0 [[buffer(0 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); + func0((*spvDescriptorSet0.ssbo_a), gl_GlobalInvocationID, (*spvDescriptorSet0.ubo_c), gl_WorkGroupID, v_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); func3(gl_GlobalInvocationID, gl_WorkGroupID, ssbo_e, ubo_g, ssbo_f, ubo_h, ssbo_i); 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 index 25ec7840..eea6a3df 100644 --- 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 @@ -76,9 +76,14 @@ layout(set = 2, binding = 0) readonly buffer SSBO_I uvec2 data[]; } ssbo_i; +layout(push_constant) uniform Registers +{ + float reg; +}; + void func0() { - ssbo_a.data[gl_GlobalInvocationID.x] = ubo_c.data[gl_WorkGroupID.x]; + ssbo_a.data[gl_GlobalInvocationID.x] = ubo_c.data[gl_WorkGroupID.x] + reg; ssbo_b.data[gl_GlobalInvocationID.x] = ubo_d.data[gl_WorkGroupID.y] + ssbo_b_readonly.data[gl_GlobalInvocationID.x]; } 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 index 25ec7840..eea6a3df 100644 --- a/shaders-msl/comp/raw-buffer-descriptor-aliasing.argument.discrete.msl2.comp +++ b/shaders-msl/comp/raw-buffer-descriptor-aliasing.argument.discrete.msl2.comp @@ -76,9 +76,14 @@ layout(set = 2, binding = 0) readonly buffer SSBO_I uvec2 data[]; } ssbo_i; +layout(push_constant) uniform Registers +{ + float reg; +}; + void func0() { - ssbo_a.data[gl_GlobalInvocationID.x] = ubo_c.data[gl_WorkGroupID.x]; + ssbo_a.data[gl_GlobalInvocationID.x] = ubo_c.data[gl_WorkGroupID.x] + reg; ssbo_b.data[gl_GlobalInvocationID.x] = ubo_d.data[gl_WorkGroupID.y] + ssbo_b_readonly.data[gl_GlobalInvocationID.x]; } diff --git a/spirv_msl.cpp b/spirv_msl.cpp index d01ffd4b..dd80fbbc 100644 --- a/spirv_msl.cpp +++ b/spirv_msl.cpp @@ -12468,25 +12468,31 @@ void CompilerMSL::entry_point_args_discrete_descriptors(string &ep_args) // 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 (var.storage == StorageClassUniform || var.storage == StorageClassStorageBuffer) { - 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) + for (auto &resource : resources) { - // 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; + 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 && + (resource.var->storage == StorageClassUniform || + resource.var->storage == StorageClassStorageBuffer)) + { + // 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; + } } } @@ -16618,20 +16624,25 @@ void CompilerMSL::analyze_argument_buffers() // 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 (var.storage == StorageClassUniform || var.storage == StorageClassStorageBuffer) { - if (get_decoration(resource.var->self, DecorationBinding) == get_decoration(var_id, DecorationBinding) && - resource.basetype == SPIRType::Struct && - type.basetype == SPIRType::Struct) + for (auto &resource : resources_in_set[desc_set]) { - 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; + if (get_decoration(resource.var->self, DecorationBinding) == + get_decoration(var_id, DecorationBinding) && + resource.basetype == SPIRType::Struct && type.basetype == SPIRType::Struct && + (resource.var->storage == StorageClassUniform || + resource.var->storage == StorageClassStorageBuffer)) + { + 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; + } } } -- cgit v1.2.3