Merge branch 'KhronosGroup:main' into SPV_QCOM_image_processing

CMakeLists.txt

@@ -243,7 +243,7 @@ set(spirv-cross-util-sources
 		${CMAKE_CURRENT_SOURCE_DIR}/spirv_cross_util.hpp)
 
 set(spirv-cross-abi-major 0)
-set(spirv-cross-abi-minor 58)
+set(spirv-cross-abi-minor 59)
 set(spirv-cross-abi-patch 0)
 set(SPIRV_CROSS_VERSION ${spirv-cross-abi-major}.${spirv-cross-abi-minor}.${spirv-cross-abi-patch})
 

reference/opt/shaders-msl/comp/argument-buffers-runtime-array-buffer.argument.device-argument-buffer.argument-tier-1.msl2.comp

@@ -0,0 +1,77 @@
+#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)
+    {
+    }
+    const device T& operator [] (size_t i) const
+    {
+        return ptr[i].value;
+    }
+    const device spvDescriptor<T>* ptr;
+};
+
+struct SSBO
+{
+    float4 v[1];
+};
+
+struct UBO
+{
+    float4 v[1024];
+};
+
+struct UBOs
+{
+    float4 v;
+};
+
+struct SSBOIn
+{
+    float4 v[1024];
+};
+
+struct SSBOIns
+{
+    float4 v;
+};
+
+constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(64u, 1u, 1u);
+
+struct spvDescriptorSetBuffer0
+{
+    device SSBO* o [[id(0)]];
+    constant UBO* v [[id(1)]];
+    spvDescriptor<constant UBOs *> vs [[id(2)]][1] /* unsized array hack */;
+};
+
+struct spvDescriptorSetBuffer1
+{
+    device SSBOIn* w [[id(0)]];
+    spvDescriptor<device SSBOIns *> ws [[id(1)]][1] /* unsized array hack */;
+};
+
+kernel void main0(const device spvDescriptorSetBuffer0& spvDescriptorSet0 [[buffer(0)]], const device spvDescriptorSetBuffer1& spvDescriptorSet1 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]])
+{
+    spvDescriptorArray<constant UBOs*> vs {spvDescriptorSet0.vs};
+    spvDescriptorArray<device SSBOIns*> ws {spvDescriptorSet1.ws};
+
+    (*spvDescriptorSet0.o).v[gl_GlobalInvocationID.x] = (*spvDescriptorSet0.v).v[gl_WorkGroupID.x];
+    (*spvDescriptorSet0.o).v[gl_GlobalInvocationID.x] = vs[gl_WorkGroupID.x]->v;
+    (*spvDescriptorSet0.o).v[gl_GlobalInvocationID.x] = (*spvDescriptorSet1.w).v[gl_WorkGroupID.x];
+    (*spvDescriptorSet0.o).v[gl_GlobalInvocationID.x] = ws[gl_WorkGroupID.x]->v;
+}
+

reference/opt/shaders-msl/comp/argument-buffers-runtime-array-buffer.rich-descriptor.argument.device-argument-buffer.argument-tier-1.msl2.comp

@@ -0,0 +1,109 @@
+#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 spvBufferDescriptor
+{
+    T value;
+    int length;
+    const device T& operator -> () const device
+    {
+        return value;
+    }
+    const device T& operator * () const device
+    {
+        return value;
+    }
+};
+
+template<typename T>
+struct spvDescriptorArray
+{
+    spvDescriptorArray(const device spvDescriptor<T>* ptr) : ptr(ptr)
+    {
+    }
+    const device T& operator [] (size_t i) const
+    {
+        return ptr[i].value;
+    }
+    const device spvDescriptor<T>* ptr;
+};
+
+template<typename T>
+struct spvDescriptorArray<device T*>
+{
+    spvDescriptorArray(const device spvBufferDescriptor<device T*>* ptr) : ptr(ptr)
+    {
+    }
+    const device T* operator [] (size_t i) const
+    {
+        return ptr[i].value;
+    }
+    const int length(int i) const
+    {
+        return ptr[i].length;
+    }
+    const device spvBufferDescriptor<device T*>* ptr;
+};
+
+struct SSBO
+{
+    float4 v[1];
+};
+
+struct UBO
+{
+    float4 v[1024];
+};
+
+struct UBOs
+{
+    float4 v;
+};
+
+struct SSBOIn
+{
+    float4 v[1024];
+};
+
+struct SSBOIns
+{
+    float4 v;
+};
+
+constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(64u, 1u, 1u);
+
+struct spvDescriptorSetBuffer0
+{
+    device SSBO* o [[id(0)]];
+    constant UBO* v [[id(1)]];
+    spvDescriptor<constant UBOs *> vs [[id(2)]][1] /* unsized array hack */;
+};
+
+struct spvDescriptorSetBuffer1
+{
+    device SSBOIn* w [[id(0)]];
+    spvBufferDescriptor<device SSBOIns *> ws [[id(1)]][1] /* unsized array hack */;
+};
+
+kernel void main0(const device spvDescriptorSetBuffer0& spvDescriptorSet0 [[buffer(0)]], const device spvDescriptorSetBuffer1& spvDescriptorSet1 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]])
+{
+    spvDescriptorArray<constant UBOs*> vs {spvDescriptorSet0.vs};
+    spvDescriptorArray<device SSBOIns*> ws {spvDescriptorSet1.ws};
+
+    (*spvDescriptorSet0.o).v[gl_GlobalInvocationID.x] = (*spvDescriptorSet0.v).v[gl_WorkGroupID.x];
+    (*spvDescriptorSet0.o).v[gl_GlobalInvocationID.x] = vs[gl_WorkGroupID.x]->v;
+    (*spvDescriptorSet0.o).v[gl_GlobalInvocationID.x] = (*spvDescriptorSet1.w).v[gl_WorkGroupID.x];
+    (*spvDescriptorSet0.o).v[gl_GlobalInvocationID.x] = ws[gl_WorkGroupID.x]->v;
+}
+

