Merge pull request #2024 from KhronosGroup/msl-descriptor-aliasing

MSL: Do not attempt to alias push constants.

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;
+					}
 				}
 			}