MSL: Implement Metal 2.0 indirect argument buffers.

CMakeLists.txt

@@ -195,7 +195,7 @@ endif()
 
 if (SPIRV_CROSS_SHARED)
 	set(spirv-cross-abi-major 0)
-	set(spirv-cross-abi-minor 1)
+	set(spirv-cross-abi-minor 2)
 	set(spirv-cross-abi-patch 0)
 	set(SPIRV_CROSS_VERSION ${spirv-cross-abi-major}.${spirv-cross-abi-minor}.${spirv-cross-abi-patch})
 	set(SPIRV_CROSS_INSTALL_LIB_DIR ${CMAKE_INSTALL_PREFIX}/lib)

main.cpp

@@ -495,6 +495,7 @@ struct CLIArguments
 	bool msl_ios = false;
 	bool msl_pad_fragment_output = false;
 	bool msl_domain_lower_left = false;
+	bool msl_argument_buffers = false;
 	vector<PLSArg> pls_in;
 	vector<PLSArg> pls_out;
 	vector<Remap> remaps;
@@ -552,6 +553,7 @@ static void print_help()
 	                "\t[--msl-ios]\n"
 	                "\t[--msl-pad-fragment-output]\n"
 	                "\t[--msl-domain-lower-left]\n"
+	                "\t[--msl-argument-buffers]\n"
 	                "\t[--hlsl]\n"
 	                "\t[--reflect]\n"
 	                "\t[--shader-model]\n"
@@ -723,6 +725,7 @@ static int main_inner(int argc, char *argv[])
 	cbs.add("--msl-ios", [&args](CLIParser &) { args.msl_ios = true; });
 	cbs.add("--msl-pad-fragment-output", [&args](CLIParser &) { args.msl_pad_fragment_output = true; });
 	cbs.add("--msl-domain-lower-left", [&args](CLIParser &) { args.msl_domain_lower_left = true; });