reference/opt/shaders-msl/comp/argument-buffers-runtime-array.argument.device-argument-buffer.argument-tier-1.msl2.comp

@@ -0,0 +1,55 @@
+#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)
+    {
+    }
+    const device T& operator [] (size_t i) const
+    {
+        return ptr[i].value;
+    }
+    const device spvDescriptor<T>* ptr;
+};
+
+struct SSBO
+{
+    float4 v[1];
+};
+
+constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(64u, 1u, 1u);
+
+struct spvDescriptorSetBuffer0
+{
+    device SSBO* m_13 [[id(0)]];
+    texture2d<float> T [[id(1)]];
+    spvDescriptor<texture2d<float>> Ts [[id(2)]][1] /* unsized array hack */;
+};
+
+struct spvDescriptorSetBuffer1
+{
+    sampler S [[id(0)]];
+    spvDescriptor<sampler> Ss [[id(1)]][1] /* unsized array hack */;
+};
+
+kernel void main0(const device spvDescriptorSetBuffer0& spvDescriptorSet0 [[buffer(0)]], const device spvDescriptorSetBuffer1& spvDescriptorSet1 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]])
+{
+    spvDescriptorArray<texture2d<float>> Ts {spvDescriptorSet0.Ts};
+    spvDescriptorArray<sampler> Ss {spvDescriptorSet1.Ss};
+
+    (*spvDescriptorSet0.m_13).v[gl_GlobalInvocationID.x] = spvDescriptorSet0.T.sample(spvDescriptorSet1.S, float2(0.5), level(0.0));
+    (*spvDescriptorSet0.m_13).v[gl_GlobalInvocationID.x] = Ts[gl_WorkGroupID.x].sample(Ss[gl_WorkGroupID.x], float2(0.5), level(0.0));
+}
+

reference/shaders-msl/comp/argument-buffers-runtime-array-buffer.argument.device-argument-buffer.argument-tier-1.msl2.comp

@@ -0,0 +1,83 @@
+#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)
+    {
+    }
+    const device T& operator [] (size_t i) const
+    {
+        return ptr[i].value;
+    }
+    const device spvDescriptor<T>* ptr;
+};
+
+struct SSBO
+{
+    float4 v[1];
+};
+
+struct UBO
+{
+    float4 v[1024];
+};
+
+struct UBOs
+{
+    float4 v;
+};
+
+struct SSBOIn
+{
+    float4 v[1024];
+};
+
+struct SSBOIns
+{
+    float4 v;
+};
+
+constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(64u, 1u, 1u);
+
+struct spvDescriptorSetBuffer0
+{
+    device SSBO* o [[id(0)]];
+    constant UBO* v [[id(1)]];
+    spvDescriptor<constant UBOs *> vs [[id(2)]][1] /* unsized array hack */;
+};
+
+struct spvDescriptorSetBuffer1
+{
+    device SSBOIn* w [[id(0)]];
+    spvDescriptor<device SSBOIns *> ws [[id(1)]][1] /* unsized array hack */;
+};
+
+static inline __attribute__((always_inline))
+void in_func(device SSBO& o, thread uint3& gl_GlobalInvocationID, constant UBO& v, thread uint3& gl_WorkGroupID, const spvDescriptorArray<constant UBOs*> vs, device SSBOIn& w, const spvDescriptorArray<device SSBOIns*> ws)
+{
+    o.v[gl_GlobalInvocationID.x] = v.v[gl_WorkGroupID.x];
+    o.v[gl_GlobalInvocationID.x] = vs[gl_WorkGroupID.x]->v;
+    o.v[gl_GlobalInvocationID.x] = w.v[gl_WorkGroupID.x];
+    o.v[gl_GlobalInvocationID.x] = ws[gl_WorkGroupID.x]->v;
+}
+
+kernel void main0(const device spvDescriptorSetBuffer0& spvDescriptorSet0 [[buffer(0)]], const device spvDescriptorSetBuffer1& spvDescriptorSet1 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]])
+{
+    spvDescriptorArray<constant UBOs*> vs {spvDescriptorSet0.vs};
+    spvDescriptorArray<device SSBOIns*> ws {spvDescriptorSet1.ws};
+
+    in_func((*spvDescriptorSet0.o), gl_GlobalInvocationID, (*spvDescriptorSet0.v), gl_WorkGroupID, vs, (*spvDescriptorSet1.w), ws);
+}
+

