MSL: Deal with resource name aliasing.

Apparently we didn't use those yet. MSL seems to be able to alias struct
types and variable types to a degree, so that's why it has escaped
testing until now.

reference/opt/shaders-msl/asm/comp/struct-resource-name-aliasing.asm.comp

@@ -0,0 +1,16 @@
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+struct bufA
+{
+    uint _data[1];
+};
+
+kernel void main0(device bufA& bufA_1 [[buffer(0)]], device bufA& bufB [[buffer(1)]])
+{
+    bufA_1._data[0] = 0u;
+    bufB._data[0] = 0u;
+}
+

reference/shaders-msl/asm/comp/struct-resource-name-aliasing.asm.comp

@@ -0,0 +1,23 @@
+#pragma clang diagnostic ignored "-Wmissing-prototypes"
+
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+struct bufA
+{
+    uint _data[1];
+};
+
+void _main(device bufA& bufA_1, device bufA& bufB)
+{
+    bufA_1._data[0] = 0u;
+    bufB._data[0] = 0u;
+}
+
+kernel void main0(device bufA& bufA_1 [[buffer(0)]], device bufA& bufB [[buffer(1)]])
+{
+    _main(bufA_1, bufB);
+}
+

shaders-msl/asm/comp/struct-resource-name-aliasing.asm.comp

@@ -0,0 +1,47 @@
+; SPIR-V
+; Version: 1.0
+; Generator: Khronos Glslang Reference Front End; 7
+; Bound: 21
+; Schema: 0
+               OpCapability Shader
+          %1 = OpExtInstImport "GLSL.std.450"
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 8 8 1
+               OpSource HLSL 500
+               OpName %main "main"
+               OpName %_main_ "@main("
+               OpName %bufA "bufA"
+               OpMemberName %bufA 0 "@data"
+               OpName %bufA_0 "bufA"
+               OpName %bufB "bufB"
+               OpDecorate %_runtimearr_uint ArrayStride 4
+               OpMemberDecorate %bufA 0 Offset 0
+               OpDecorate %bufA BufferBlock
+               OpDecorate %bufA_0 DescriptorSet 0
+               OpDecorate %bufB DescriptorSet 0
+       %void = OpTypeVoid
+          %3 = OpTypeFunction %void
+       %uint = OpTypeInt 32 0
+%_runtimearr_uint = OpTypeRuntimeArray %uint
+       %bufA = OpTypeStruct %_runtimearr_uint
+%_ptr_Uniform_bufA = OpTypePointer Uniform %bufA
+     %bufA_0 = OpVariable %_ptr_Uniform_bufA Uniform
+        %int = OpTypeInt 32 1
+      %int_0 = OpConstant %int 0
+     %uint_0 = OpConstant %uint 0
+%_ptr_Uniform_uint = OpTypePointer Uniform %uint
+       %bufB = OpVariable %_ptr_Uniform_bufA Uniform
+       %main = OpFunction %void None %3
+          %5 = OpLabel
+         %20 = OpFunctionCall %void %_main_
+               OpReturn
+               OpFunctionEnd
+     %_main_ = OpFunction %void None %3
+          %7 = OpLabel
+         %17 = OpAccessChain %_ptr_Uniform_uint %bufA_0 %int_0 %int_0
+               OpStore %17 %uint_0
+         %19 = OpAccessChain %_ptr_Uniform_uint %bufB %int_0 %int_0
+               OpStore %19 %uint_0
+               OpReturn
+               OpFunctionEnd

spirv_msl.cpp

@@ -4424,6 +4424,7 @@ string CompilerMSL::entry_point_args(bool append_comma)
 		if (!ep_args.empty())
 			ep_args += ", ";
 
+		add_resource_name(var.self);
 		ep_args += type_to_glsl(type) + " " + to_name(var.self) + " [[stage_in]]";
 	}
 
@@ -4452,6 +4453,7 @@ string CompilerMSL::entry_point_args(bool append_comma)
 		{
 			if (type.basetype == SPIRType::SampledImage)
 			{
+				add_resource_name(var_id);
 				resources.push_back(
 				    { &id, to_name(var_id), SPIRType::Image, get_metal_resource_index(var, SPIRType::Image) });
 
@@ -4464,6 +4466,7 @@ string CompilerMSL::entry_point_args(bool append_comma)
 			else if (constexpr_samplers.count(var_id) == 0)
 			{
 				// constexpr samplers are not declared as resources.
+				add_resource_name(var_id);
 				resources.push_back(
 				    { &id, to_name(var_id), type.basetype, get_metal_resource_index(var, type.basetype) });
 			}