+	cbs.add("--msl-argument-buffers", [&args](CLIParser &) { args.msl_argument_buffers = true; });
 	cbs.add("--extension", [&args](CLIParser &parser) { args.extensions.push_back(parser.next_string()); });
 	cbs.add("--rename-entry-point", [&args](CLIParser &parser) {
 		auto old_name = parser.next_string();
@@ -855,6 +858,7 @@ static int main_inner(int argc, char *argv[])
 			msl_opts.platform = CompilerMSL::Options::iOS;
 		msl_opts.pad_fragment_output_components = args.msl_pad_fragment_output;
 		msl_opts.tess_domain_origin_lower_left = args.msl_domain_lower_left;
+		msl_opts.argument_buffers = args.msl_argument_buffers;
 		msl_comp->set_msl_options(msl_opts);
 	}
 	else if (args.hlsl)

reference/opt/shaders-msl/comp/argument-buffers-image-load-store.msl2.argument.comp

@@ -0,0 +1,17 @@
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+struct spvDescriptorSetBuffer0
+{
+    texture2d<float, access::write> uImage [[id(1)]];
+    texture2d<float> uImageRead [[id(2)]];
+};
+
+kernel void main0(constant spvDescriptorSetBuffer0& spvDescriptorSet0 [[buffer(0)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
+{
+    int2 _17 = int2(gl_GlobalInvocationID.xy);
+    spvDescriptorSet0.uImage.write(spvDescriptorSet0.uImageRead.read(uint2(_17)), uint2(_17));
+}
+

reference/opt/shaders-msl/frag/argument-buffers.msl2.argument.frag

@@ -0,0 +1,67 @@
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+struct SSBO
+{
+    float4 ssbo;
+};
+
+struct Push
+{
+    float4 push;
+};
+
+struct UBO
+{
+    float4 ubo;
+};
+
+struct UBOs
+{
+    float4 ubo;
+};
+
+struct spvDescriptorSetBuffer0
+{
+    texture2d<float> uTexture [[id(2)]];
+    sampler uTextureSmplr [[id(3)]];
+    constant UBO* m_82 [[id(5)]];
+    array<texture2d<float>, 2> uTextures [[id(6)]];
+    array<sampler, 2> uTexturesSmplr [[id(8)]];
+};
+
+struct spvDescriptorSetBuffer1
+{
+    array<texture2d<float>, 4> uTexture2 [[id(3)]];
+    device SSBO* m_60 [[id(7)]];
+    array<sampler, 2> uSampler [[id(8)]];
+};
+
+struct spvDescriptorSetBuffer2
+{
+    constant UBOs* ubos [[id(4)]][4];
+};
+
+struct main0_out
+{
+    float4 FragColor [[color(0)]];
+};
+
+struct main0_in
+{
+    float2 vUV [[user(locn0)]];
+};
+
+fragment main0_out main0(main0_in in [[stage_in]], constant spvDescriptorSetBuffer0& spvDescriptorSet0 [[buffer(0)]], constant spvDescriptorSetBuffer1& spvDescriptorSet1 [[buffer(1)]], constant spvDescriptorSetBuffer2& spvDescriptorSet2 [[buffer(2)]], constant Push& registers [[buffer(3)]])
+{
+    main0_out out = {};
+    out.FragColor = (((((spvDescriptorSet0.uTexture.sample(spvDescriptorSet0.uTextureSmplr, in.vUV) + spvDescriptorSet1.uTexture2[2].sample(spvDescriptorSet1.uSampler[1], in.vUV)) + spvDescriptorSet0.uTextures[1].sample(spvDescriptorSet0.uTexturesSmplr[1], in.vUV)) + (*spvDescriptorSet1.m_60).ssbo) + registers.push) + (*spvDescriptorSet0.m_82).ubo) + spvDescriptorSet2.ubos[0]->ubo;
+    out.FragColor += (*spvDescriptorSet0.m_82).ubo;
+    out.FragColor += (*spvDescriptorSet1.m_60).ssbo;
+    out.FragColor += spvDescriptorSet2.ubos[1]->ubo;
+    out.FragColor += registers.push;
+    return out;
+}
+

reference/shaders-msl/comp/argument-buffers-image-load-store.msl2.argument.comp

@@ -0,0 +1,17 @@
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+struct spvDescriptorSetBuffer0
+{
+    texture2d<float, access::write> uImage [[id(1)]];
+    texture2d<float> uImageRead [[id(2)]];
+};
+
+kernel void main0(constant spvDescriptorSetBuffer0& spvDescriptorSet0 [[buffer(0)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
+{
+    int2 coord = int2(gl_GlobalInvocationID.xy);
+    spvDescriptorSet0.uImage.write(spvDescriptorSet0.uImageRead.read(uint2(coord)), uint2(coord));
+}
+

reference/shaders-msl/frag/argument-buffers.msl2.argument.frag

@@ -0,0 +1,87 @@
+#pragma clang diagnostic ignored "-Wmissing-prototypes"
+
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+struct SSBO
+{
+    float4 ssbo;
+};
+
+struct Push
+{
+    float4 push;
+};
+
+struct UBO
+{
+    float4 ubo;
+};
+
+struct UBOs
+{
+    float4 ubo;
+};
+
+struct spvDescriptorSetBuffer0
+{
+    texture2d<float> uTexture [[id(2)]];
+    sampler uTextureSmplr [[id(3)]];
+    constant UBO* v_82 [[id(5)]];
+    array<texture2d<float>, 2> uTextures [[id(6)]];
+    array<sampler, 2> uTexturesSmplr [[id(8)]];
+};
+
+struct spvDescriptorSetBuffer1
+{
+    array<texture2d<float>, 4> uTexture2 [[id(3)]];
+    device SSBO* v_60 [[id(7)]];
+    array<sampler, 2> uSampler [[id(8)]];
+};
+
+struct spvDescriptorSetBuffer2
+{
+    constant UBOs* ubos [[id(4)]][4];
+};
+
+struct main0_out
+{
+    float4 FragColor [[color(0)]];
+};
+
+struct main0_in
+{
+    float2 vUV [[user(locn0)]];
+};
+
+float4 sample_in_function2(thread texture2d<float> uTexture, thread const sampler uTextureSmplr, thread float2& vUV, thread const array<texture2d<float>, 4> uTexture2, thread const array<sampler, 2> uSampler, thread const array<texture2d<float>, 2> uTextures, thread const array<sampler, 2> uTexturesSmplr, device SSBO& v_60, constant Push& registers)
+{
+    float4 ret = uTexture.sample(uTextureSmplr, vUV);
+    ret += uTexture2[2].sample(uSampler[1], vUV);
+    ret += uTextures[1].sample(uTexturesSmplr[1], vUV);
+    ret += v_60.ssbo;
+    ret += registers.push;
+    return ret;
+}
+
+float4 sample_in_function(thread texture2d<float> uTexture, thread const sampler uTextureSmplr, thread float2& vUV, thread const array<texture2d<float>, 4> uTexture2, thread const array<sampler, 2> uSampler, thread const array<texture2d<float>, 2> uTextures, thread const array<sampler, 2> uTexturesSmplr, device SSBO& v_60, constant Push& registers, constant UBO& v_82, constant UBOs* (&ubos)[4])
+{
+    float4 ret = sample_in_function2(uTexture, uTextureSmplr, vUV, uTexture2, uSampler, uTextures, uTexturesSmplr, v_60, registers);
+    ret += v_82.ubo;
+    ret += ubos[0]->ubo;
+    return ret;
+}
+
+fragment main0_out main0(main0_in in [[stage_in]], constant spvDescriptorSetBuffer0& spvDescriptorSet0 [[buffer(0)]], constant spvDescriptorSetBuffer1& spvDescriptorSet1 [[buffer(1)]], constant spvDescriptorSetBuffer2& spvDescriptorSet2 [[buffer(2)]], constant Push& registers [[buffer(3)]])
+{
+    main0_out out = {};
+    out.FragColor = sample_in_function(spvDescriptorSet0.uTexture, spvDescriptorSet0.uTextureSmplr, in.vUV, spvDescriptorSet1.uTexture2, spvDescriptorSet1.uSampler, spvDescriptorSet0.uTextures, spvDescriptorSet0.uTexturesSmplr, (*spvDescriptorSet1.v_60), registers, (*spvDescriptorSet0.v_82), spvDescriptorSet2.ubos);
+    out.FragColor += (*spvDescriptorSet0.v_82).ubo;
+    out.FragColor += (*spvDescriptorSet1.v_60).ssbo;
+    out.FragColor += spvDescriptorSet2.ubos[1]->ubo;
+    out.FragColor += registers.push;
+    return out;
+}
+

shaders-msl/comp/argument-buffers-image-load-store.msl2.argument.comp

@@ -0,0 +1,10 @@
+#version 450
+
+layout(set = 0, binding = 1, r32f) writeonly uniform image2D uImage;
+layout(set = 0, binding = 2, r32f) readonly uniform image2D uImageRead;
+
+void main()
+{
+   ivec2 coord = ivec2(gl_GlobalInvocationID.xy);
+   imageStore(uImage, coord, imageLoad(uImageRead, coord));
+}

shaders-msl/frag/argument-buffers.msl2.argument.frag

@@ -0,0 +1,55 @@
+#version 450
+
+layout(std430, push_constant) uniform Push
+{
+   vec4 push;
+} registers;
+
+layout(std140, set = 0, binding = 5) uniform UBO
+{
+   vec4 ubo;
+};
+
+layout(std430, set = 1, binding = 7) buffer SSBO
+{
+   vec4 ssbo;
+};
+
+layout(std140, set = 2, binding = 4) uniform UBOs
+{
+   vec4 ubo;
+} ubos[4];
+
+layout(set = 0, binding = 2) uniform sampler2D uTexture;
+layout(set = 0, binding = 6) uniform sampler2D uTextures[2];
+layout(set = 1, binding = 3) uniform texture2D uTexture2[4];
+layout(set = 1, binding = 8) uniform sampler uSampler[2];
+layout(location = 0) in vec2 vUV;
+layout(location = 0) out vec4 FragColor;
+
+vec4 sample_in_function2()
+{
+   vec4 ret = texture(uTexture, vUV);
+   ret += texture(sampler2D(uTexture2[2], uSampler[1]), vUV);
+   ret += texture(uTextures[1], vUV);
+   ret += ssbo;
+   ret += registers.push;
+   return ret;
+}
+
+vec4 sample_in_function()
+{
+   vec4 ret = sample_in_function2();
+   ret += ubo;
+   ret += ubos[0].ubo;
+   return ret;
+}
+
+void main()
+{
+   FragColor = sample_in_function();
+   FragColor += ubo;
+   FragColor += ssbo;
+   FragColor += ubos[1].ubo;
+   FragColor += registers.push;
+}

spirv_common.hpp

@@ -1404,8 +1404,9 @@ struct Meta
 		{
 			uint32_t packed_type = 0;
 			bool packed = false;
-			uint32_t ib_member_index = static_cast<uint32_t>(-1);
+			uint32_t ib_member_index = ~(0u);
 			uint32_t ib_orig_id = 0;
+			uint32_t argument_buffer_id = ~(0u);
 		} extended;
 	};
 

spirv_cross.cpp

@@ -17,6 +17,7 @@
 #include "spirv_cross.hpp"
 #include "GLSL.std.450.h"
 #include "spirv_cfg.hpp"
+#include "spirv_common.hpp"
 #include "spirv_parser.hpp"
 #include <algorithm>
 #include <cstring>
@@ -1091,6 +1092,11 @@ const std::string &Compiler::get_member_name(uint32_t id, uint32_t index) const
 	return ir.get_member_name(id, index);
 }
 
+void Compiler::set_qualified_name(uint32_t id, const string &name)
+{
+	ir.meta[id].decoration.qualified_alias = name;
+}
+
 void Compiler::set_member_qualified_name(uint32_t type_id, uint32_t index, const std::string &name)
 {
 	ir.meta[type_id].members.resize(max(ir.meta[type_id].members.size(), size_t(index) + 1));
@@ -1156,6 +1162,10 @@ void Compiler::set_extended_decoration(uint32_t id, ExtendedDecorations decorati
 	case SPIRVCrossDecorationInterfaceOrigID:
 		dec.extended.ib_orig_id = value;
 		break;
+
+	case SPIRVCrossDecorationArgumentBufferID:
+		dec.extended.argument_buffer_id = value;
+		break;
 	}
 }
 
@@ -1182,6 +1192,10 @@ void Compiler::set_extended_member_decoration(uint32_t type, uint32_t index, Ext
 	case SPIRVCrossDecorationInterfaceOrigID:
 		dec.extended.ib_orig_id = value;
 		break;
+
+	case SPIRVCrossDecorationArgumentBufferID:
+		dec.extended.argument_buffer_id = value;
+		break;
 	}
 }
 
@@ -1205,6 +1219,9 @@ uint32_t Compiler::get_extended_decoration(uint32_t id, ExtendedDecorations deco
 
 	case SPIRVCrossDecorationInterfaceOrigID:
 		return dec.extended.ib_orig_id;
+
+	case SPIRVCrossDecorationArgumentBufferID:
+		return dec.extended.argument_buffer_id;
 	}
 
 	return 0;
@@ -1233,6 +1250,9 @@ uint32_t Compiler::get_extended_member_decoration(uint32_t type, uint32_t index,
 
 	case SPIRVCrossDecorationInterfaceOrigID:
 		return dec.extended.ib_orig_id;
+
+	case SPIRVCrossDecorationArgumentBufferID:
+		return dec.extended.argument_buffer_id;
 	}
 
 	return 0;
@@ -1258,6 +1278,9 @@ bool Compiler::has_extended_decoration(uint32_t id, ExtendedDecorations decorati
 
 	case SPIRVCrossDecorationInterfaceOrigID:
 		return dec.extended.ib_orig_id != 0;
+
+	case SPIRVCrossDecorationArgumentBufferID:
+		return dec.extended.argument_buffer_id != 0;
 	}
 
 	return false;
@@ -1286,6 +1309,9 @@ bool Compiler::has_extended_member_decoration(uint32_t type, uint32_t index, Ext
 
 	case SPIRVCrossDecorationInterfaceOrigID:
 		return dec.extended.ib_orig_id != 0;
+
+	case SPIRVCrossDecorationArgumentBufferID:
+		return dec.extended.argument_buffer_id != uint32_t(-1);
 	}
 
 	return false;
@@ -1305,12 +1331,16 @@ void Compiler::unset_extended_decoration(uint32_t id, ExtendedDecorations decora
 		break;
 
 	case SPIRVCrossDecorationInterfaceMemberIndex:
-		dec.extended.ib_member_index = -1;
+		dec.extended.ib_member_index = ~(0u);
 		break;
 
 	case SPIRVCrossDecorationInterfaceOrigID:
 		dec.extended.ib_orig_id = 0;
 		break;
+
+	case SPIRVCrossDecorationArgumentBufferID:
+		dec.extended.argument_buffer_id = 0;
+		break;
 	}
 }
 
@@ -1330,12 +1360,16 @@ void Compiler::unset_extended_member_decoration(uint32_t type, uint32_t index, E
 		break;
 
 	case SPIRVCrossDecorationInterfaceMemberIndex:
-		dec.extended.ib_member_index = -1;
+		dec.extended.ib_member_index = ~(0u);
 		break;
 
 	case SPIRVCrossDecorationInterfaceOrigID:
 		dec.extended.ib_orig_id = 0;
 		break;
+
+	case SPIRVCrossDecorationArgumentBufferID:
+		dec.extended.argument_buffer_id = 0;
+		break;
 	}
 }
 

spirv_cross.hpp

@@ -120,6 +120,7 @@ enum ExtendedDecorations
 	SPIRVCrossDecorationPackedType,
 	SPIRVCrossDecorationInterfaceMemberIndex,
 	SPIRVCrossDecorationInterfaceOrigID,
+	SPIRVCrossDecorationArgumentBufferID
 };
 
 class Compiler
@@ -209,9 +210,6 @@ public:
 	// or an empty string if no qualified alias exists
 	const std::string &get_member_qualified_name(uint32_t type_id, uint32_t index) const;
 
-	// Sets the qualified member identifier for OpTypeStruct ID, member number "index".
-	void set_member_qualified_name(uint32_t type_id, uint32_t index, const std::string &name);
-
 	// Gets the decoration mask for a member of a struct, similar to get_decoration_mask.
 	const Bitset &get_member_decoration_bitset(uint32_t id, uint32_t index) const;
 
@@ -587,6 +585,10 @@ protected:
 	// Gets the SPIR-V element type underlying an array variable.
 	const SPIRType &get_variable_element_type(const SPIRVariable &var) const;
 
+	// Sets the qualified member identifier for OpTypeStruct ID, member number "index".
+	void set_member_qualified_name(uint32_t type_id, uint32_t index, const std::string &name);
+	void set_qualified_name(uint32_t id, const std::string &name);
+
 	// Returns if the given type refers to a sampled image.
 	bool is_sampled_image_type(const SPIRType &type);
 

spirv_cross_c.cpp

@@ -470,6 +470,10 @@ spvc_result spvc_compiler_options_set_uint(spvc_compiler_options options, spvc_c
 		options->msl.platform = static_cast<CompilerMSL::Options::Platform>(value);
 		break;
 
+	case SPVC_COMPILER_OPTION_MSL_ARGUMENT_BUFFERS:
+		options->msl.argument_buffers = value != 0;
+		break;
+
 	default:
 		options->context->report_error("Unknown option.");
 		return SPVC_ERROR_INVALID_ARGUMENT;

spirv_cross_c.h

@@ -33,7 +33,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 1
+#define SPVC_C_API_VERSION_MINOR 2
 /* Bumped if internal implementation details change. */
 #define SPVC_C_API_VERSION_PATCH 0
 
@@ -418,6 +418,7 @@ typedef enum spvc_compiler_option
 	SPVC_COMPILER_OPTION_MSL_PAD_FRAGMENT_OUTPUT_COMPONENTS = 29 | SPVC_COMPILER_OPTION_MSL_BIT,
 	SPVC_COMPILER_OPTION_MSL_TESS_DOMAIN_ORIGIN_LOWER_LEFT = 30 | SPVC_COMPILER_OPTION_MSL_BIT,
 	SPVC_COMPILER_OPTION_MSL_PLATFORM = 31 | SPVC_COMPILER_OPTION_MSL_BIT,
+	SPVC_COMPILER_OPTION_MSL_ARGUMENT_BUFFERS = 32 | SPVC_COMPILER_OPTION_MSL_BIT,
 
 	SPVC_COMPILER_OPTION_INT_MAX = 0x7fffffff
 } spvc_compiler_option;

spirv_msl.cpp

@@ -630,6 +630,15 @@ string CompilerMSL::compile()
 	// the loop, so the hooks aren't added multiple times.
 	fix_up_shader_inputs_outputs();
 
+	// If we are using argument buffers, we create argument buffer structures for them here.
+	// These buffers will be used in the entry point, not the individual resources.
+	if (msl_options.argument_buffers)
+	{
+		if (!msl_options.supports_msl_version(2, 0))
+			SPIRV_CROSS_THROW("Argument buffers can only be used with MSL 2.0 and up.");
+		analyze_argument_buffers();
+	}
+
 	uint32_t pass_count = 0;
 	do
 	{
@@ -4271,7 +4280,10 @@ void CompilerMSL::emit_function_prototype(SPIRFunction &func, const Bitset &)
 
 	if (processing_entry_point)
 	{
-		decl += entry_point_args(!func.arguments.empty());
+		if (msl_options.argument_buffers)
+			decl += entry_point_args_argument_buffer(!func.arguments.empty());
+		else
+			decl += entry_point_args_classic(!func.arguments.empty());
 
 		// If entry point function has variables that require early declaration,
 		// ensure they each have an empty initializer, creating one if needed.
@@ -4778,7 +4790,15 @@ string CompilerMSL::to_func_call_arg(uint32_t id)
 	// Manufacture automatic sampler arg if the arg is a SampledImage texture.
 	auto &type = expression_type(id);
 	if (type.basetype == SPIRType::SampledImage && type.image.dim != DimBuffer)
-		arg_str += ", " + to_sampler_expression(id);
+	{
+		// Need to check the base variable in case we need to apply a qualified alias.
+		uint32_t var_id = 0;
+		auto *sampler_var = maybe_get<SPIRVariable>(id);
+		if (sampler_var)
+			var_id = sampler_var->basevariable;
+
+		arg_str += ", " + to_sampler_expression(var_id ? var_id : id);
+	}
 	if (msl_options.swizzle_texture_samples && has_sampled_images && is_sampled_image_type(type))
 		arg_str += ", " + to_swizzle_expression(id);
 
@@ -4963,6 +4983,10 @@ string CompilerMSL::to_struct_member(const SPIRType &type, uint32_t member_type_
 	const SPIRType *effective_membertype = &membertype;
 	SPIRType override_type;
 
+	uint32_t orig_id = 0;
+	if (has_extended_member_decoration(type.self, index, SPIRVCrossDecorationInterfaceOrigID))
+		orig_id = get_extended_member_decoration(type.self, index, SPIRVCrossDecorationInterfaceOrigID);
+
 	if (member_is_packed_type(type, index))
 	{
 		// If we're packing a matrix, output an appropriate typedef
@@ -4992,8 +5016,16 @@ string CompilerMSL::to_struct_member(const SPIRType &type, uint32_t member_type_
 			pack_pfx = "packed_";
 	}
 
-	return join(pack_pfx, type_to_glsl(*effective_membertype), " ", qualifier, to_member_name(type, index),
-	            member_attribute_qualifier(type, index), type_to_array_glsl(membertype), ";");
+	// Array information is baked into these types.
+	string array_type;
+	if (membertype.basetype != SPIRType::Image && membertype.basetype != SPIRType::Sampler &&
+	    membertype.basetype != SPIRType::SampledImage)
+	{
+		array_type = type_to_array_glsl(membertype);
+	}
+
+	return join(pack_pfx, type_to_glsl(*effective_membertype, orig_id), " ", qualifier, to_member_name(type, index),
+	            member_attribute_qualifier(type, index), array_type, ";");
 }
 
 // Emit a structure member, padding and packing to maintain the correct memeber alignments.
@@ -5014,6 +5046,10 @@ string CompilerMSL::member_attribute_qualifier(const SPIRType &type, uint32_t in
 	BuiltIn builtin = BuiltInMax;
 	bool is_builtin = is_member_builtin(type, index, &builtin);
 
+	if (has_extended_member_decoration(type.self, index, SPIRVCrossDecorationArgumentBufferID))
+		return join(" [[id(", get_extended_member_decoration(type.self, index, SPIRVCrossDecorationArgumentBufferID),
+		            ")]]");
+
 	// Vertex function inputs
 	if (execution.model == ExecutionModelVertex && type.storage == StorageClassInput)
 	{
@@ -5435,10 +5471,9 @@ string CompilerMSL::get_type_address_space(const SPIRType &type)
 	return "thread";
 }
 
-// Returns a string containing a comma-delimited list of args for the entry point function
-string CompilerMSL::entry_point_args(bool append_comma)
+string CompilerMSL::entry_point_arg_stage_in()
 {
-	string ep_args;
+	string decl;
 
 	// Stage-in structure
 	uint32_t stage_in_id;
@@ -5452,126 +5487,15 @@ string CompilerMSL::entry_point_args(bool append_comma)
 		auto &var = get<SPIRVariable>(stage_in_id);
 		auto &type = get_variable_data_type(var);
 
-		if (!ep_args.empty())
-			ep_args += ", ";
-
 		add_resource_name(var.self);
-		ep_args += join(type_to_glsl(type), " ", to_name(var.self), " [[stage_in]]");
+		decl = join(type_to_glsl(type), " ", to_name(var.self), " [[stage_in]]");
 	}
 
-	// Output resources, sorted by resource index & type
-	// We need to sort to work around a bug on macOS 10.13 with NVidia drivers where switching between shaders
-	// with different order of buffers can result in issues with buffer assignments inside the driver.
-	struct Resource
-	{
-		Variant *id;
-		string name;
-		SPIRType::BaseType basetype;
-		uint32_t index;
-	};
-
-	vector<Resource> resources;
-
-	ir.for_each_typed_id<SPIRVariable>([&](uint32_t self, SPIRVariable &var) {
-		auto &id = ir.ids[self];
-		auto &type = get_variable_data_type(var);
-
-		uint32_t var_id = var.self;
-
-		if ((var.storage == StorageClassUniform || var.storage == StorageClassUniformConstant ||
-		     var.storage == StorageClassPushConstant || var.storage == StorageClassStorageBuffer) &&
-		    !is_hidden_variable(var))
-		{
-			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) });
-
-				if (type.image.dim != DimBuffer && constexpr_samplers.count(var_id) == 0)
-				{
-					resources.push_back({ &id, to_sampler_expression(var_id), SPIRType::Sampler,
-					                      get_metal_resource_index(var, SPIRType::Sampler) });
-				}
-			}
-			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) });
-			}
-		}
-	});
-
-	std::sort(resources.begin(), resources.end(), [](const Resource &lhs, const Resource &rhs) {
-		return tie(lhs.basetype, lhs.index) < tie(rhs.basetype, rhs.index);
-	});
-
-	for (auto &r : resources)
-	{
-		auto &var = r.id->get<SPIRVariable>();
-		auto &type = get_variable_data_type(var);
-
-		uint32_t var_id = var.self;
-
-		switch (r.basetype)
-		{
-		case SPIRType::Struct:
-		{
-			auto &m = ir.meta[type.self];
-			if (m.members.size() == 0)
-				break;
-			if (!type.array.empty())
-			{
-				if (type.array.size() > 1)
-					SPIRV_CROSS_THROW("Arrays of arrays of buffers are not supported.");
-
-				// Metal doesn't directly support this, so we must expand the
-				// array. We'll declare a local array to hold these elements
-				// later.
-				uint32_t array_size = to_array_size_literal(type);
-
-				if (array_size == 0)
-					SPIRV_CROSS_THROW("Unsized arrays of buffers are not supported in MSL.");
-
-				buffer_arrays.push_back(var_id);
-				for (uint32_t i = 0; i < array_size; ++i)
-				{
-					if (!ep_args.empty())
-						ep_args += ", ";
-					ep_args += get_argument_address_space(var) + " " + type_to_glsl(type) + "* " + r.name + "_" +
-					           convert_to_string(i);
-					ep_args += " [[buffer(" + convert_to_string(r.index + i) + ")]]";
-				}
-			}
-			else
-			{
-				if (!ep_args.empty())
-					ep_args += ", ";
-				ep_args += get_argument_address_space(var) + " " + type_to_glsl(type) + "& " + r.name;
-				ep_args += " [[buffer(" + convert_to_string(r.index) + ")]]";
-			}
-			break;
-		}
-		case SPIRType::Sampler:
-			if (!ep_args.empty())
-				ep_args += ", ";
-			ep_args += sampler_type(type) + " " + r.name;
-			ep_args += " [[sampler(" + convert_to_string(r.index) + ")]]";
-			break;
-		case SPIRType::Image:
-			if (!ep_args.empty())
-				ep_args += ", ";
-			ep_args += image_type_glsl(type, var_id) + " " + r.name;
-			ep_args += " [[texture(" + convert_to_string(r.index) + ")]]";
-			break;
-		default:
-			SPIRV_CROSS_THROW("Unexpected resource type");
-			break;
-		}
-	}
+	return decl;
+}
 