reference/shaders-msl/comp/argument-buffers-runtime-array-buffer.rich-descriptor.argument.device-argument-buffer.argument-tier-1.msl2.comp

@@ -0,0 +1,115 @@
+#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 spvBufferDescriptor
+{
+    T value;
+    int length;
+    const device T& operator -> () const device
+    {
+        return value;
+    }
+    const device T& operator * () const device
+    {
+        return value;
+    }
+};
+
+template<typename T>
+struct spvDescriptorArray
+{
+    spvDescriptorArray(const device spvDescriptor<T>* ptr) : ptr(ptr)
+    {
+    }
+    const device T& operator [] (size_t i) const
+    {
+        return ptr[i].value;
+    }
+    const device spvDescriptor<T>* ptr;
+};
+
+template<typename T>
+struct spvDescriptorArray<device T*>
+{
+    spvDescriptorArray(const device spvBufferDescriptor<device T*>* ptr) : ptr(ptr)
+    {
+    }
+    const device T* operator [] (size_t i) const
+    {
+        return ptr[i].value;
+    }
+    const int length(int i) const
+    {
+        return ptr[i].length;
+    }
+    const device spvBufferDescriptor<device T*>* ptr;
+};
+
+struct SSBO
+{
+    float4 v[1];
+};
+
+struct UBO
+{
+    float4 v[1024];
+};
+
+struct UBOs
+{
+    float4 v;
+};
+
+struct SSBOIn
+{
+    float4 v[1024];
+};
+
+struct SSBOIns
+{
+    float4 v;
+};
+
+constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(64u, 1u, 1u);
+
+struct spvDescriptorSetBuffer0
+{
+    device SSBO* o [[id(0)]];
+    constant UBO* v [[id(1)]];
+    spvDescriptor<constant UBOs *> vs [[id(2)]][1] /* unsized array hack */;
+};
+
+struct spvDescriptorSetBuffer1
+{
+    device SSBOIn* w [[id(0)]];
+    spvBufferDescriptor<device SSBOIns *> ws [[id(1)]][1] /* unsized array hack */;
+};
+
+static inline __attribute__((always_inline))
+void in_func(device SSBO& o, thread uint3& gl_GlobalInvocationID, constant UBO& v, thread uint3& gl_WorkGroupID, const spvDescriptorArray<constant UBOs*> vs, device SSBOIn& w, const spvDescriptorArray<device SSBOIns*> ws)
+{
+    o.v[gl_GlobalInvocationID.x] = v.v[gl_WorkGroupID.x];
+    o.v[gl_GlobalInvocationID.x] = vs[gl_WorkGroupID.x]->v;
+    o.v[gl_GlobalInvocationID.x] = w.v[gl_WorkGroupID.x];
+    o.v[gl_GlobalInvocationID.x] = ws[gl_WorkGroupID.x]->v;
+}
+
+kernel void main0(const device spvDescriptorSetBuffer0& spvDescriptorSet0 [[buffer(0)]], const device spvDescriptorSetBuffer1& spvDescriptorSet1 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]])
+{
+    spvDescriptorArray<constant UBOs*> vs {spvDescriptorSet0.vs};
+    spvDescriptorArray<device SSBOIns*> ws {spvDescriptorSet1.ws};
+
+    in_func((*spvDescriptorSet0.o), gl_GlobalInvocationID, (*spvDescriptorSet0.v), gl_WorkGroupID, vs, (*spvDescriptorSet1.w), ws);
+}
+

reference/shaders-msl/comp/argument-buffers-runtime-array.argument.device-argument-buffer.argument-tier-1.msl2.comp

@@ -0,0 +1,61 @@
+#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)
+    {
+    }
+    const device T& operator [] (size_t i) const
+    {
+        return ptr[i].value;
+    }
+    const device spvDescriptor<T>* ptr;
+};
+
+struct SSBO
+{
+    float4 v[1];
+};
+
+constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(64u, 1u, 1u);
+
+struct spvDescriptorSetBuffer0
+{
+    device SSBO* m_13 [[id(0)]];
+    texture2d<float> T [[id(1)]];
+    spvDescriptor<texture2d<float>> Ts [[id(2)]][1] /* unsized array hack */;
+};
+
+struct spvDescriptorSetBuffer1
+{
+    sampler S [[id(0)]];
+    spvDescriptor<sampler> Ss [[id(1)]][1] /* unsized array hack */;
+};
+
+static inline __attribute__((always_inline))
+void in_func(device SSBO& _13, thread uint3& gl_GlobalInvocationID, texture2d<float> T, sampler S, const spvDescriptorArray<texture2d<float>> Ts, thread uint3& gl_WorkGroupID, const spvDescriptorArray<sampler> Ss)
+{
+    _13.v[gl_GlobalInvocationID.x] = T.sample(S, float2(0.5), level(0.0));
+    _13.v[gl_GlobalInvocationID.x] = Ts[gl_WorkGroupID.x].sample(Ss[gl_WorkGroupID.x], float2(0.5), level(0.0));
+}
+
+kernel void main0(const device spvDescriptorSetBuffer0& spvDescriptorSet0 [[buffer(0)]], const device spvDescriptorSetBuffer1& spvDescriptorSet1 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]])
+{
+    spvDescriptorArray<texture2d<float>> Ts {spvDescriptorSet0.Ts};
+    spvDescriptorArray<sampler> Ss {spvDescriptorSet1.Ss};
+
+    in_func((*spvDescriptorSet0.m_13), gl_GlobalInvocationID, spvDescriptorSet0.T, spvDescriptorSet1.S, Ts, gl_WorkGroupID, Ss);
+}
+

