From 1b97ecd2171cdd5551099ab8e06c8625cb3b8a2a Mon Sep 17 00:00:00 2001
From: Hans-Kristian Arntzen <post@arntzen-software.no>
Date: Mon, 14 Oct 2024 12:19:55 +0200
Subject: [PATCH] MSL: Fix emission of bindless helper template for bindless
 SSBO.

---
 ...rgument-tier-1.device-argument-buffer.comp | 45 +++++++++++++++++++
 ...rgument-tier-1.device-argument-buffer.comp | 13 ++++++
 spirv_msl.cpp                                 |  1 +
 3 files changed, 59 insertions(+)
 create mode 100644 reference/shaders-msl-no-opt/comp/simple-bindless-ssbo.msl2.argument.argument-tier-1.device-argument-buffer.comp
 create mode 100644 shaders-msl-no-opt/comp/simple-bindless-ssbo.msl2.argument.argument-tier-1.device-argument-buffer.comp

diff --git a/reference/shaders-msl-no-opt/comp/simple-bindless-ssbo.msl2.argument.argument-tier-1.device-argument-buffer.comp b/reference/shaders-msl-no-opt/comp/simple-bindless-ssbo.msl2.argument.argument-tier-1.device-argument-buffer.comp
new file mode 100644
index 000000000..d8107ec47
--- /dev/null
+++ b/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);
+}
+
diff --git a/shaders-msl-no-opt/comp/simple-bindless-ssbo.msl2.argument.argument-tier-1.device-argument-buffer.comp b/shaders-msl-no-opt/comp/simple-bindless-ssbo.msl2.argument.argument-tier-1.device-argument-buffer.comp
new file mode 100644
index 000000000..0db56342c
--- /dev/null
+++ b/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;
+}
diff --git a/spirv_msl.cpp b/spirv_msl.cpp
index 06552be92..9e30be864 100644
--- a/spirv_msl.cpp
+++ b/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) + "*>* ";
 					}