+void CompilerMSL::entry_point_args_builtin(string &ep_args)
+{
 	// Builtin variables
 	ir.for_each_typed_id<SPIRVariable>([&](uint32_t var_id, SPIRVariable &var) {
 		BuiltIn bi_type = ir.meta[var_id].decoration.builtin_type;
@@ -5650,6 +5574,178 @@ string CompilerMSL::entry_point_args(bool append_comma)
 			}
 		}
 	}
+}
+
+string CompilerMSL::entry_point_args_argument_buffer(bool append_comma)
+{
+	string ep_args = entry_point_arg_stage_in();
+
+	for (uint32_t i = 0; i < kMaxArgumentBuffers; i++)
+	{
+		uint32_t id = argument_buffer_ids[i];
+		if (id == 0)
+			continue;
+
+		add_resource_name(id);
+		auto &var = get<SPIRVariable>(id);
+		auto &type = get_variable_data_type(var);
+
+		if (!ep_args.empty())
+			ep_args += ", ";
+
+		ep_args += get_argument_address_space(var) + " " + type_to_glsl(type) + "& " + to_name(id);
+		ep_args += " [[buffer(" + convert_to_string(i) + ")]]";
+
+		// Makes it more practical for testing, since the push constant block can occupy the first available
+		// buffer slot if it's not bound explicitly.
+		next_metal_resource_index_buffer = i + 1;
+	}
+
+	// Declare the push constant block separately, it does not belong in an indirect argument buffer.
+	uint32_t push_constant_id = 0;
+	ir.for_each_typed_id<SPIRVariable>([&](uint32_t self, SPIRVariable &var) {
+		if (var.storage == StorageClassPushConstant && !is_hidden_variable(var))
+			push_constant_id = self;
+	});
+
+	if (push_constant_id)
+	{
+		add_resource_name(push_constant_id);
+		auto &var = get<SPIRVariable>(push_constant_id);
+		auto &type = get_variable_data_type(var);
+		uint32_t resource_index = get_metal_resource_index(var, type.basetype);
+		if (!ep_args.empty())
+			ep_args += ", ";
+		ep_args += get_argument_address_space(var) + " " + type_to_glsl(type) + "& " + to_name(push_constant_id);
+		ep_args += " [[buffer(" + convert_to_string(resource_index) + ")]]";
+	}
+
+	entry_point_args_builtin(ep_args);
+
+	if (!ep_args.empty() && append_comma)
+		ep_args += ", ";
+
+	return ep_args;
+}
+
+// Returns a string containing a comma-delimited list of args for the entry point function
+// This is the "classic" method of MSL 1 when we don't have argument buffer support.
+string CompilerMSL::entry_point_args_classic(bool append_comma)
+{
+	string ep_args = entry_point_arg_stage_in();
+
+	// Output resources, sorted by resource index & type
+	// We need to sort to work around a bug on macOS 10.13 with NVidia drivers where switching between shaders
+	// with different order of buffers can result in issues with buffer assignments inside the driver.
+	struct Resource
+	{
+		SPIRVariable *var;
+		string name;
+		SPIRType::BaseType basetype;
+		uint32_t index;
+	};
+
+	vector<Resource> resources;
+
+	ir.for_each_typed_id<SPIRVariable>([&](uint32_t, SPIRVariable &var) {
+		if ((var.storage == StorageClassUniform || var.storage == StorageClassUniformConstant ||
+		     var.storage == StorageClassPushConstant || var.storage == StorageClassStorageBuffer) &&
+		    !is_hidden_variable(var))
+		{
+			auto &type = get_variable_data_type(var);
+			uint32_t var_id = var.self;
+
+			if (type.basetype == SPIRType::SampledImage)
+			{
+				add_resource_name(var_id);
+				resources.push_back(
+				    { &var, to_name(var_id), SPIRType::Image, get_metal_resource_index(var, SPIRType::Image) });
+
+				if (type.image.dim != DimBuffer && constexpr_samplers.count(var_id) == 0)
+				{
+					resources.push_back({ &var, to_sampler_expression(var_id), SPIRType::Sampler,
+					                      get_metal_resource_index(var, SPIRType::Sampler) });
+				}
+			}
+			else if (constexpr_samplers.count(var_id) == 0)
+			{
+				// constexpr samplers are not declared as resources.
+				add_resource_name(var_id);
+				resources.push_back(
+				    { &var, to_name(var_id), type.basetype, get_metal_resource_index(var, type.basetype) });
+			}
+		}
+	});
+
+	std::sort(resources.begin(), resources.end(), [](const Resource &lhs, const Resource &rhs) {
+		return tie(lhs.basetype, lhs.index) < tie(rhs.basetype, rhs.index);
+	});
+
+	for (auto &r : resources)
+	{
+		auto &var = *r.var;
+		auto &type = get_variable_data_type(var);
+
+		uint32_t var_id = var.self;
+
+		switch (r.basetype)
+		{
+		case SPIRType::Struct:
+		{
+			auto &m = ir.meta[type.self];
+			if (m.members.size() == 0)
+				break;
+			if (!type.array.empty())
+			{
+				if (type.array.size() > 1)
+					SPIRV_CROSS_THROW("Arrays of arrays of buffers are not supported.");
+
+				// Metal doesn't directly support this, so we must expand the
+				// array. We'll declare a local array to hold these elements
+				// later.
+				uint32_t array_size = to_array_size_literal(type);
+
+				if (array_size == 0)
+					SPIRV_CROSS_THROW("Unsized arrays of buffers are not supported in MSL.");
+
+				buffer_arrays.push_back(var_id);
+				for (uint32_t i = 0; i < array_size; ++i)
+				{
+					if (!ep_args.empty())
+						ep_args += ", ";
+					ep_args += get_argument_address_space(var) + " " + type_to_glsl(type) + "* " + r.name + "_" +
+					           convert_to_string(i);
+					ep_args += " [[buffer(" + convert_to_string(r.index + i) + ")]]";
+				}
+			}
+			else
+			{
+				if (!ep_args.empty())
+					ep_args += ", ";
+				ep_args += get_argument_address_space(var) + " " + type_to_glsl(type) + "& " + r.name;
+				ep_args += " [[buffer(" + convert_to_string(r.index) + ")]]";
+			}
+			break;
+		}
+		case SPIRType::Sampler:
+			if (!ep_args.empty())
+				ep_args += ", ";
+			ep_args += sampler_type(type) + " " + r.name;
+			ep_args += " [[sampler(" + convert_to_string(r.index) + ")]]";
+			break;
+		case SPIRType::Image:
+			if (!ep_args.empty())
+				ep_args += ", ";
+			ep_args += image_type_glsl(type, var_id) + " " + r.name;
+			ep_args += " [[texture(" + convert_to_string(r.index) + ")]]";
+			break;
+		default:
+			SPIRV_CROSS_THROW("Unexpected resource type");
+			break;
+		}
+	}
+
+	entry_point_args_builtin(ep_args);
 
 	if (!ep_args.empty() && append_comma)
 		ep_args += ", ";
@@ -6230,9 +6326,16 @@ string CompilerMSL::to_member_reference(uint32_t base, const SPIRType &type, uin
 	auto *var = maybe_get<SPIRVariable>(base);
 	// If this is a buffer array, we have to dereference the buffer pointers.
 	// Otherwise, if this is a pointer expression, dereference it.
-	if ((var && ((var->storage == StorageClassUniform || var->storage == StorageClassStorageBuffer) &&
-	             is_array(get<SPIRType>(var->basetype)))) ||
-	    (!ptr_chain && should_dereference(base)))
+
+	bool declared_as_pointer = false;
+
+	if (var)
+	{
+		bool is_buffer_variable = var->storage == StorageClassUniform || var->storage == StorageClassStorageBuffer;
+		declared_as_pointer = is_buffer_variable && is_array(get<SPIRType>(var->basetype));
+	}
+
+	if (declared_as_pointer || (!ptr_chain && should_dereference(base)))
 		return join("->", to_member_name(type, index));
 	else
 		return join(".", to_member_name(type, index));
@@ -7391,3 +7494,158 @@ std::string CompilerMSL::to_initializer_expression(const SPIRVariable &var)
 	else
 		return CompilerGLSL::to_initializer_expression(var);
 }
+
+void CompilerMSL::analyze_argument_buffers()
+{
+	// Gather all used resources and sort them out into argument buffers.
+	// Each argument buffer corresponds to a descriptor set in SPIR-V.
+	// The [[id(N)]] values used correspond to the resource mapping we have for MSL.
+	// Otherwise, the binding number is used, but this is generally not safe some types like
+	// combined image samplers and arrays of resources. Metal needs different indices here,
+	// while SPIR-V can have one descriptor set binding. To use argument buffers in practice,
+	// you will need to use the remapping from the API.
+	for (auto &id : argument_buffer_ids)
+		id = 0;
+
+	// Output resources, sorted by resource index & type.
+	struct Resource
+	{
+		SPIRVariable *var;
+		string name;
+		SPIRType::BaseType basetype;
+		uint32_t index;
+	};
+	vector<Resource> resources_in_set[kMaxArgumentBuffers];
+
+	ir.for_each_typed_id<SPIRVariable>([&](uint32_t self, SPIRVariable &var) {
+		if ((var.storage == StorageClassUniform || var.storage == StorageClassUniformConstant ||
+		     var.storage == StorageClassStorageBuffer) &&
+		    !is_hidden_variable(var))
+		{
+			uint32_t desc_set = get_decoration(self, DecorationDescriptorSet);
+			uint32_t var_id = var.self;
+			auto &type = get_variable_data_type(var);
+
+			if (desc_set >= kMaxArgumentBuffers)
+				SPIRV_CROSS_THROW("Descriptor set index is out of range.");
+
+			if (type.basetype == SPIRType::SampledImage)
+			{
+				add_resource_name(var_id);
+
+				uint32_t image_resource_index = get_metal_resource_index(var, SPIRType::Image);
+				uint32_t sampler_resource_index = get_metal_resource_index(var, SPIRType::Sampler);
+
+				// Avoid trivial conflicts where we didn't remap.
+				// This will let us at least compile test cases without having to instrument remaps.
+				if (sampler_resource_index == image_resource_index)
+					sampler_resource_index += type.array.empty() ? 1 : to_array_size_literal(type);
+
+				resources_in_set[desc_set].push_back({ &var, to_name(var_id), SPIRType::Image, image_resource_index });
+
+				if (type.image.dim != DimBuffer && constexpr_samplers.count(var_id) == 0)
+				{
+					resources_in_set[desc_set].push_back(
+					    { &var, to_sampler_expression(var_id), SPIRType::Sampler, sampler_resource_index });
+				}
+			}
+			else if (constexpr_samplers.count(var_id) == 0)
+			{
+				// constexpr samplers are not declared as resources.
+				add_resource_name(var_id);
+				resources_in_set[desc_set].push_back(
+				    { &var, to_name(var_id), type.basetype, get_metal_resource_index(var, type.basetype) });
+			}
+		}
+	});
+
+	for (uint32_t desc_set = 0; desc_set < kMaxArgumentBuffers; desc_set++)
+	{
+		auto &resources = resources_in_set[desc_set];
+		if (resources.empty())
+			continue;
+
+		uint32_t next_id = ir.increase_bound_by(3);
+		uint32_t type_id = next_id + 1;
+		uint32_t ptr_type_id = next_id + 2;
+		argument_buffer_ids[desc_set] = next_id;
+
+		auto &buffer_type = set<SPIRType>(type_id);
+		buffer_type.storage = StorageClassUniform;
+		buffer_type.basetype = SPIRType::Struct;
+		set_name(type_id, join("spvDescriptorSetBuffer", desc_set));
+
+		auto &ptr_type = set<SPIRType>(ptr_type_id);
+		ptr_type = buffer_type;
+		ptr_type.pointer = true;
+		ptr_type.pointer_depth = 1;
+		ptr_type.parent_type = type_id;
+
+		uint32_t buffer_variable_id = next_id;
+		set<SPIRVariable>(buffer_variable_id, ptr_type_id, StorageClassUniform);
+		set_name(buffer_variable_id, join("spvDescriptorSet", desc_set));
+
+		// Ids must be emitted in ID order.
+		sort(begin(resources), end(resources), [&](const Resource &lhs, const Resource &rhs) -> bool {
+			return tie(lhs.index, lhs.basetype) < tie(rhs.index, rhs.basetype);
+		});
+
+		uint32_t member_index = 0;
+		for (auto &resource : resources)
+		{
+			auto &var = *resource.var;
+			auto &type = get_variable_data_type(var);
+			string mbr_name = ensure_valid_name(resource.name, "m");
+			set_member_name(buffer_type.self, member_index, mbr_name);
+
+			if (resource.basetype == SPIRType::Sampler && type.basetype != SPIRType::Sampler)
+			{
+				// Have to synthesize a sampler type here.
+
+				bool type_is_array = !type.array.empty();
+				uint32_t sampler_type_id = ir.increase_bound_by(type_is_array ? 2 : 1);
+				auto &new_sampler_type = set<SPIRType>(sampler_type_id);
+				new_sampler_type.basetype = SPIRType::Sampler;
+				new_sampler_type.storage = StorageClassUniformConstant;
+
+				if (type_is_array)
+				{
+					uint32_t sampler_type_array_id = sampler_type_id + 1;
+					auto &sampler_type_array = set<SPIRType>(sampler_type_array_id);
+					sampler_type_array = new_sampler_type;
+					sampler_type_array.array = type.array;
+					sampler_type_array.array_size_literal = type.array_size_literal;
+					sampler_type_array.parent_type = sampler_type_id;
+					buffer_type.member_types.push_back(sampler_type_array_id);
+				}
+				else
+					buffer_type.member_types.push_back(sampler_type_id);
+			}
+			else
+			{
+				if (resource.basetype == SPIRType::Image || resource.basetype == SPIRType::Sampler ||
+				    resource.basetype == SPIRType::SampledImage)
+				{
+					// Drop pointer information when we emit the resources into a struct.
+					buffer_type.member_types.push_back(get_variable_data_type_id(var));
+					set_qualified_name(var.self, join(to_name(buffer_variable_id), ".", mbr_name));
+				}
+				else
+				{
+					// Resources will be declared as pointers not references, so automatically dereference as appropriate.
+					buffer_type.member_types.push_back(var.basetype);
+					if (type.array.empty())
+						set_qualified_name(var.self, join("(*", to_name(buffer_variable_id), ".", mbr_name, ")"));
+					else
+						set_qualified_name(var.self, join(to_name(buffer_variable_id), ".", mbr_name));
+				}
+			}
+
+			set_extended_member_decoration(buffer_type.self, member_index, SPIRVCrossDecorationArgumentBufferID,
+			                               resource.index);
+			set_extended_member_decoration(buffer_type.self, member_index, SPIRVCrossDecorationInterfaceOrigID,
+			                               var.self);
+			member_index++;
+		}
+	}
+}

spirv_msl.hpp

@@ -148,6 +148,8 @@ static const uint32_t kPushConstDescSet = ~(0u);
 // element to indicate the bindings for the push constants.
 static const uint32_t kPushConstBinding = 0;
 
+static const uint32_t kMaxArgumentBuffers = 8;
+
 // The current version of the aux buffer structure. It must be incremented any time a
 // new field is added to the aux buffer.
 #define SPIRV_CROSS_MSL_AUX_BUFFER_STRUCT_VERSION 1
@@ -179,6 +181,7 @@ public:
 		bool capture_output_to_buffer = false;
 		bool swizzle_texture_samples = false;
 		bool tess_domain_origin_lower_left = false;
+		bool argument_buffers = false;
 
 		// Fragment output in MSL must have at least as many components as the render pass.
 		// Add support to explicit pad out components.
@@ -413,7 +416,10 @@ protected:
 	void fix_up_shader_inputs_outputs();
 
 	std::string func_type_decl(SPIRType &type);
-	std::string entry_point_args(bool append_comma);
+	std::string entry_point_args_classic(bool append_comma);
+	std::string entry_point_args_argument_buffer(bool append_comma);
+	std::string entry_point_arg_stage_in();
+	void entry_point_args_builtin(std::string &args);
 	std::string to_qualified_member_name(const SPIRType &type, uint32_t index);
 	std::string ensure_valid_name(std::string name, std::string pfx);
 	std::string to_sampler_expression(uint32_t id);
@@ -513,6 +519,9 @@ protected:
 	std::unordered_map<uint32_t, MSLConstexprSampler> constexpr_samplers;
 	std::vector<uint32_t> buffer_arrays;
 
+	uint32_t argument_buffer_ids[kMaxArgumentBuffers];
+	void analyze_argument_buffers();
+
 	uint32_t get_target_components_for_fragment_location(uint32_t location) const;
 	uint32_t build_extended_vector_type(uint32_t type_id, uint32_t components);
 

test_shaders.py

@@ -169,6 +169,8 @@ def cross_compile_msl(shader, spirv, opt, paths):
         msl_args.append('--msl-capture-output')
     if '.domain.' in shader:
         msl_args.append('--msl-domain-lower-left')
+    if '.argument.' in shader:
+        msl_args.append('--msl-argument-buffers')
 
     subprocess.check_call(msl_args)