shaders-msl/comp/argument-buffers-runtime-array-buffer.argument.device-argument-buffer.argument-tier-1.msl2.comp

@@ -0,0 +1,28 @@
+#version 450
+#extension GL_EXT_nonuniform_qualifier : require
+layout(local_size_x = 64) in;
+
+layout(set = 0, binding = 0) buffer SSBO
+{
+	vec4 v[];
+} o;
+
+layout(set = 0, binding = 1) uniform UBO { vec4 v[1024]; } v;
+layout(set = 0, binding = 2) uniform UBOs { vec4 v; } vs[];
+
+layout(set = 1, binding = 0) buffer SSBOIn { vec4 v[1024]; } w;
+layout(set = 1, binding = 1) buffer SSBOIns { vec4 v; } ws[];
+
+void in_func()
+{
+	o.v[gl_GlobalInvocationID.x] = v.v[gl_WorkGroupID.x];
+	o.v[gl_GlobalInvocationID.x] = vs[gl_WorkGroupID.x].v;
+
+	o.v[gl_GlobalInvocationID.x] = w.v[gl_WorkGroupID.x];
+	o.v[gl_GlobalInvocationID.x] = ws[gl_WorkGroupID.x].v;
+}
+
+void main()
+{
+	in_func();
+}

shaders-msl/comp/argument-buffers-runtime-array-buffer.rich-descriptor.argument.device-argument-buffer.argument-tier-1.msl2.comp

@@ -0,0 +1,28 @@
+#version 450
+#extension GL_EXT_nonuniform_qualifier : require
+layout(local_size_x = 64) in;
+
+layout(set = 0, binding = 0) buffer SSBO
+{
+	vec4 v[];
+} o;
+
+layout(set = 0, binding = 1) uniform UBO { vec4 v[1024]; } v;
+layout(set = 0, binding = 2) uniform UBOs { vec4 v; } vs[];
+
+layout(set = 1, binding = 0) buffer SSBOIn { vec4 v[1024]; } w;
+layout(set = 1, binding = 1) buffer SSBOIns { vec4 v; } ws[];
+
+void in_func()
+{
+	o.v[gl_GlobalInvocationID.x] = v.v[gl_WorkGroupID.x];
+	o.v[gl_GlobalInvocationID.x] = vs[gl_WorkGroupID.x].v;
+
+	o.v[gl_GlobalInvocationID.x] = w.v[gl_WorkGroupID.x];
+	o.v[gl_GlobalInvocationID.x] = ws[gl_WorkGroupID.x].v;
+}
+
+void main()
+{
+	in_func();
+}

shaders-msl/comp/argument-buffers-runtime-array.argument.device-argument-buffer.argument-tier-1.msl2.comp

@@ -0,0 +1,25 @@
+#version 450
+#extension GL_EXT_nonuniform_qualifier : require
+layout(local_size_x = 64) in;
+
+layout(set = 0, binding = 0) buffer SSBO
+{
+	vec4 v[];
+};
+
+layout(set = 0, binding = 1) uniform texture2D T;
+layout(set = 0, binding = 2) uniform texture2D Ts[];
+
+layout(set = 1, binding = 0) uniform sampler S;
+layout(set = 1, binding = 1) uniform sampler Ss[];
+
+void in_func()
+{
+	v[gl_GlobalInvocationID.x] = textureLod(sampler2D(T, S), vec2(0.5), 0.0);
+	v[gl_GlobalInvocationID.x] = textureLod(sampler2D(Ts[gl_WorkGroupID.x], Ss[gl_WorkGroupID.x]), vec2(0.5), 0.0);
+}
+
+void main()
+{
+	in_func();
+}

spirv_common.hpp

@@ -1696,6 +1696,7 @@ struct Meta
 		uint32_t index = 0;
 		spv::FPRoundingMode fp_rounding_mode = spv::FPRoundingModeMax;
 		bool builtin = false;
