MSL: Fix emission of bindless helper template for bindless SSBO.

reference/shaders-msl-no-opt/comp/simple-bindless-ssbo.msl2.argument.argument-tier-1.device-argument-buffer.comp

@@ -0,0 +1,45 @@
+#pragma clang diagnostic ignored "-Wmissing-prototypes"
+
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+template<typename T>
+struct spvDescriptor
+{
+    T value;
+};
+
+template<typename T>
+struct spvDescriptorArray
+{
+    spvDescriptorArray(const device spvDescriptor<T>* ptr) : ptr(&ptr->value)
+    {
+    }
+    const device T& operator [] (size_t i) const
+    {
+        return ptr[i];
+    }
+    const device T* ptr;
+};
+
+struct SSBO
+{
+    float4 a;
+};
+
+constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u);
+
+struct spvDescriptorSetBuffer0
+{
+    spvDescriptor<device SSBO *> ssbos [[id(0)]][1] /* unsized array hack */;
+};
+
+kernel void main0(const device spvDescriptorSetBuffer0& spvDescriptorSet0 [[buffer(0)]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]])
+{
+    spvDescriptorArray<device SSBO*> ssbos {spvDescriptorSet0.ssbos};
+
+    ssbos[gl_WorkGroupID.x]->a += float4(2.0);
+}
+

shaders-msl-no-opt/comp/simple-bindless-ssbo.msl2.argument.argument-tier-1.device-argument-buffer.comp

@@ -0,0 +1,13 @@
+#version 450
+#extension GL_EXT_nonuniform_qualifier : require
+layout(local_size_x = 1) in;
+
+layout(set = 0, binding = 0) buffer SSBO
+{
+	vec4 a;
+} ssbos[];
+
+void main()
+{
+	ssbos[gl_WorkGroupID.x].a += 2.0;
+}

spirv_msl.cpp

@@ -13874,6 +13874,7 @@ void CompilerMSL::entry_point_args_discrete_descriptors(string &ep_args)
 					}
 					else
 					{
+						add_spv_func_and_recompile(SPVFuncImplVariableDescriptor);
 						ep_args += "const device spvDescriptor<" + get_argument_address_space(var) + " " +
 						           type_to_glsl(type) + "*>* ";
 					}