+		bool qualified_alias_explicit_override = false;
 
 		struct Extended
 		{

spirv_cross_c.cpp

@@ -2552,6 +2552,16 @@ int spvc_constant_get_scalar_i32(spvc_constant constant, unsigned column, unsign
 	return constant->scalar_i32(column, row);
 }
 
+unsigned long long spvc_constant_get_scalar_u64(spvc_constant constant, unsigned column, unsigned row)
+{
+	return constant->scalar_u64(column, row);
+}
+
+long long spvc_constant_get_scalar_i64(spvc_constant constant, unsigned column, unsigned row)
+{
+	return constant->scalar_i64(column, row);
+}
+
 unsigned spvc_constant_get_scalar_u16(spvc_constant constant, unsigned column, unsigned row)
 {
 	return constant->scalar_u16(column, row);
@@ -2609,6 +2619,16 @@ void spvc_constant_set_scalar_i32(spvc_constant constant, unsigned column, unsig
 	constant->m.c[column].r[row].i32 = value;
 }
 
+void spvc_constant_set_scalar_u64(spvc_constant constant, unsigned column, unsigned row, unsigned long long value)
+{
+	constant->m.c[column].r[row].u64 = value;
+}
+
+void spvc_constant_set_scalar_i64(spvc_constant constant, unsigned column, unsigned row, long long value)
+{
+	constant->m.c[column].r[row].i64 = value;
+}
+
 void spvc_constant_set_scalar_u16(spvc_constant constant, unsigned column, unsigned row, unsigned short value)
 {
 	constant->m.c[column].r[row].u32 = uint32_t(value);

spirv_cross_c.h

@@ -40,7 +40,7 @@ extern "C" {
 /* Bumped if ABI or API breaks backwards compatibility. */
 #define SPVC_C_API_VERSION_MAJOR 0
 /* Bumped if APIs or enumerations are added in a backwards compatible way. */
-#define SPVC_C_API_VERSION_MINOR 58
+#define SPVC_C_API_VERSION_MINOR 59
 /* Bumped if internal implementation details change. */
 #define SPVC_C_API_VERSION_PATCH 0
 
@@ -1052,6 +1052,8 @@ SPVC_PUBLIC_API int spvc_constant_get_scalar_i16(spvc_constant constant, unsigne
 SPVC_PUBLIC_API unsigned spvc_constant_get_scalar_u8(spvc_constant constant, unsigned column, unsigned row);
 SPVC_PUBLIC_API int spvc_constant_get_scalar_i8(spvc_constant constant, unsigned column, unsigned row);
 SPVC_PUBLIC_API void spvc_constant_get_subconstants(spvc_constant constant, const spvc_constant_id **constituents, size_t *count);
+SPVC_PUBLIC_API unsigned long long spvc_constant_get_scalar_u64(spvc_constant constant, unsigned column, unsigned row);
+SPVC_PUBLIC_API long long spvc_constant_get_scalar_i64(spvc_constant constant, unsigned column, unsigned row);
 SPVC_PUBLIC_API spvc_type_id spvc_constant_get_type(spvc_constant constant);
 
 /*
@@ -1062,6 +1064,8 @@ SPVC_PUBLIC_API void spvc_constant_set_scalar_fp32(spvc_constant constant, unsig
 SPVC_PUBLIC_API void spvc_constant_set_scalar_fp64(spvc_constant constant, unsigned column, unsigned row, double value);
 SPVC_PUBLIC_API void spvc_constant_set_scalar_u32(spvc_constant constant, unsigned column, unsigned row, unsigned value);
 SPVC_PUBLIC_API void spvc_constant_set_scalar_i32(spvc_constant constant, unsigned column, unsigned row, int value);
+SPVC_PUBLIC_API void spvc_constant_set_scalar_u64(spvc_constant constant, unsigned column, unsigned row, unsigned long long value);
+SPVC_PUBLIC_API void spvc_constant_set_scalar_i64(spvc_constant constant, unsigned column, unsigned row, long long value);
 SPVC_PUBLIC_API void spvc_constant_set_scalar_u16(spvc_constant constant, unsigned column, unsigned row, unsigned short value);
 SPVC_PUBLIC_API void spvc_constant_set_scalar_i16(spvc_constant constant, unsigned column, unsigned row, signed short value);
 SPVC_PUBLIC_API void spvc_constant_set_scalar_u8(spvc_constant constant, unsigned column, unsigned row, unsigned char value);

spirv_glsl.cpp

@@ -15543,7 +15543,7 @@ string CompilerGLSL::type_to_array_glsl(const SPIRType &type)
 	}
 }
 
-string CompilerGLSL::image_type_glsl(const SPIRType &type, uint32_t id)
+string CompilerGLSL::image_type_glsl(const SPIRType &type, uint32_t id, bool /*member*/)
 {
 	auto &imagetype = get<SPIRType>(type.image.type);
 	string res;

spirv_glsl.hpp

@@ -438,7 +438,7 @@ protected:
 	virtual void emit_struct_member(const SPIRType &type, uint32_t member_type_id, uint32_t index,
 	                                const std::string &qualifier = "", uint32_t base_offset = 0);
 	virtual void emit_struct_padding_target(const SPIRType &type);
-	virtual std::string image_type_glsl(const SPIRType &type, uint32_t id = 0);
+	virtual std::string image_type_glsl(const SPIRType &type, uint32_t id = 0, bool member = false);
 	std::string constant_expression(const SPIRConstant &c,
 	                                bool inside_block_like_struct_scope = false,
 	                                bool inside_struct_scope = false);

spirv_msl.cpp

@@ -1379,20 +1379,28 @@ void CompilerMSL::emit_entry_point_declarations()
 			{
 				SPIRV_CROSS_THROW("Unsized array of descriptors requires argument buffer tier 2");
 			}
+
+			string resource_name;
+			if (descriptor_set_is_argument_buffer(get_decoration(var.self, DecorationDescriptorSet)))
+				resource_name = ir.meta[var.self].decoration.qualified_alias;
+			else
+				resource_name = name + "_";
+
 			switch (type.basetype)
 			{
 			case SPIRType::Image:
 			case SPIRType::Sampler:
 			case SPIRType::AccelerationStructure:
-				statement("spvDescriptorArray<", type_to_glsl(buffer_type), "> ", name, " {", name, "_};");
+				statement("spvDescriptorArray<", type_to_glsl(buffer_type), "> ", name, " {", resource_name, "};");
 				break;
 			case SPIRType::SampledImage:
-				statement("spvDescriptorArray<", type_to_glsl(buffer_type), "> ", name, " {", name, "_};");
+				statement("spvDescriptorArray<", type_to_glsl(buffer_type), "> ", name, " {", resource_name, "};");
+				// Unsupported with argument buffer for now.
 				statement("spvDescriptorArray<sampler> ", name, "Smplr {", name, "Smplr_};");
 				break;
 			case SPIRType::Struct:
 				statement("spvDescriptorArray<", get_argument_address_space(var), " ", type_to_glsl(buffer_type), "*> ",
-				          name, " {", name, "_};");
+				          name, " {", resource_name, "};");
 				break;
 			default:
 				break;
@@ -10665,7 +10673,7 @@ void CompilerMSL::emit_function_prototype(SPIRFunction &func, const Bitset &)
 			{
 				if (arg_type.array.empty() || (var ? is_var_runtime_size_array(*var) : is_runtime_size_array(arg_type)))
 				{
-					decl += join(", ", sampler_type(arg_type, arg.id), " ", to_sampler_expression(name_id));
+					decl += join(", ", sampler_type(arg_type, arg.id, false), " ", to_sampler_expression(name_id));
 				}
 				else
 				{
@@ -10673,7 +10681,7 @@ void CompilerMSL::emit_function_prototype(SPIRFunction &func, const Bitset &)
 							descriptor_address_space(name_id,
 							                         StorageClassUniformConstant,
 							                         "thread const");
-					decl += join(", ", sampler_address_space, " ", sampler_type(arg_type, name_id), "& ",
+					decl += join(", ", sampler_address_space, " ", sampler_type(arg_type, name_id, false), "& ",
 					             to_sampler_expression(name_id));
 				}
 			}
@@ -12063,6 +12071,20 @@ string CompilerMSL::to_struct_member(const SPIRType &type, uint32_t member_type_
 		array_type = type_to_array_glsl(physical_type);
 	}
 
+	if (orig_id)
+	{
+		auto *data_type = declared_type;
+		if (is_pointer(*data_type))
+			data_type = &get_pointee_type(*data_type);
+
+		if (is_array(*data_type) && get_resource_array_size(*data_type, orig_id) == 0)
+		{
+			// Hack for declaring unsized array of resources. Need to declare dummy sized array by value inline.
+			// This can then be wrapped in spvDescriptorArray as usual.
+			array_type = "[1] /* unsized array hack */";
+		}
+	}
+
 	string decl_type;
 	if (declared_type->vecsize > 4)
 	{
@@ -13283,12 +13305,22 @@ void CompilerMSL::entry_point_args_discrete_descriptors(string &ep_args)
 		    !is_hidden_variable(var))
 		{
 			auto &type = get_variable_data_type(var);
+			uint32_t desc_set = get_decoration(var_id, DecorationDescriptorSet);
 
 			if (is_supported_argument_buffer_type(type) && var.storage != StorageClassPushConstant)
 			{
-				uint32_t desc_set = get_decoration(var_id, DecorationDescriptorSet);
 				if (descriptor_set_is_argument_buffer(desc_set))
+				{
+					if (is_var_runtime_size_array(var))
+					{
+						// Runtime arrays need to be wrapped in spvDescriptorArray from argument buffer payload.
+						entry_point_bindings.push_back(&var);
+						// We'll wrap this, so to_name() will always use non-qualified name.
+						// We'll need the qualified name to create temporary variable instead.
+						ir.meta[var_id].decoration.qualified_alias_explicit_override = true;
+					}
 					return;
+				}
 			}
 
 			// Handle descriptor aliasing. We can handle aliasing of buffers by casting pointers,
@@ -13480,7 +13512,7 @@ void CompilerMSL::entry_point_args_discrete_descriptors(string &ep_args)
 		case SPIRType::Sampler:
 			if (!ep_args.empty())
 				ep_args += ", ";
-			ep_args += sampler_type(type, var_id) + " " + r.name;
+			ep_args += sampler_type(type, var_id, false) + " " + r.name;
 			if (is_var_runtime_size_array(var))
 				ep_args += "_ [[buffer(" + convert_to_string(r.index) + ")]]";
 			else
@@ -13495,7 +13527,7 @@ void CompilerMSL::entry_point_args_discrete_descriptors(string &ep_args)
 			const auto &basetype = get<SPIRType>(var.basetype);
 			if (!type_is_msl_framebuffer_fetch(basetype))
 			{
-				ep_args += image_type_glsl(type, var_id) + " " + r.name;
+				ep_args += image_type_glsl(type, var_id, false) + " " + r.name;
 				if (r.plane > 0)
 					ep_args += join(plane_name_suffix, r.plane);
 
@@ -13512,7 +13544,7 @@ void CompilerMSL::entry_point_args_discrete_descriptors(string &ep_args)
 			{
 				if (msl_options.is_macos() && !msl_options.supports_msl_version(2, 3))
 					SPIRV_CROSS_THROW("Framebuffer fetch on Mac is not supported before MSL 2.3.");
-				ep_args += image_type_glsl(type, var_id) + " " + r.name;
+				ep_args += image_type_glsl(type, var_id, false) + " " + r.name;
 				ep_args += " [[color(" + convert_to_string(r.index) + ")]]";
 			}
 
@@ -14500,11 +14532,15 @@ string CompilerMSL::argument_decl(const SPIRFunction::Parameter &arg)
 		if (!address_space.empty())
 			decl = join(address_space, " ", decl);
 
-		const char *argument_buffer_space = descriptor_address_space(name_id, type_storage, nullptr);
+		// spvDescriptorArray absorbs the address space inside the template.
-		if (argument_buffer_space)
+		if (!is_var_runtime_size_array(var))
 		{
-			decl += " ";
+			const char *argument_buffer_space = descriptor_address_space(name_id, type_storage, nullptr);
-			decl += argument_buffer_space;
+			if (argument_buffer_space)
+			{
+				decl += " ";
+				decl += argument_buffer_space;
+			}
 		}
 
 		// Special case, need to override the array size here if we're using tess level as an argument.
@@ -14626,7 +14662,7 @@ string CompilerMSL::to_name(uint32_t id, bool allow_alias) const
 	if (current_function && (current_function->self == ir.default_entry_point))
 	{
 		auto *m = ir.find_meta(id);
-		if (m && !m->decoration.qualified_alias.empty())
+		if (m && !m->decoration.qualified_alias_explicit_override && !m->decoration.qualified_alias.empty())
 			return m->decoration.qualified_alias;
 	}
 	return Compiler::to_name(id, allow_alias);
@@ -15049,6 +15085,26 @@ string CompilerMSL::type_to_glsl(const SPIRType &type, uint32_t id, bool member)
 		auto type_address_space = get_type_address_space(type, id);
 		const auto *p_parent_type = &get<SPIRType>(type.parent_type);
 
+		// If we're wrapping buffer descriptors in a spvDescriptorArray, we'll have to handle it as a special case.
+		if (member && id)
+		{
+			auto &var = get<SPIRVariable>(id);
+			if (is_var_runtime_size_array(var) && is_runtime_size_array(*p_parent_type))
+			{
+				const bool ssbo = has_decoration(p_parent_type->self, DecorationBufferBlock);
+				bool buffer_desc =
+						(var.storage == StorageClassStorageBuffer || ssbo) &&
+						msl_options.runtime_array_rich_descriptor;
+
+				const char *wrapper_type = buffer_desc ? "spvBufferDescriptor" : "spvDescriptor";
+				add_spv_func_and_recompile(SPVFuncImplVariableDescriptorArray);
+				add_spv_func_and_recompile(buffer_desc ? SPVFuncImplVariableSizedDescriptor : SPVFuncImplVariableDescriptor);
+
+				type_name = join(wrapper_type, "<", type_address_space, " ", type_to_glsl(*p_parent_type, id), " *>");
+				return type_name;
+			}
+		}
+
 		// Work around C pointer qualifier rules. If glsl_type is a pointer type as well
 		// we'll need to emit the address space to the right.
 		// We could always go this route, but it makes the code unnatural.
@@ -15107,10 +15163,10 @@ string CompilerMSL::type_to_glsl(const SPIRType &type, uint32_t id, bool member)
 
 	case SPIRType::Image:
 	case SPIRType::SampledImage:
-		return image_type_glsl(type, id);
+		return image_type_glsl(type, id, member);
 
 	case SPIRType::Sampler:
-		return sampler_type(type, id);
+		return sampler_type(type, id, member);
 
 	case SPIRType::Void:
 		return "void";
@@ -15326,7 +15382,7 @@ std::string CompilerMSL::variable_decl(const SPIRType &type, const std::string &
 	return CompilerGLSL::variable_decl(type, name, id);
 }
 
-std::string CompilerMSL::sampler_type(const SPIRType &type, uint32_t id)
+std::string CompilerMSL::sampler_type(const SPIRType &type, uint32_t id, bool member)
 {
 	auto *var = maybe_get<SPIRVariable>(id);
 	if (var && var->basevariable)
@@ -15345,26 +15401,31 @@ std::string CompilerMSL::sampler_type(const SPIRType &type, uint32_t id)
 
 		// Arrays of samplers in MSL must be declared with a special array<T, N> syntax ala C++11 std::array.
 		// If we have a runtime array, it could be a variable-count descriptor set binding.
+		auto &parent = get<SPIRType>(get_pointee_type(type).parent_type);
 		uint32_t array_size = get_resource_array_size(type, id);
+
 		if (array_size == 0)
 		{
 			add_spv_func_and_recompile(SPVFuncImplVariableDescriptor);
 			add_spv_func_and_recompile(SPVFuncImplVariableDescriptorArray);
-			auto &parent = get<SPIRType>(get_pointee_type(type).parent_type);
-			if (processing_entry_point)
-				return join("const device spvDescriptor<", sampler_type(parent, id), ">*");
-			return join("const spvDescriptorArray<", sampler_type(parent, id), ">");
-		}
 
-		auto &parent = get<SPIRType>(get_pointee_type(type).parent_type);
+			const char *descriptor_wrapper = processing_entry_point ? "const device spvDescriptor" : "const spvDescriptorArray";
-		return join("array<", sampler_type(parent, id), ", ", array_size, ">");
+			if (member)
+				descriptor_wrapper = "spvDescriptor";
+			return join(descriptor_wrapper, "<", sampler_type(parent, id, false), ">",
+			            processing_entry_point ? "*" : "");
+		}
+		else
+		{
+			return join("array<", sampler_type(parent, id, false), ", ", array_size, ">");
+		}
 	}
 	else
 		return "sampler";
 }
 
 // Returns an MSL string describing the SPIR-V image type
-string CompilerMSL::image_type_glsl(const SPIRType &type, uint32_t id)
+string CompilerMSL::image_type_glsl(const SPIRType &type, uint32_t id, bool member)
 {
 	auto *var = maybe_get<SPIRVariable>(id);
 	if (var && var->basevariable)
@@ -15395,17 +15456,29 @@ string CompilerMSL::image_type_glsl(const SPIRType &type, uint32_t id)
 
 		// Arrays of images in MSL must be declared with a special array<T, N> syntax ala C++11 std::array.
 		// If we have a runtime array, it could be a variable-count descriptor set binding.
+		auto &parent = get<SPIRType>(get_pointee_type(type).parent_type);
 		uint32_t array_size = get_resource_array_size(type, id);
+
 		if (array_size == 0)
 		{
 			add_spv_func_and_recompile(SPVFuncImplVariableDescriptor);
 			add_spv_func_and_recompile(SPVFuncImplVariableDescriptorArray);
-			auto &parent = get<SPIRType>(get_pointee_type(type).parent_type);
+			const char *descriptor_wrapper = processing_entry_point ? "const device spvDescriptor" : "const spvDescriptorArray";
-			return join("const device spvDescriptor<", image_type_glsl(parent, id), ">*");
+			if (member)
+			{
+				descriptor_wrapper = "spvDescriptor";
+				// This requires a specialized wrapper type that packs image and sampler side by side.
+				// It is possible in theory.
+				if (type.basetype == SPIRType::SampledImage)
+					SPIRV_CROSS_THROW("Argument buffer runtime array currently not supported for combined image sampler.");
+			}
+			return join(descriptor_wrapper, "<", image_type_glsl(parent, id, false), ">",
+			            processing_entry_point ? "*" : "");
+		}
+		else
+		{
+			return join("array<", image_type_glsl(parent, id, false), ", ", array_size, ">");
 		}
-
-		auto &parent = get<SPIRType>(get_pointee_type(type).parent_type);
-		return join("array<", image_type_glsl(parent, id), ", ", array_size, ">");
 	}
 
 	string img_type_name;
@@ -17842,6 +17915,9 @@ void CompilerMSL::analyze_argument_buffers()
 			auto &var = *resource.var;
 			auto &type = get_variable_data_type(var);
 
+			if (is_var_runtime_size_array(var) && (argument_buffer_device_storage_mask & (1u << desc_set)) == 0)
+				SPIRV_CROSS_THROW("Runtime sized variables must be in device storage argument buffers.");
+
 			// If needed, synthesize and add padding members.
 			// member_index and next_arg_buff_index are incremented when padding members are added.
 			if (msl_options.pad_argument_buffer_resources)

spirv_msl.hpp

@@ -862,8 +862,8 @@ protected:
 	// GCC workaround of lambdas calling protected functions (for older GCC versions)
 	std::string variable_decl(const SPIRType &type, const std::string &name, uint32_t id = 0) override;
 
-	std::string image_type_glsl(const SPIRType &type, uint32_t id = 0) override;
+	std::string image_type_glsl(const SPIRType &type, uint32_t id, bool member) override;
-	std::string sampler_type(const SPIRType &type, uint32_t id);
+	std::string sampler_type(const SPIRType &type, uint32_t id, bool member);
 	std::string builtin_to_glsl(spv::BuiltIn builtin, spv::StorageClass storage) override;
 	std::string to_func_call_arg(const SPIRFunction::Parameter &arg, uint32_t id) override;
 	std::string to_name(uint32_t id, bool allow_alias = true) const override;