spirv_msl.cpp

/*
 * Copyright 2016-2021 The Brenwill Workshop Ltd.
 * SPDX-License-Identifier: Apache-2.0 OR MIT
 *
 * Licensed under the Apache License, Version 2.0 (the "License");
 * you may not use this file except in compliance with the License.
 * You may obtain a copy of the License at
 *
 *     http://www.apache.org/licenses/LICENSE-2.0
 *
 * Unless required by applicable law or agreed to in writing, software
 * distributed under the License is distributed on an "AS IS" BASIS,
 * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
 * See the License for the specific language governing permissions and
 * limitations under the License.
 */

/*
 * At your option, you may choose to accept this material under either:
 *  1. The Apache License, Version 2.0, found at <http://www.apache.org/licenses/LICENSE-2.0>, or
 *  2. The MIT License, found at <http://opensource.org/licenses/MIT>.
 */

#include "spirv_msl.hpp"
#include "GLSL.std.450.h"

#include <algorithm>
#include <assert.h>
#include <numeric>

using namespace spv;
using namespace SPIRV_CROSS_NAMESPACE;
using namespace std;

static const uint32_t k_unknown_location = ~0u;
static const uint32_t k_unknown_component = ~0u;
static const char *force_inline = "static inline __attribute__((always_inline))";

CompilerMSL::CompilerMSL(std::vector<uint32_t> spirv_)
    : CompilerGLSL(std::move(spirv_))
{
}

CompilerMSL::CompilerMSL(const uint32_t *ir_, size_t word_count)
    : CompilerGLSL(ir_, word_count)
{
}

CompilerMSL::CompilerMSL(const ParsedIR &ir_)
    : CompilerGLSL(ir_)
{
}

CompilerMSL::CompilerMSL(ParsedIR &&ir_)
    : CompilerGLSL(std::move(ir_))
{
}

void CompilerMSL::add_msl_shader_input(const MSLShaderInterfaceVariable &si)
{
	inputs_by_location[{si.location, si.component}] = si;
	if (si.builtin != BuiltInMax && !inputs_by_builtin.count(si.builtin))
		inputs_by_builtin[si.builtin] = si;
}

void CompilerMSL::add_msl_shader_output(const MSLShaderInterfaceVariable &so)
{
	outputs_by_location[{so.location, so.component}] = so;
	if (so.builtin != BuiltInMax && !outputs_by_builtin.count(so.builtin))
		outputs_by_builtin[so.builtin] = so;
}

void CompilerMSL::add_msl_resource_binding(const MSLResourceBinding &binding)
{
	StageSetBinding tuple = { binding.stage, binding.desc_set, binding.binding };
	resource_bindings[tuple] = { binding, false };

	// If we might need to pad argument buffer members to positionally align
	// arg buffer indexes, also maintain a lookup by argument buffer index.
	if (msl_options.pad_argument_buffer_resources)
	{
		StageSetBinding arg_idx_tuple = { binding.stage, binding.desc_set, k_unknown_component };

#define ADD_ARG_IDX_TO_BINDING_NUM_LOOKUP(rez) \
	arg_idx_tuple.binding = binding.msl_##rez; \
	resource_arg_buff_idx_to_binding_number[arg_idx_tuple] = binding.binding

		switch (binding.basetype)
		{
		case SPIRType::Void:
		case SPIRType::Boolean:
		case SPIRType::SByte:
		case SPIRType::UByte:
		case SPIRType::Short:
		case SPIRType::UShort:
		case SPIRType::Int:
		case SPIRType::UInt:
		case SPIRType::Int64:
		case SPIRType::UInt64:
		case SPIRType::AtomicCounter:
		case SPIRType::Half:
		case SPIRType::Float:
		case SPIRType::Double:
			ADD_ARG_IDX_TO_BINDING_NUM_LOOKUP(buffer);
			break;
		case SPIRType::Image:
			ADD_ARG_IDX_TO_BINDING_NUM_LOOKUP(texture);
			break;
		case SPIRType::Sampler:
			ADD_ARG_IDX_TO_BINDING_NUM_LOOKUP(sampler);
			break;
		case SPIRType::SampledImage:
			ADD_ARG_IDX_TO_BINDING_NUM_LOOKUP(texture);
			ADD_ARG_IDX_TO_BINDING_NUM_LOOKUP(sampler);
			break;
		default:
			SPIRV_CROSS_THROW("Unexpected argument buffer resource base type. When padding argument buffer elements, "
			                  "all descriptor set resources must be supplied with a base type by the app.");
		}
#undef ADD_ARG_IDX_TO_BINDING_NUM_LOOKUP
	}
}

void CompilerMSL::add_dynamic_buffer(uint32_t desc_set, uint32_t binding, uint32_t index)
{
	SetBindingPair pair = { desc_set, binding };
	buffers_requiring_dynamic_offset[pair] = { index, 0 };
}

void CompilerMSL::add_inline_uniform_block(uint32_t desc_set, uint32_t binding)
{
	SetBindingPair pair = { desc_set, binding };
	inline_uniform_blocks.insert(pair);
}

void CompilerMSL::add_discrete_descriptor_set(uint32_t desc_set)
{
	if (desc_set < kMaxArgumentBuffers)
		argument_buffer_discrete_mask |= 1u << desc_set;
}

void CompilerMSL::set_argument_buffer_device_address_space(uint32_t desc_set, bool device_storage)
{
	if (desc_set < kMaxArgumentBuffers)
	{
		if (device_storage)
			argument_buffer_device_storage_mask |= 1u << desc_set;
		else
			argument_buffer_device_storage_mask &= ~(1u << desc_set);
	}
}

bool CompilerMSL::is_msl_shader_input_used(uint32_t location)
{
	// Don't report internal location allocations to app.
	return location_inputs_in_use.count(location) != 0 &&
	       location_inputs_in_use_fallback.count(location) == 0;
}

bool CompilerMSL::is_msl_shader_output_used(uint32_t location)
{
	// Don't report internal location allocations to app.
	return location_outputs_in_use.count(location) != 0 &&
	       location_outputs_in_use_fallback.count(location) == 0;
}

uint32_t CompilerMSL::get_automatic_builtin_input_location(spv::BuiltIn builtin) const
{
	auto itr = builtin_to_automatic_input_location.find(builtin);
	if (itr == builtin_to_automatic_input_location.end())
		return k_unknown_location;
	else
		return itr->second;
}

uint32_t CompilerMSL::get_automatic_builtin_output_location(spv::BuiltIn builtin) const
{
	auto itr = builtin_to_automatic_output_location.find(builtin);
	if (itr == builtin_to_automatic_output_location.end())
		return k_unknown_location;
	else
		return itr->second;
}

bool CompilerMSL::is_msl_resource_binding_used(ExecutionModel model, uint32_t desc_set, uint32_t binding) const
{
	StageSetBinding tuple = { model, desc_set, binding };
	auto itr = resource_bindings.find(tuple);
	return itr != end(resource_bindings) && itr->second.second;
}

// Returns the size of the array of resources used by the variable with the specified id.
// The returned value is retrieved from the resource binding added using add_msl_resource_binding().
uint32_t CompilerMSL::get_resource_array_size(uint32_t id) const
{
	StageSetBinding tuple = { get_entry_point().model, get_decoration(id, DecorationDescriptorSet),
		                      get_decoration(id, DecorationBinding) };
	auto itr = resource_bindings.find(tuple);
	return itr != end(resource_bindings) ? itr->second.first.count : 0;
}

uint32_t CompilerMSL::get_automatic_msl_resource_binding(uint32_t id) const
{
	return get_extended_decoration(id, SPIRVCrossDecorationResourceIndexPrimary);
}

uint32_t CompilerMSL::get_automatic_msl_resource_binding_secondary(uint32_t id) const
{
	return get_extended_decoration(id, SPIRVCrossDecorationResourceIndexSecondary);
}

uint32_t CompilerMSL::get_automatic_msl_resource_binding_tertiary(uint32_t id) const
{
	return get_extended_decoration(id, SPIRVCrossDecorationResourceIndexTertiary);
}

uint32_t CompilerMSL::get_automatic_msl_resource_binding_quaternary(uint32_t id) const
{
	return get_extended_decoration(id, SPIRVCrossDecorationResourceIndexQuaternary);
}

void CompilerMSL::set_fragment_output_components(uint32_t location, uint32_t components)
{
	fragment_output_components[location] = components;
}

bool CompilerMSL::builtin_translates_to_nonarray(spv::BuiltIn builtin) const
{
	return (builtin == BuiltInSampleMask);
}

void CompilerMSL::build_implicit_builtins()
{
	bool need_sample_pos = active_input_builtins.get(BuiltInSamplePosition);
	bool need_vertex_params = capture_output_to_buffer && get_execution_model() == ExecutionModelVertex &&
	                          !msl_options.vertex_for_tessellation;
	bool need_tesc_params = is_tesc_shader();
	bool need_tese_params = is_tese_shader() && msl_options.raw_buffer_tese_input;
	bool need_subgroup_mask =
	    active_input_builtins.get(BuiltInSubgroupEqMask) || active_input_builtins.get(BuiltInSubgroupGeMask) ||
	    active_input_builtins.get(BuiltInSubgroupGtMask) || active_input_builtins.get(BuiltInSubgroupLeMask) ||
	    active_input_builtins.get(BuiltInSubgroupLtMask);
	bool need_subgroup_ge_mask = !msl_options.is_ios() && (active_input_builtins.get(BuiltInSubgroupGeMask) ||
	                                                       active_input_builtins.get(BuiltInSubgroupGtMask));
	bool need_multiview = get_execution_model() == ExecutionModelVertex && !msl_options.view_index_from_device_index &&
	                      msl_options.multiview_layered_rendering &&
	                      (msl_options.multiview || active_input_builtins.get(BuiltInViewIndex));
	bool need_dispatch_base =
	    msl_options.dispatch_base && get_execution_model() == ExecutionModelGLCompute &&
	    (active_input_builtins.get(BuiltInWorkgroupId) || active_input_builtins.get(BuiltInGlobalInvocationId));
	bool need_grid_params = get_execution_model() == ExecutionModelVertex && msl_options.vertex_for_tessellation;
	bool need_vertex_base_params =
	    need_grid_params &&
	    (active_input_builtins.get(BuiltInVertexId) || active_input_builtins.get(BuiltInVertexIndex) ||
	     active_input_builtins.get(BuiltInBaseVertex) || active_input_builtins.get(BuiltInInstanceId) ||
	     active_input_builtins.get(BuiltInInstanceIndex) || active_input_builtins.get(BuiltInBaseInstance));
	bool need_local_invocation_index = msl_options.emulate_subgroups && active_input_builtins.get(BuiltInSubgroupId);
	bool need_workgroup_size = msl_options.emulate_subgroups && active_input_builtins.get(BuiltInNumSubgroups);

	if (need_subpass_input || need_sample_pos || need_subgroup_mask || need_vertex_params || need_tesc_params ||
	    need_tese_params || need_multiview || need_dispatch_base || need_vertex_base_params || need_grid_params ||
	    needs_sample_id || needs_subgroup_invocation_id || needs_subgroup_size || needs_helper_invocation ||
		has_additional_fixed_sample_mask() || need_local_invocation_index || need_workgroup_size)
	{
		bool has_frag_coord = false;
		bool has_sample_id = false;
		bool has_vertex_idx = false;
		bool has_base_vertex = false;
		bool has_instance_idx = false;
		bool has_base_instance = false;
		bool has_invocation_id = false;
		bool has_primitive_id = false;
		bool has_subgroup_invocation_id = false;
		bool has_subgroup_size = false;
		bool has_view_idx = false;
		bool has_layer = false;
		bool has_helper_invocation = false;
		bool has_local_invocation_index = false;
		bool has_workgroup_size = false;
		uint32_t workgroup_id_type = 0;

		ir.for_each_typed_id<SPIRVariable>([&](uint32_t, SPIRVariable &var) {
			if (var.storage != StorageClassInput && var.storage != StorageClassOutput)
				return;
			if (!interface_variable_exists_in_entry_point(var.self))
				return;
			if (!has_decoration(var.self, DecorationBuiltIn))
				return;

			BuiltIn builtin = ir.meta[var.self].decoration.builtin_type;

			if (var.storage == StorageClassOutput)
			{
				if (has_additional_fixed_sample_mask() && builtin == BuiltInSampleMask)
				{
					builtin_sample_mask_id = var.self;
					mark_implicit_builtin(StorageClassOutput, BuiltInSampleMask, var.self);
					does_shader_write_sample_mask = true;
				}
			}

			if (var.storage != StorageClassInput)
				return;

			// Use Metal's native frame-buffer fetch API for subpass inputs.
			if (need_subpass_input && (!msl_options.use_framebuffer_fetch_subpasses))
			{
				switch (builtin)
				{
				case BuiltInFragCoord:
					mark_implicit_builtin(StorageClassInput, BuiltInFragCoord, var.self);
					builtin_frag_coord_id = var.self;
					has_frag_coord = true;
					break;
				case BuiltInLayer:
					if (!msl_options.arrayed_subpass_input || msl_options.multiview)
						break;
					mark_implicit_builtin(StorageClassInput, BuiltInLayer, var.self);
					builtin_layer_id = var.self;
					has_layer = true;
					break;
				case BuiltInViewIndex:
					if (!msl_options.multiview)
						break;
					mark_implicit_builtin(StorageClassInput, BuiltInViewIndex, var.self);
					builtin_view_idx_id = var.self;
					has_view_idx = true;
					break;
				default:
					break;
				}
			}

			if ((need_sample_pos || needs_sample_id) && builtin == BuiltInSampleId)
			{
				builtin_sample_id_id = var.self;
				mark_implicit_builtin(StorageClassInput, BuiltInSampleId, var.self);
				has_sample_id = true;
			}

			if (need_vertex_params)
			{
				switch (builtin)
				{
				case BuiltInVertexIndex:
					builtin_vertex_idx_id = var.self;
					mark_implicit_builtin(StorageClassInput, BuiltInVertexIndex, var.self);
					has_vertex_idx = true;
					break;
				case BuiltInBaseVertex:
					builtin_base_vertex_id = var.self;
					mark_implicit_builtin(StorageClassInput, BuiltInBaseVertex, var.self);
					has_base_vertex = true;
					break;
				case BuiltInInstanceIndex:
					builtin_instance_idx_id = var.self;
					mark_implicit_builtin(StorageClassInput, BuiltInInstanceIndex, var.self);
					has_instance_idx = true;
					break;
				case BuiltInBaseInstance:
					builtin_base_instance_id = var.self;
					mark_implicit_builtin(StorageClassInput, BuiltInBaseInstance, var.self);
					has_base_instance = true;
					break;
				default:
					break;
				}
			}

			if (need_tesc_params && builtin == BuiltInInvocationId)
			{
				builtin_invocation_id_id = var.self;
				mark_implicit_builtin(StorageClassInput, BuiltInInvocationId, var.self);
				has_invocation_id = true;
			}

			if ((need_tesc_params || need_tese_params) && builtin == BuiltInPrimitiveId)
			{
				builtin_primitive_id_id = var.self;
				mark_implicit_builtin(StorageClassInput, BuiltInPrimitiveId, var.self);
				has_primitive_id = true;
			}

			if (need_tese_params && builtin == BuiltInTessLevelOuter)
			{
				tess_level_outer_var_id = var.self;
			}

			if (need_tese_params && builtin == BuiltInTessLevelInner)
			{
				tess_level_inner_var_id = var.self;
			}

			if ((need_subgroup_mask || needs_subgroup_invocation_id) && builtin == BuiltInSubgroupLocalInvocationId)
			{
				builtin_subgroup_invocation_id_id = var.self;
				mark_implicit_builtin(StorageClassInput, BuiltInSubgroupLocalInvocationId, var.self);
				has_subgroup_invocation_id = true;
			}

			if ((need_subgroup_ge_mask || needs_subgroup_size) && builtin == BuiltInSubgroupSize)
			{
				builtin_subgroup_size_id = var.self;
				mark_implicit_builtin(StorageClassInput, BuiltInSubgroupSize, var.self);
				has_subgroup_size = true;
			}

			if (need_multiview)
			{
				switch (builtin)
				{
				case BuiltInInstanceIndex:
					// The view index here is derived from the instance index.
					builtin_instance_idx_id = var.self;
					mark_implicit_builtin(StorageClassInput, BuiltInInstanceIndex, var.self);
					has_instance_idx = true;
					break;
				case BuiltInBaseInstance:
					// If a non-zero base instance is used, we need to adjust for it when calculating the view index.
					builtin_base_instance_id = var.self;
					mark_implicit_builtin(StorageClassInput, BuiltInBaseInstance, var.self);
					has_base_instance = true;
					break;
				case BuiltInViewIndex:
					builtin_view_idx_id = var.self;
					mark_implicit_builtin(StorageClassInput, BuiltInViewIndex, var.self);
					has_view_idx = true;
					break;
				default:
					break;
				}
			}

			if (needs_helper_invocation && builtin == BuiltInHelperInvocation)
			{
				builtin_helper_invocation_id = var.self;
				mark_implicit_builtin(StorageClassInput, BuiltInHelperInvocation, var.self);
				has_helper_invocation = true;
			}

			if (need_local_invocation_index && builtin == BuiltInLocalInvocationIndex)
			{
				builtin_local_invocation_index_id = var.self;
				mark_implicit_builtin(StorageClassInput, BuiltInLocalInvocationIndex, var.self);
				has_local_invocation_index = true;
			}

			if (need_workgroup_size && builtin == BuiltInLocalInvocationId)
			{
				builtin_workgroup_size_id = var.self;
				mark_implicit_builtin(StorageClassInput, BuiltInWorkgroupSize, var.self);
				has_workgroup_size = true;
			}

			// The base workgroup needs to have the same type and vector size
			// as the workgroup or invocation ID, so keep track of the type that
			// was used.
			if (need_dispatch_base && workgroup_id_type == 0 &&
			    (builtin == BuiltInWorkgroupId || builtin == BuiltInGlobalInvocationId))
				workgroup_id_type = var.basetype;
		});

		// Use Metal's native frame-buffer fetch API for subpass inputs.
		if ((!has_frag_coord || (msl_options.multiview && !has_view_idx) ||
		     (msl_options.arrayed_subpass_input && !msl_options.multiview && !has_layer)) &&
		    (!msl_options.use_framebuffer_fetch_subpasses) && need_subpass_input)
		{
			if (!has_frag_coord)
			{
				uint32_t offset = ir.increase_bound_by(3);
				uint32_t type_id = offset;
				uint32_t type_ptr_id = offset + 1;
				uint32_t var_id = offset + 2;

				// Create gl_FragCoord.
				SPIRType vec4_type;
				vec4_type.basetype = SPIRType::Float;
				vec4_type.width = 32;
				vec4_type.vecsize = 4;
				set<SPIRType>(type_id, vec4_type);

				SPIRType vec4_type_ptr;
				vec4_type_ptr = vec4_type;
				vec4_type_ptr.pointer = true;
				vec4_type_ptr.pointer_depth++;
				vec4_type_ptr.parent_type = type_id;
				vec4_type_ptr.storage = StorageClassInput;
				auto &ptr_type = set<SPIRType>(type_ptr_id, vec4_type_ptr);
				ptr_type.self = type_id;

				set<SPIRVariable>(var_id, type_ptr_id, StorageClassInput);
				set_decoration(var_id, DecorationBuiltIn, BuiltInFragCoord);
				builtin_frag_coord_id = var_id;
				mark_implicit_builtin(StorageClassInput, BuiltInFragCoord, var_id);
			}

			if (!has_layer && msl_options.arrayed_subpass_input && !msl_options.multiview)
			{
				uint32_t offset = ir.increase_bound_by(2);
				uint32_t type_ptr_id = offset;
				uint32_t var_id = offset + 1;

				// Create gl_Layer.
				SPIRType uint_type_ptr;
				uint_type_ptr = get_uint_type();
				uint_type_ptr.pointer = true;
				uint_type_ptr.pointer_depth++;
				uint_type_ptr.parent_type = get_uint_type_id();
				uint_type_ptr.storage = StorageClassInput;
				auto &ptr_type = set<SPIRType>(type_ptr_id, uint_type_ptr);
				ptr_type.self = get_uint_type_id();

				set<SPIRVariable>(var_id, type_ptr_id, StorageClassInput);
				set_decoration(var_id, DecorationBuiltIn, BuiltInLayer);
				builtin_layer_id = var_id;
				mark_implicit_builtin(StorageClassInput, BuiltInLayer, var_id);
			}

			if (!has_view_idx && msl_options.multiview)
			{
				uint32_t offset = ir.increase_bound_by(2);
				uint32_t type_ptr_id = offset;
				uint32_t var_id = offset + 1;

				// Create gl_ViewIndex.
				SPIRType uint_type_ptr;
				uint_type_ptr = get_uint_type();
				uint_type_ptr.pointer = true;
				uint_type_ptr.pointer_depth++;
				uint_type_ptr.parent_type = get_uint_type_id();
				uint_type_ptr.storage = StorageClassInput;
				auto &ptr_type = set<SPIRType>(type_ptr_id, uint_type_ptr);
				ptr_type.self = get_uint_type_id();

				set<SPIRVariable>(var_id, type_ptr_id, StorageClassInput);
				set_decoration(var_id, DecorationBuiltIn, BuiltInViewIndex);
				builtin_view_idx_id = var_id;
				mark_implicit_builtin(StorageClassInput, BuiltInViewIndex, var_id);
			}
		}

		if (!has_sample_id && (need_sample_pos || needs_sample_id))
		{
			uint32_t offset = ir.increase_bound_by(2);
			uint32_t type_ptr_id = offset;
			uint32_t var_id = offset + 1;

			// Create gl_SampleID.
			SPIRType uint_type_ptr;
			uint_type_ptr = get_uint_type();
			uint_type_ptr.pointer = true;
			uint_type_ptr.pointer_depth++;
			uint_type_ptr.parent_type = get_uint_type_id();
			uint_type_ptr.storage = StorageClassInput;
			auto &ptr_type = set<SPIRType>(type_ptr_id, uint_type_ptr);
			ptr_type.self = get_uint_type_id();

			set<SPIRVariable>(var_id, type_ptr_id, StorageClassInput);
			set_decoration(var_id, DecorationBuiltIn, BuiltInSampleId);
			builtin_sample_id_id = var_id;
			mark_implicit_builtin(StorageClassInput, BuiltInSampleId, var_id);
		}

		if ((need_vertex_params && (!has_vertex_idx || !has_base_vertex || !has_instance_idx || !has_base_instance)) ||
		    (need_multiview && (!has_instance_idx || !has_base_instance || !has_view_idx)))
		{
			uint32_t type_ptr_id = ir.increase_bound_by(1);

			SPIRType uint_type_ptr;
			uint_type_ptr = get_uint_type();
			uint_type_ptr.pointer = true;
			uint_type_ptr.pointer_depth++;
			uint_type_ptr.parent_type = get_uint_type_id();
			uint_type_ptr.storage = StorageClassInput;
			auto &ptr_type = set<SPIRType>(type_ptr_id, uint_type_ptr);
			ptr_type.self = get_uint_type_id();

			if (need_vertex_params && !has_vertex_idx)
			{
				uint32_t var_id = ir.increase_bound_by(1);

				// Create gl_VertexIndex.
				set<SPIRVariable>(var_id, type_ptr_id, StorageClassInput);
				set_decoration(var_id, DecorationBuiltIn, BuiltInVertexIndex);
				builtin_vertex_idx_id = var_id;
				mark_implicit_builtin(StorageClassInput, BuiltInVertexIndex, var_id);
			}

			if (need_vertex_params && !has_base_vertex)
			{
				uint32_t var_id = ir.increase_bound_by(1);

				// Create gl_BaseVertex.
				set<SPIRVariable>(var_id, type_ptr_id, StorageClassInput);
				set_decoration(var_id, DecorationBuiltIn, BuiltInBaseVertex);
				builtin_base_vertex_id = var_id;
				mark_implicit_builtin(StorageClassInput, BuiltInBaseVertex, var_id);
			}

			if (!has_instance_idx) // Needed by both multiview and tessellation
			{
				uint32_t var_id = ir.increase_bound_by(1);

				// Create gl_InstanceIndex.
				set<SPIRVariable>(var_id, type_ptr_id, StorageClassInput);
				set_decoration(var_id, DecorationBuiltIn, BuiltInInstanceIndex);
				builtin_instance_idx_id = var_id;
				mark_implicit_builtin(StorageClassInput, BuiltInInstanceIndex, var_id);
			}

			if (!has_base_instance) // Needed by both multiview and tessellation
			{
				uint32_t var_id = ir.increase_bound_by(1);

				// Create gl_BaseInstance.
				set<SPIRVariable>(var_id, type_ptr_id, StorageClassInput);
				set_decoration(var_id, DecorationBuiltIn, BuiltInBaseInstance);
				builtin_base_instance_id = var_id;
				mark_implicit_builtin(StorageClassInput, BuiltInBaseInstance, var_id);
			}

			if (need_multiview)
			{
				// Multiview shaders are not allowed to write to gl_Layer, ostensibly because
				// it is implicitly written from gl_ViewIndex, but we have to do that explicitly.
				// Note that we can't just abuse gl_ViewIndex for this purpose: it's an input, but
				// gl_Layer is an output in vertex-pipeline shaders.
				uint32_t type_ptr_out_id = ir.increase_bound_by(2);
				SPIRType uint_type_ptr_out;
				uint_type_ptr_out = get_uint_type();
				uint_type_ptr_out.pointer = true;
				uint_type_ptr_out.pointer_depth++;
				uint_type_ptr_out.parent_type = get_uint_type_id();
				uint_type_ptr_out.storage = StorageClassOutput;
				auto &ptr_out_type = set<SPIRType>(type_ptr_out_id, uint_type_ptr_out);
				ptr_out_type.self = get_uint_type_id();
				uint32_t var_id = type_ptr_out_id + 1;
				set<SPIRVariable>(var_id, type_ptr_out_id, StorageClassOutput);
				set_decoration(var_id, DecorationBuiltIn, BuiltInLayer);
				builtin_layer_id = var_id;
				mark_implicit_builtin(StorageClassOutput, BuiltInLayer, var_id);
			}

			if (need_multiview && !has_view_idx)
			{
				uint32_t var_id = ir.increase_bound_by(1);

				// Create gl_ViewIndex.
				set<SPIRVariable>(var_id, type_ptr_id, StorageClassInput);
				set_decoration(var_id, DecorationBuiltIn, BuiltInViewIndex);
				builtin_view_idx_id = var_id;
				mark_implicit_builtin(StorageClassInput, BuiltInViewIndex, var_id);
			}
		}

		if ((need_tesc_params && (msl_options.multi_patch_workgroup || !has_invocation_id || !has_primitive_id)) ||
		    (need_tese_params && !has_primitive_id) || need_grid_params)
		{
			uint32_t type_ptr_id = ir.increase_bound_by(1);

			SPIRType uint_type_ptr;
			uint_type_ptr = get_uint_type();
			uint_type_ptr.pointer = true;
			uint_type_ptr.pointer_depth++;
			uint_type_ptr.parent_type = get_uint_type_id();
			uint_type_ptr.storage = StorageClassInput;
			auto &ptr_type = set<SPIRType>(type_ptr_id, uint_type_ptr);
			ptr_type.self = get_uint_type_id();

			if ((need_tesc_params && msl_options.multi_patch_workgroup) || need_grid_params)
			{
				uint32_t var_id = ir.increase_bound_by(1);

				// Create gl_GlobalInvocationID.
				set<SPIRVariable>(var_id, type_ptr_id, StorageClassInput);
				set_decoration(var_id, DecorationBuiltIn, BuiltInGlobalInvocationId);
				builtin_invocation_id_id = var_id;
				mark_implicit_builtin(StorageClassInput, BuiltInGlobalInvocationId, var_id);
			}
			else if (need_tesc_params && !has_invocation_id)
			{
				uint32_t var_id = ir.increase_bound_by(1);

				// Create gl_InvocationID.
				set<SPIRVariable>(var_id, type_ptr_id, StorageClassInput);
				set_decoration(var_id, DecorationBuiltIn, BuiltInInvocationId);
				builtin_invocation_id_id = var_id;
				mark_implicit_builtin(StorageClassInput, BuiltInInvocationId, var_id);
			}

			if ((need_tesc_params || need_tese_params) && !has_primitive_id)
			{
				uint32_t var_id = ir.increase_bound_by(1);

				// Create gl_PrimitiveID.
				set<SPIRVariable>(var_id, type_ptr_id, StorageClassInput);
				set_decoration(var_id, DecorationBuiltIn, BuiltInPrimitiveId);
				builtin_primitive_id_id = var_id;
				mark_implicit_builtin(StorageClassInput, BuiltInPrimitiveId, var_id);
			}

			if (need_grid_params)
			{
				uint32_t var_id = ir.increase_bound_by(1);

				set<SPIRVariable>(var_id, build_extended_vector_type(get_uint_type_id(), 3), StorageClassInput);
				set_extended_decoration(var_id, SPIRVCrossDecorationBuiltInStageInputSize);
				get_entry_point().interface_variables.push_back(var_id);
				set_name(var_id, "spvStageInputSize");
				builtin_stage_input_size_id = var_id;
			}
		}

		if (!has_subgroup_invocation_id && (need_subgroup_mask || needs_subgroup_invocation_id))
		{
			uint32_t offset = ir.increase_bound_by(2);
			uint32_t type_ptr_id = offset;
			uint32_t var_id = offset + 1;

			// Create gl_SubgroupInvocationID.
			SPIRType uint_type_ptr;
			uint_type_ptr = get_uint_type();
			uint_type_ptr.pointer = true;
			uint_type_ptr.pointer_depth++;
			uint_type_ptr.parent_type = get_uint_type_id();
			uint_type_ptr.storage = StorageClassInput;
			auto &ptr_type = set<SPIRType>(type_ptr_id, uint_type_ptr);
			ptr_type.self = get_uint_type_id();

			set<SPIRVariable>(var_id, type_ptr_id, StorageClassInput);
			set_decoration(var_id, DecorationBuiltIn, BuiltInSubgroupLocalInvocationId);
			builtin_subgroup_invocation_id_id = var_id;
			mark_implicit_builtin(StorageClassInput, BuiltInSubgroupLocalInvocationId, var_id);
		}

		if (!has_subgroup_size && (need_subgroup_ge_mask || needs_subgroup_size))
		{
			uint32_t offset = ir.increase_bound_by(2);
			uint32_t type_ptr_id = offset;
			uint32_t var_id = offset + 1;

			// Create gl_SubgroupSize.
			SPIRType uint_type_ptr;
			uint_type_ptr = get_uint_type();
			uint_type_ptr.pointer = true;
			uint_type_ptr.pointer_depth++;
			uint_type_ptr.parent_type = get_uint_type_id();
			uint_type_ptr.storage = StorageClassInput;
			auto &ptr_type = set<SPIRType>(type_ptr_id, uint_type_ptr);
			ptr_type.self = get_uint_type_id();

			set<SPIRVariable>(var_id, type_ptr_id, StorageClassInput);
			set_decoration(var_id, DecorationBuiltIn, BuiltInSubgroupSize);
			builtin_subgroup_size_id = var_id;
			mark_implicit_builtin(StorageClassInput, BuiltInSubgroupSize, var_id);
		}

		if (need_dispatch_base || need_vertex_base_params)
		{
			if (workgroup_id_type == 0)
				workgroup_id_type = build_extended_vector_type(get_uint_type_id(), 3);
			uint32_t var_id;
			if (msl_options.supports_msl_version(1, 2))
			{
				// If we have MSL 1.2, we can (ab)use the [[grid_origin]] builtin
				// to convey this information and save a buffer slot.
				uint32_t offset = ir.increase_bound_by(1);
				var_id = offset;

				set<SPIRVariable>(var_id, workgroup_id_type, StorageClassInput);
				set_extended_decoration(var_id, SPIRVCrossDecorationBuiltInDispatchBase);
				get_entry_point().interface_variables.push_back(var_id);
			}
			else
			{
				// Otherwise, we need to fall back to a good ol' fashioned buffer.
				uint32_t offset = ir.increase_bound_by(2);
				var_id = offset;
				uint32_t type_id = offset + 1;

				SPIRType var_type = get<SPIRType>(workgroup_id_type);
				var_type.storage = StorageClassUniform;
				set<SPIRType>(type_id, var_type);

				set<SPIRVariable>(var_id, type_id, StorageClassUniform);
				// This should never match anything.
				set_decoration(var_id, DecorationDescriptorSet, ~(5u));
				set_decoration(var_id, DecorationBinding, msl_options.indirect_params_buffer_index);
				set_extended_decoration(var_id, SPIRVCrossDecorationResourceIndexPrimary,
				                        msl_options.indirect_params_buffer_index);
			}
			set_name(var_id, "spvDispatchBase");
			builtin_dispatch_base_id = var_id;
		}

		if (has_additional_fixed_sample_mask() && !does_shader_write_sample_mask)
		{
			uint32_t offset = ir.increase_bound_by(2);
			uint32_t var_id = offset + 1;

			// Create gl_SampleMask.
			SPIRType uint_type_ptr_out;
			uint_type_ptr_out = get_uint_type();
			uint_type_ptr_out.pointer = true;
			uint_type_ptr_out.pointer_depth++;
			uint_type_ptr_out.parent_type = get_uint_type_id();
			uint_type_ptr_out.storage = StorageClassOutput;

			auto &ptr_out_type = set<SPIRType>(offset, uint_type_ptr_out);
			ptr_out_type.self = get_uint_type_id();
			set<SPIRVariable>(var_id, offset, StorageClassOutput);
			set_decoration(var_id, DecorationBuiltIn, BuiltInSampleMask);
			builtin_sample_mask_id = var_id;
			mark_implicit_builtin(StorageClassOutput, BuiltInSampleMask, var_id);
		}

		if (!has_helper_invocation && needs_helper_invocation)
		{
			uint32_t offset = ir.increase_bound_by(3);
			uint32_t type_id = offset;
			uint32_t type_ptr_id = offset + 1;
			uint32_t var_id = offset + 2;

			// Create gl_HelperInvocation.
			SPIRType bool_type;
			bool_type.basetype = SPIRType::Boolean;
			bool_type.width = 8;
			bool_type.vecsize = 1;
			set<SPIRType>(type_id, bool_type);

			SPIRType bool_type_ptr_in;
			bool_type_ptr_in = bool_type;
			bool_type_ptr_in.pointer = true;
			bool_type_ptr_in.pointer_depth++;
			bool_type_ptr_in.parent_type = type_id;
			bool_type_ptr_in.storage = StorageClassInput;

			auto &ptr_in_type = set<SPIRType>(type_ptr_id, bool_type_ptr_in);
			ptr_in_type.self = type_id;
			set<SPIRVariable>(var_id, type_ptr_id, StorageClassInput);
			set_decoration(var_id, DecorationBuiltIn, BuiltInHelperInvocation);
			builtin_helper_invocation_id = var_id;
			mark_implicit_builtin(StorageClassInput, BuiltInHelperInvocation, var_id);
		}

		if (need_local_invocation_index && !has_local_invocation_index)
		{
			uint32_t offset = ir.increase_bound_by(2);
			uint32_t type_ptr_id = offset;
			uint32_t var_id = offset + 1;

			// Create gl_LocalInvocationIndex.
			SPIRType uint_type_ptr;
			uint_type_ptr = get_uint_type();
			uint_type_ptr.pointer = true;
			uint_type_ptr.pointer_depth++;
			uint_type_ptr.parent_type = get_uint_type_id();
			uint_type_ptr.storage = StorageClassInput;

			auto &ptr_type = set<SPIRType>(type_ptr_id, uint_type_ptr);
			ptr_type.self = get_uint_type_id();
			set<SPIRVariable>(var_id, type_ptr_id, StorageClassInput);
			set_decoration(var_id, DecorationBuiltIn, BuiltInLocalInvocationIndex);
			builtin_local_invocation_index_id = var_id;
			mark_implicit_builtin(StorageClassInput, BuiltInLocalInvocationIndex, var_id);
		}

		if (need_workgroup_size && !has_workgroup_size)
		{
			uint32_t offset = ir.increase_bound_by(2);
			uint32_t type_ptr_id = offset;
			uint32_t var_id = offset + 1;

			// Create gl_WorkgroupSize.
			uint32_t type_id = build_extended_vector_type(get_uint_type_id(), 3);
			SPIRType uint_type_ptr = get<SPIRType>(type_id);
			uint_type_ptr.pointer = true;
			uint_type_ptr.pointer_depth++;
			uint_type_ptr.parent_type = type_id;
			uint_type_ptr.storage = StorageClassInput;

			auto &ptr_type = set<SPIRType>(type_ptr_id, uint_type_ptr);
			ptr_type.self = type_id;
			set<SPIRVariable>(var_id, type_ptr_id, StorageClassInput);
			set_decoration(var_id, DecorationBuiltIn, BuiltInWorkgroupSize);
			builtin_workgroup_size_id = var_id;
			mark_implicit_builtin(StorageClassInput, BuiltInWorkgroupSize, var_id);
		}
	}

	if (needs_swizzle_buffer_def)
	{
		uint32_t var_id = build_constant_uint_array_pointer();
		set_name(var_id, "spvSwizzleConstants");
		// This should never match anything.
		set_decoration(var_id, DecorationDescriptorSet, kSwizzleBufferBinding);
		set_decoration(var_id, DecorationBinding, msl_options.swizzle_buffer_index);
		set_extended_decoration(var_id, SPIRVCrossDecorationResourceIndexPrimary, msl_options.swizzle_buffer_index);
		swizzle_buffer_id = var_id;
	}

	if (needs_buffer_size_buffer())
	{
		uint32_t var_id = build_constant_uint_array_pointer();
		set_name(var_id, "spvBufferSizeConstants");
		// This should never match anything.
		set_decoration(var_id, DecorationDescriptorSet, kBufferSizeBufferBinding);
		set_decoration(var_id, DecorationBinding, msl_options.buffer_size_buffer_index);
		set_extended_decoration(var_id, SPIRVCrossDecorationResourceIndexPrimary, msl_options.buffer_size_buffer_index);
		buffer_size_buffer_id = var_id;
	}

	if (needs_view_mask_buffer())
	{
		uint32_t var_id = build_constant_uint_array_pointer();
		set_name(var_id, "spvViewMask");
		// This should never match anything.
		set_decoration(var_id, DecorationDescriptorSet, ~(4u));
		set_decoration(var_id, DecorationBinding, msl_options.view_mask_buffer_index);
		set_extended_decoration(var_id, SPIRVCrossDecorationResourceIndexPrimary, msl_options.view_mask_buffer_index);
		view_mask_buffer_id = var_id;
	}

	if (!buffers_requiring_dynamic_offset.empty())
	{
		uint32_t var_id = build_constant_uint_array_pointer();
		set_name(var_id, "spvDynamicOffsets");
		// This should never match anything.
		set_decoration(var_id, DecorationDescriptorSet, ~(5u));
		set_decoration(var_id, DecorationBinding, msl_options.dynamic_offsets_buffer_index);
		set_extended_decoration(var_id, SPIRVCrossDecorationResourceIndexPrimary,
		                        msl_options.dynamic_offsets_buffer_index);
		dynamic_offsets_buffer_id = var_id;
	}

	// If we're returning a struct from a vertex-like entry point, we must return a position attribute.
	bool need_position = (get_execution_model() == ExecutionModelVertex || is_tese_shader()) &&
	                     !capture_output_to_buffer && !get_is_rasterization_disabled() &&
	                     !active_output_builtins.get(BuiltInPosition);

	if (need_position)
	{
		// If we can get away with returning void from entry point, we don't need to care.
		// If there is at least one other stage output, we need to return [[position]],
		// so we need to create one if it doesn't appear in the SPIR-V. Before adding the
		// implicit variable, check if it actually exists already, but just has not been used
		// or initialized, and if so, mark it as active, and do not create the implicit variable.
		bool has_output = false;
		ir.for_each_typed_id<SPIRVariable>([&](uint32_t, SPIRVariable &var) {
			if (var.storage == StorageClassOutput && interface_variable_exists_in_entry_point(var.self))
			{
				has_output = true;

				// Check if the var is the Position builtin
				if (has_decoration(var.self, DecorationBuiltIn) && get_decoration(var.self, DecorationBuiltIn) == BuiltInPosition)
					active_output_builtins.set(BuiltInPosition);

				// If the var is a struct, check if any members is the Position builtin
				auto &var_type = get_variable_element_type(var);
				if (var_type.basetype == SPIRType::Struct)
				{
					auto mbr_cnt = var_type.member_types.size();
					for (uint32_t mbr_idx = 0; mbr_idx < mbr_cnt; mbr_idx++)
					{
						auto builtin = BuiltInMax;
						bool is_builtin = is_member_builtin(var_type, mbr_idx, &builtin);
						if (is_builtin && builtin == BuiltInPosition)
							active_output_builtins.set(BuiltInPosition);
					}
				}
			}
		});
		need_position = has_output && !active_output_builtins.get(BuiltInPosition);
	}

	if (need_position)
	{
		uint32_t offset = ir.increase_bound_by(3);
		uint32_t type_id = offset;
		uint32_t type_ptr_id = offset + 1;
		uint32_t var_id = offset + 2;

		// Create gl_Position.
		SPIRType vec4_type;
		vec4_type.basetype = SPIRType::Float;
		vec4_type.width = 32;
		vec4_type.vecsize = 4;
		set<SPIRType>(type_id, vec4_type);

		SPIRType vec4_type_ptr;
		vec4_type_ptr = vec4_type;
		vec4_type_ptr.pointer = true;
		vec4_type_ptr.pointer_depth++;
		vec4_type_ptr.parent_type = type_id;
		vec4_type_ptr.storage = StorageClassOutput;
		auto &ptr_type = set<SPIRType>(type_ptr_id, vec4_type_ptr);
		ptr_type.self = type_id;

		set<SPIRVariable>(var_id, type_ptr_id, StorageClassOutput);
		set_decoration(var_id, DecorationBuiltIn, BuiltInPosition);
		mark_implicit_builtin(StorageClassOutput, BuiltInPosition, var_id);
	}
}

// Checks if the specified builtin variable (e.g. gl_InstanceIndex) is marked as active.
// If not, it marks it as active and forces a recompilation.
// This might be used when the optimization of inactive builtins was too optimistic (e.g. when "spvOut" is emitted).
void CompilerMSL::ensure_builtin(spv::StorageClass storage, spv::BuiltIn builtin)
{
	Bitset *active_builtins = nullptr;
	switch (storage)
	{
	case StorageClassInput:
		active_builtins = &active_input_builtins;
		break;

	case StorageClassOutput:
		active_builtins = &active_output_builtins;
		break;

	default:
		break;
	}

	// At this point, the specified builtin variable must have already been declared in the entry point.
	// If not, mark as active and force recompile.
	if (active_builtins != nullptr && !active_builtins->get(builtin))
	{
		active_builtins->set(builtin);
		force_recompile();
	}
}

void CompilerMSL::mark_implicit_builtin(StorageClass storage, BuiltIn builtin, uint32_t id)
{
	Bitset *active_builtins = nullptr;
	switch (storage)
	{
	case StorageClassInput:
		active_builtins = &active_input_builtins;
		break;

	case StorageClassOutput:
		active_builtins = &active_output_builtins;
		break;

	default:
		break;
	}

	assert(active_builtins != nullptr);
	active_builtins->set(builtin);

	auto &var = get_entry_point().interface_variables;
	if (find(begin(var), end(var), VariableID(id)) == end(var))
		var.push_back(id);
}

uint32_t CompilerMSL::build_constant_uint_array_pointer()
{
	uint32_t offset = ir.increase_bound_by(3);
	uint32_t type_ptr_id = offset;
	uint32_t type_ptr_ptr_id = offset + 1;
	uint32_t var_id = offset + 2;

	// Create a buffer to hold extra data, including the swizzle constants.
	SPIRType uint_type_pointer = get_uint_type();
	uint_type_pointer.pointer = true;
	uint_type_pointer.pointer_depth++;
	uint_type_pointer.parent_type = get_uint_type_id();
	uint_type_pointer.storage = StorageClassUniform;
	set<SPIRType>(type_ptr_id, uint_type_pointer);
	set_decoration(type_ptr_id, DecorationArrayStride, 4);

	SPIRType uint_type_pointer2 = uint_type_pointer;
	uint_type_pointer2.pointer_depth++;
	uint_type_pointer2.parent_type = type_ptr_id;
	set<SPIRType>(type_ptr_ptr_id, uint_type_pointer2);

	set<SPIRVariable>(var_id, type_ptr_ptr_id, StorageClassUniformConstant);
	return var_id;
}

static string create_sampler_address(const char *prefix, MSLSamplerAddress addr)
{
	switch (addr)
	{
	case MSL_SAMPLER_ADDRESS_CLAMP_TO_EDGE:
		return join(prefix, "address::clamp_to_edge");
	case MSL_SAMPLER_ADDRESS_CLAMP_TO_ZERO:
		return join(prefix, "address::clamp_to_zero");
	case MSL_SAMPLER_ADDRESS_CLAMP_TO_BORDER:
		return join(prefix, "address::clamp_to_border");
	case MSL_SAMPLER_ADDRESS_REPEAT:
		return join(prefix, "address::repeat");
	case MSL_SAMPLER_ADDRESS_MIRRORED_REPEAT:
		return join(prefix, "address::mirrored_repeat");
	default:
		SPIRV_CROSS_THROW("Invalid sampler addressing mode.");
	}
}

SPIRType &CompilerMSL::get_stage_in_struct_type()
{
	auto &si_var = get<SPIRVariable>(stage_in_var_id);
	return get_variable_data_type(si_var);
}

SPIRType &CompilerMSL::get_stage_out_struct_type()
{
	auto &so_var = get<SPIRVariable>(stage_out_var_id);
	return get_variable_data_type(so_var);
}

SPIRType &CompilerMSL::get_patch_stage_in_struct_type()
{
	auto &si_var = get<SPIRVariable>(patch_stage_in_var_id);
	return get_variable_data_type(si_var);
}

SPIRType &CompilerMSL::get_patch_stage_out_struct_type()
{
	auto &so_var = get<SPIRVariable>(patch_stage_out_var_id);
	return get_variable_data_type(so_var);
}

std::string CompilerMSL::get_tess_factor_struct_name()
{
	if (is_tessellating_triangles())
		return "MTLTriangleTessellationFactorsHalf";
	return "MTLQuadTessellationFactorsHalf";
}

SPIRType &CompilerMSL::get_uint_type()
{
	return get<SPIRType>(get_uint_type_id());
}

uint32_t CompilerMSL::get_uint_type_id()
{
	if (uint_type_id != 0)
		return uint_type_id;

	uint_type_id = ir.increase_bound_by(1);

	SPIRType type;
	type.basetype = SPIRType::UInt;
	type.width = 32;
	set<SPIRType>(uint_type_id, type);
	return uint_type_id;
}

void CompilerMSL::emit_entry_point_declarations()
{
	// FIXME: Get test coverage here ...
	// Constant arrays of non-primitive types (i.e. matrices) won't link properly into Metal libraries
	declare_complex_constant_arrays();

	// Emit constexpr samplers here.
	for (auto &samp : constexpr_samplers_by_id)
	{
		auto &var = get<SPIRVariable>(samp.first);
		auto &type = get<SPIRType>(var.basetype);
		if (type.basetype == SPIRType::Sampler)
			add_resource_name(samp.first);

		SmallVector<string> args;
		auto &s = samp.second;

		if (s.coord != MSL_SAMPLER_COORD_NORMALIZED)
			args.push_back("coord::pixel");

		if (s.min_filter == s.mag_filter)
		{
			if (s.min_filter != MSL_SAMPLER_FILTER_NEAREST)
				args.push_back("filter::linear");
		}
		else
		{
			if (s.min_filter != MSL_SAMPLER_FILTER_NEAREST)
				args.push_back("min_filter::linear");
			if (s.mag_filter != MSL_SAMPLER_FILTER_NEAREST)
				args.push_back("mag_filter::linear");
		}

		switch (s.mip_filter)
		{
		case MSL_SAMPLER_MIP_FILTER_NONE:
			// Default
			break;
		case MSL_SAMPLER_MIP_FILTER_NEAREST:
			args.push_back("mip_filter::nearest");
			break;
		case MSL_SAMPLER_MIP_FILTER_LINEAR:
			args.push_back("mip_filter::linear");
			break;
		default:
			SPIRV_CROSS_THROW("Invalid mip filter.");
		}

		if (s.s_address == s.t_address && s.s_address == s.r_address)
		{
			if (s.s_address != MSL_SAMPLER_ADDRESS_CLAMP_TO_EDGE)
				args.push_back(create_sampler_address("", s.s_address));
		}
		else
		{
			if (s.s_address != MSL_SAMPLER_ADDRESS_CLAMP_TO_EDGE)
				args.push_back(create_sampler_address("s_", s.s_address));
			if (s.t_address != MSL_SAMPLER_ADDRESS_CLAMP_TO_EDGE)
				args.push_back(create_sampler_address("t_", s.t_address));
			if (s.r_address != MSL_SAMPLER_ADDRESS_CLAMP_TO_EDGE)
				args.push_back(create_sampler_address("r_", s.r_address));
		}

		if (s.compare_enable)
		{
			switch (s.compare_func)
			{
			case MSL_SAMPLER_COMPARE_FUNC_ALWAYS:
				args.push_back("compare_func::always");
				break;
			case MSL_SAMPLER_COMPARE_FUNC_NEVER:
				args.push_back("compare_func::never");
				break;
			case MSL_SAMPLER_COMPARE_FUNC_EQUAL:
				args.push_back("compare_func::equal");
				break;
			case MSL_SAMPLER_COMPARE_FUNC_NOT_EQUAL:
				args.push_back("compare_func::not_equal");
				break;
			case MSL_SAMPLER_COMPARE_FUNC_LESS:
				args.push_back("compare_func::less");
				break;
			case MSL_SAMPLER_COMPARE_FUNC_LESS_EQUAL:
				args.push_back("compare_func::less_equal");
				break;
			case MSL_SAMPLER_COMPARE_FUNC_GREATER:
				args.push_back("compare_func::greater");
				break;
			case MSL_SAMPLER_COMPARE_FUNC_GREATER_EQUAL:
				args.push_back("compare_func::greater_equal");
				break;
			default:
				SPIRV_CROSS_THROW("Invalid sampler compare function.");
			}
		}

		if (s.s_address == MSL_SAMPLER_ADDRESS_CLAMP_TO_BORDER || s.t_address == MSL_SAMPLER_ADDRESS_CLAMP_TO_BORDER ||
		    s.r_address == MSL_SAMPLER_ADDRESS_CLAMP_TO_BORDER)
		{
			switch (s.border_color)
			{
			case MSL_SAMPLER_BORDER_COLOR_OPAQUE_BLACK:
				args.push_back("border_color::opaque_black");
				break;
			case MSL_SAMPLER_BORDER_COLOR_OPAQUE_WHITE:
				args.push_back("border_color::opaque_white");
				break;
			case MSL_SAMPLER_BORDER_COLOR_TRANSPARENT_BLACK:
				args.push_back("border_color::transparent_black");
				break;
			default:
				SPIRV_CROSS_THROW("Invalid sampler border color.");
			}
		}

		if (s.anisotropy_enable)
			args.push_back(join("max_anisotropy(", s.max_anisotropy, ")"));
		if (s.lod_clamp_enable)
		{
			args.push_back(join("lod_clamp(", convert_to_string(s.lod_clamp_min, current_locale_radix_character), ", ",
			                    convert_to_string(s.lod_clamp_max, current_locale_radix_character), ")"));
		}

		// If we would emit no arguments, then omit the parentheses entirely. Otherwise,
		// we'll wind up with a "most vexing parse" situation.
		if (args.empty())
			statement("constexpr sampler ",
			          type.basetype == SPIRType::SampledImage ? to_sampler_expression(samp.first) : to_name(samp.first),
			          ";");
		else
			statement("constexpr sampler ",
			          type.basetype == SPIRType::SampledImage ? to_sampler_expression(samp.first) : to_name(samp.first),
			          "(", merge(args), ");");
	}

	// Emit dynamic buffers here.
	for (auto &dynamic_buffer : buffers_requiring_dynamic_offset)
	{
		if (!dynamic_buffer.second.second)
		{
			// Could happen if no buffer was used at requested binding point.
			continue;
		}

		const auto &var = get<SPIRVariable>(dynamic_buffer.second.second);
		uint32_t var_id = var.self;
		const auto &type = get_variable_data_type(var);
		string name = to_name(var.self);
		uint32_t desc_set = get_decoration(var.self, DecorationDescriptorSet);
		uint32_t arg_id = argument_buffer_ids[desc_set];
		uint32_t base_index = dynamic_buffer.second.first;

		if (!type.array.empty())
		{
			// This is complicated, because we need to support arrays of arrays.
			// And it's even worse if the outermost dimension is a runtime array, because now
			// all this complicated goop has to go into the shader itself. (FIXME)
			if (!type.array[type.array.size() - 1])
				SPIRV_CROSS_THROW("Runtime arrays with dynamic offsets are not supported yet.");
			else
			{
				is_using_builtin_array = true;
				statement(get_argument_address_space(var), " ", type_to_glsl(type), "* ", to_restrict(var_id, true), name,
				          type_to_array_glsl(type), " =");

				uint32_t dim = uint32_t(type.array.size());
				uint32_t j = 0;
				for (SmallVector<uint32_t> indices(type.array.size());
				     indices[type.array.size() - 1] < to_array_size_literal(type); j++)
				{
					while (dim > 0)
					{
						begin_scope();
						--dim;
					}

					string arrays;
					for (uint32_t i = uint32_t(type.array.size()); i; --i)
						arrays += join("[", indices[i - 1], "]");
					statement("(", get_argument_address_space(var), " ", type_to_glsl(type), "* ",
					          to_restrict(var_id, false), ")((", get_argument_address_space(var), " char* ",
					          to_restrict(var_id, false), ")", to_name(arg_id), ".", ensure_valid_name(name, "m"),
					          arrays, " + ", to_name(dynamic_offsets_buffer_id), "[", base_index + j, "]),");

					while (++indices[dim] >= to_array_size_literal(type, dim) && dim < type.array.size() - 1)
					{
						end_scope(",");
						indices[dim++] = 0;
					}
				}
				end_scope_decl();
				statement_no_indent("");
				is_using_builtin_array = false;
			}
		}
		else
		{
			statement(get_argument_address_space(var), " auto& ", to_restrict(var_id, true), name, " = *(",
			          get_argument_address_space(var), " ", type_to_glsl(type), "* ", to_restrict(var_id, false), ")((",
			          get_argument_address_space(var), " char* ", to_restrict(var_id, false), ")", to_name(arg_id), ".",
			          ensure_valid_name(name, "m"), " + ", to_name(dynamic_offsets_buffer_id), "[", base_index, "]);");
		}
	}

	// Emit buffer arrays here.
	for (uint32_t array_id : buffer_arrays_discrete)
	{
		const auto &var = get<SPIRVariable>(array_id);
		const auto &type = get_variable_data_type(var);
		const auto &buffer_type = get_variable_element_type(var);
		string name = to_name(array_id);
		statement(get_argument_address_space(var), " ", type_to_glsl(buffer_type), "* ", to_restrict(array_id, true), name,
		          "[] =");
		begin_scope();
		for (uint32_t i = 0; i < to_array_size_literal(type); ++i)
			statement(name, "_", i, ",");
		end_scope_decl();
		statement_no_indent("");
	}
	// Discrete descriptors are processed in entry point emission every compiler iteration.
	buffer_arrays_discrete.clear();

	// Emit buffer aliases here.
	for (auto &var_id : buffer_aliases_discrete)
	{
		const auto &var = get<SPIRVariable>(var_id);
		const auto &type = get_variable_data_type(var);
		auto addr_space = get_argument_address_space(var);
		auto name = to_name(var_id);

		uint32_t desc_set = get_decoration(var_id, DecorationDescriptorSet);
		uint32_t desc_binding = get_decoration(var_id, DecorationBinding);
		auto alias_name = join("spvBufferAliasSet", desc_set, "Binding", desc_binding);

		statement(addr_space, " auto& ", to_restrict(var_id, true),
		          name,
		          " = *(", addr_space, " ", type_to_glsl(type), "*)", alias_name, ";");
	}
	// Discrete descriptors are processed in entry point emission every compiler iteration.
	buffer_aliases_discrete.clear();

	for (auto &var_pair : buffer_aliases_argument)
	{
		uint32_t var_id = var_pair.first;
		uint32_t alias_id = var_pair.second;

		const auto &var = get<SPIRVariable>(var_id);
		const auto &type = get_variable_data_type(var);
		auto addr_space = get_argument_address_space(var);

		if (type.array.empty())
		{
			statement(addr_space, " auto& ", to_restrict(var_id, true), to_name(var_id), " = (", addr_space, " ",
			          type_to_glsl(type), "&)", ir.meta[alias_id].decoration.qualified_alias, ";");
		}
		else
		{
			const char *desc_addr_space = descriptor_address_space(var_id, var.storage, "thread");

			// Esoteric type cast. Reference to array of pointers.
			// Auto here defers to UBO or SSBO. The address space of the reference needs to refer to the
			// address space of the argument buffer itself, which is usually constant, but can be const device for
			// large argument buffers.
			is_using_builtin_array = true;
			statement(desc_addr_space, " auto& ", to_restrict(var_id, true), to_name(var_id), " = (", addr_space, " ",
			          type_to_glsl(type), "* ", desc_addr_space, " (&)",
			          type_to_array_glsl(type), ")", ir.meta[alias_id].decoration.qualified_alias, ";");
			is_using_builtin_array = false;
		}
	}

	// Emit disabled fragment outputs.
	std::sort(disabled_frag_outputs.begin(), disabled_frag_outputs.end());
	for (uint32_t var_id : disabled_frag_outputs)
	{
		auto &var = get<SPIRVariable>(var_id);
		add_local_variable_name(var_id);
		statement(variable_decl(var), ";");
		var.deferred_declaration = false;
	}
}

string CompilerMSL::compile()
{
	replace_illegal_entry_point_names();
	ir.fixup_reserved_names();

	// Do not deal with GLES-isms like precision, older extensions and such.
	options.vulkan_semantics = true;
	options.es = false;
	options.version = 450;
	backend.null_pointer_literal = "nullptr";
	backend.float_literal_suffix = false;
	backend.uint32_t_literal_suffix = true;
	backend.int16_t_literal_suffix = "";
	backend.uint16_t_literal_suffix = "";
	backend.basic_int_type = "int";
	backend.basic_uint_type = "uint";
	backend.basic_int8_type = "char";
	backend.basic_uint8_type = "uchar";
	backend.basic_int16_type = "short";
	backend.basic_uint16_type = "ushort";
	backend.boolean_mix_function = "select";
	backend.swizzle_is_function = false;
	backend.shared_is_implied = false;
	backend.use_initializer_list = true;
	backend.use_typed_initializer_list = true;
	backend.native_row_major_matrix = false;
	backend.unsized_array_supported = false;
	backend.can_declare_arrays_inline = false;
	backend.allow_truncated_access_chain = true;
	backend.comparison_image_samples_scalar = true;
	backend.native_pointers = true;
	backend.nonuniform_qualifier = "";
	backend.support_small_type_sampling_result = true;
	backend.supports_empty_struct = true;
	backend.support_64bit_switch = true;

	// Allow Metal to use the array<T> template unless we force it off.
	backend.can_return_array = !msl_options.force_native_arrays;
	backend.array_is_value_type = !msl_options.force_native_arrays;
	// Arrays which are part of buffer objects are never considered to be value types (just plain C-style).
	backend.array_is_value_type_in_buffer_blocks = false;
	backend.support_pointer_to_pointer = true;
	backend.implicit_c_integer_promotion_rules = true;

	capture_output_to_buffer = msl_options.capture_output_to_buffer;
	is_rasterization_disabled = msl_options.disable_rasterization || capture_output_to_buffer;

	// Initialize array here rather than constructor, MSVC 2013 workaround.
	for (auto &id : next_metal_resource_ids)
		id = 0;

	fixup_anonymous_struct_names();
	fixup_type_alias();
	replace_illegal_names();
	sync_entry_point_aliases_and_names();

	build_function_control_flow_graphs_and_analyze();
	update_active_builtins();
	analyze_image_and_sampler_usage();
	analyze_sampled_image_usage();
	analyze_interlocked_resource_usage();
	preprocess_op_codes();
	build_implicit_builtins();

	if (needs_manual_helper_invocation_updates() &&
	    (active_input_builtins.get(BuiltInHelperInvocation) || needs_helper_invocation))
	{
		string discard_expr =
		    join(builtin_to_glsl(BuiltInHelperInvocation, StorageClassInput), " = true, discard_fragment()");
		backend.discard_literal = discard_expr;
		backend.demote_literal = discard_expr;
	}
	else
	{
		backend.discard_literal = "discard_fragment()";
		backend.demote_literal = "discard_fragment()";
	}

	fixup_image_load_store_access();

	set_enabled_interface_variables(get_active_interface_variables());
	if (msl_options.force_active_argument_buffer_resources)
		activate_argument_buffer_resources();

	if (swizzle_buffer_id)
		add_active_interface_variable(swizzle_buffer_id);
	if (buffer_size_buffer_id)
		add_active_interface_variable(buffer_size_buffer_id);
	if (view_mask_buffer_id)
		add_active_interface_variable(view_mask_buffer_id);
	if (dynamic_offsets_buffer_id)
		add_active_interface_variable(dynamic_offsets_buffer_id);
	if (builtin_layer_id)
		add_active_interface_variable(builtin_layer_id);
	if (builtin_dispatch_base_id && !msl_options.supports_msl_version(1, 2))
		add_active_interface_variable(builtin_dispatch_base_id);
	if (builtin_sample_mask_id)
		add_active_interface_variable(builtin_sample_mask_id);

	// Create structs to hold input, output and uniform variables.
	// Do output first to ensure out. is declared at top of entry function.
	qual_pos_var_name = "";
	stage_out_var_id = add_interface_block(StorageClassOutput);
	patch_stage_out_var_id = add_interface_block(StorageClassOutput, true);
	stage_in_var_id = add_interface_block(StorageClassInput);
	if (is_tese_shader())
		patch_stage_in_var_id = add_interface_block(StorageClassInput, true);

	if (is_tesc_shader())
		stage_out_ptr_var_id = add_interface_block_pointer(stage_out_var_id, StorageClassOutput);
	if (is_tessellation_shader())
		stage_in_ptr_var_id = add_interface_block_pointer(stage_in_var_id, StorageClassInput);

	// Metal vertex functions that define no output must disable rasterization and return void.
	if (!stage_out_var_id)
		is_rasterization_disabled = true;

	// Convert the use of global variables to recursively-passed function parameters
	localize_global_variables();
	extract_global_variables_from_functions();

	// Mark any non-stage-in structs to be tightly packed.
	mark_packable_structs();
	reorder_type_alias();

	// Add fixup hooks required by shader inputs and outputs. This needs to happen before
	// 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
	{
		reset(pass_count);

		// Start bindings at zero.
		next_metal_resource_index_buffer = 0;
		next_metal_resource_index_texture = 0;
		next_metal_resource_index_sampler = 0;
		for (auto &id : next_metal_resource_ids)
			id = 0;

		// Move constructor for this type is broken on GCC 4.9 ...
		buffer.reset();

		emit_header();
		emit_custom_templates();
		emit_custom_functions();
		emit_specialization_constants_and_structs();
		emit_resources();
		emit_function(get<SPIRFunction>(ir.default_entry_point), Bitset());

		pass_count++;
	} while (is_forcing_recompilation());

	return buffer.str();
}

// Register the need to output any custom functions.
void CompilerMSL::preprocess_op_codes()
{
	OpCodePreprocessor preproc(*this);
	traverse_all_reachable_opcodes(get<SPIRFunction>(ir.default_entry_point), preproc);

	suppress_missing_prototypes = preproc.suppress_missing_prototypes;

	if (preproc.uses_atomics)
	{
		add_header_line("#include <metal_atomic>");
		add_pragma_line("#pragma clang diagnostic ignored \"-Wunused-variable\"");
	}

	// Before MSL 2.1 (2.2 for textures), Metal vertex functions that write to
	// resources must disable rasterization and return void.
	if ((preproc.uses_buffer_write && !msl_options.supports_msl_version(2, 1)) ||
	    (preproc.uses_image_write && !msl_options.supports_msl_version(2, 2)))
		is_rasterization_disabled = true;

	// Tessellation control shaders are run as compute functions in Metal, and so
	// must capture their output to a buffer.
	if (is_tesc_shader() || (get_execution_model() == ExecutionModelVertex && msl_options.vertex_for_tessellation))
	{
		is_rasterization_disabled = true;
		capture_output_to_buffer = true;
	}

	if (preproc.needs_subgroup_invocation_id)
		needs_subgroup_invocation_id = true;
	if (preproc.needs_subgroup_size)
		needs_subgroup_size = true;
	// build_implicit_builtins() hasn't run yet, and in fact, this needs to execute
	// before then so that gl_SampleID will get added; so we also need to check if
	// that function would add gl_FragCoord.
	if (preproc.needs_sample_id || msl_options.force_sample_rate_shading ||
	    (is_sample_rate() && (active_input_builtins.get(BuiltInFragCoord) ||
	                          (need_subpass_input_ms && !msl_options.use_framebuffer_fetch_subpasses))))
		needs_sample_id = true;
	if (preproc.needs_helper_invocation)
		needs_helper_invocation = true;

	// OpKill is removed by the parser, so we need to identify those by inspecting
	// blocks.
	ir.for_each_typed_id<SPIRBlock>([&preproc](uint32_t, SPIRBlock &block) {
		if (block.terminator == SPIRBlock::Kill)
			preproc.uses_discard = true;
	});

	// Fragment shaders that both write to storage resources and discard fragments
	// need checks on the writes, to work around Metal allowing these writes despite
	// the fragment being dead.
	if (msl_options.check_discarded_frag_stores && preproc.uses_discard &&
	    (preproc.uses_buffer_write || preproc.uses_image_write))
	{
		frag_shader_needs_discard_checks = true;
		needs_helper_invocation = true;
		// Fragment discard store checks imply manual HelperInvocation updates.
		msl_options.manual_helper_invocation_updates = true;
	}

	if (is_intersection_query())
	{
		add_header_line("#if __METAL_VERSION__ >= 230");
		add_header_line("#include <metal_raytracing>");
		add_header_line("using namespace metal::raytracing;");
		add_header_line("#endif");
	}
}

// Move the Private and Workgroup global variables to the entry function.
// Non-constant variables cannot have global scope in Metal.
void CompilerMSL::localize_global_variables()
{
	auto &entry_func = get<SPIRFunction>(ir.default_entry_point);
	auto iter = global_variables.begin();
	while (iter != global_variables.end())
	{
		uint32_t v_id = *iter;
		auto &var = get<SPIRVariable>(v_id);
		if (var.storage == StorageClassPrivate || var.storage == StorageClassWorkgroup)
		{
			if (!variable_is_lut(var))
				entry_func.add_local_variable(v_id);
			iter = global_variables.erase(iter);
		}
		else
			iter++;
	}
}

// For any global variable accessed directly by a function,
// extract that variable and add it as an argument to that function.
void CompilerMSL::extract_global_variables_from_functions()
{
	// Uniforms
	unordered_set<uint32_t> global_var_ids;
	ir.for_each_typed_id<SPIRVariable>([&](uint32_t, SPIRVariable &var) {
		// Some builtins resolve directly to a function call which does not need any declared variables.
		// Skip these.
		if (var.storage == StorageClassInput && has_decoration(var.self, DecorationBuiltIn))
		{
			auto bi_type = BuiltIn(get_decoration(var.self, DecorationBuiltIn));
			if (bi_type == BuiltInHelperInvocation && !needs_manual_helper_invocation_updates())
				return;
			if (bi_type == BuiltInHelperInvocation && needs_manual_helper_invocation_updates())
			{
				if (msl_options.is_ios() && !msl_options.supports_msl_version(2, 3))
					SPIRV_CROSS_THROW("simd_is_helper_thread() requires version 2.3 on iOS.");
				else if (msl_options.is_macos() && !msl_options.supports_msl_version(2, 1))
					SPIRV_CROSS_THROW("simd_is_helper_thread() requires version 2.1 on macOS.");
				// Make sure this is declared and initialized.
				// Force this to have the proper name.
				set_name(var.self, builtin_to_glsl(BuiltInHelperInvocation, StorageClassInput));
				auto &entry_func = this->get<SPIRFunction>(ir.default_entry_point);
				entry_func.add_local_variable(var.self);
				vars_needing_early_declaration.push_back(var.self);
				entry_func.fixup_hooks_in.push_back([this, &var]()
				                                    { statement(to_name(var.self), " = simd_is_helper_thread();"); });
			}
		}

		if (var.storage == StorageClassInput || var.storage == StorageClassOutput ||
		    var.storage == StorageClassUniform || var.storage == StorageClassUniformConstant ||
		    var.storage == StorageClassPushConstant || var.storage == StorageClassStorageBuffer)
		{
			global_var_ids.insert(var.self);
		}
	});

	// Local vars that are declared in the main function and accessed directly by a function
	auto &entry_func = get<SPIRFunction>(ir.default_entry_point);
	for (auto &var : entry_func.local_variables)
		if (get<SPIRVariable>(var).storage != StorageClassFunction)
			global_var_ids.insert(var);

	std::set<uint32_t> added_arg_ids;
	unordered_set<uint32_t> processed_func_ids;
	extract_global_variables_from_function(ir.default_entry_point, added_arg_ids, global_var_ids, processed_func_ids);
}

// MSL does not support the use of global variables for shader input content.
// For any global variable accessed directly by the specified function, extract that variable,
// add it as an argument to that function, and the arg to the added_arg_ids collection.
void CompilerMSL::extract_global_variables_from_function(uint32_t func_id, std::set<uint32_t> &added_arg_ids,
                                                         unordered_set<uint32_t> &global_var_ids,
                                                         unordered_set<uint32_t> &processed_func_ids)
{
	// Avoid processing a function more than once
	if (processed_func_ids.find(func_id) != processed_func_ids.end())
	{
		// Return function global variables
		added_arg_ids = function_global_vars[func_id];
		return;
	}

	processed_func_ids.insert(func_id);

	auto &func = get<SPIRFunction>(func_id);

	// Recursively establish global args added to functions on which we depend.
	for (auto block : func.blocks)
	{
		auto &b = get<SPIRBlock>(block);
		for (auto &i : b.ops)
		{
			auto ops = stream(i);
			auto op = static_cast<Op>(i.op);

			switch (op)
			{
			case OpLoad:
			case OpInBoundsAccessChain:
			case OpAccessChain:
			case OpPtrAccessChain:
			case OpArrayLength:
			{
				uint32_t base_id = ops[2];
				if (global_var_ids.find(base_id) != global_var_ids.end())
					added_arg_ids.insert(base_id);

				// Use Metal's native frame-buffer fetch API for subpass inputs.
				auto &type = get<SPIRType>(ops[0]);
				if (type.basetype == SPIRType::Image && type.image.dim == DimSubpassData &&
				    (!msl_options.use_framebuffer_fetch_subpasses))
				{
					// Implicitly reads gl_FragCoord.
					assert(builtin_frag_coord_id != 0);
					added_arg_ids.insert(builtin_frag_coord_id);
					if (msl_options.multiview)
					{
						// Implicitly reads gl_ViewIndex.
						assert(builtin_view_idx_id != 0);
						added_arg_ids.insert(builtin_view_idx_id);
					}
					else if (msl_options.arrayed_subpass_input)
					{
						// Implicitly reads gl_Layer.
						assert(builtin_layer_id != 0);
						added_arg_ids.insert(builtin_layer_id);
					}
				}

				break;
			}

			case OpFunctionCall:
			{
				// First see if any of the function call args are globals
				for (uint32_t arg_idx = 3; arg_idx < i.length; arg_idx++)
				{
					uint32_t arg_id = ops[arg_idx];
					if (global_var_ids.find(arg_id) != global_var_ids.end())
						added_arg_ids.insert(arg_id);
				}

				// Then recurse into the function itself to extract globals used internally in the function
				uint32_t inner_func_id = ops[2];
				std::set<uint32_t> inner_func_args;
				extract_global_variables_from_function(inner_func_id, inner_func_args, global_var_ids,
				                                       processed_func_ids);
				added_arg_ids.insert(inner_func_args.begin(), inner_func_args.end());
				break;
			}

			case OpStore:
			{
				uint32_t base_id = ops[0];
				if (global_var_ids.find(base_id) != global_var_ids.end())
					added_arg_ids.insert(base_id);

				uint32_t rvalue_id = ops[1];
				if (global_var_ids.find(rvalue_id) != global_var_ids.end())
					added_arg_ids.insert(rvalue_id);

				if (needs_frag_discard_checks())
					added_arg_ids.insert(builtin_helper_invocation_id);

				break;
			}

			case OpSelect:
			{
				uint32_t base_id = ops[3];
				if (global_var_ids.find(base_id) != global_var_ids.end())
					added_arg_ids.insert(base_id);
				base_id = ops[4];
				if (global_var_ids.find(base_id) != global_var_ids.end())
					added_arg_ids.insert(base_id);
				break;
			}

			case OpAtomicExchange:
			case OpAtomicCompareExchange:
			case OpAtomicStore:
			case OpAtomicIIncrement:
			case OpAtomicIDecrement:
			case OpAtomicIAdd:
			case OpAtomicFAddEXT:
			case OpAtomicISub:
			case OpAtomicSMin:
			case OpAtomicUMin:
			case OpAtomicSMax:
			case OpAtomicUMax:
			case OpAtomicAnd:
			case OpAtomicOr:
			case OpAtomicXor:
			case OpImageWrite:
				if (needs_frag_discard_checks())
					added_arg_ids.insert(builtin_helper_invocation_id);
				break;

			// Emulate texture2D atomic operations
			case OpImageTexelPointer:
			{
				// When using the pointer, we need to know which variable it is actually loaded from.
				uint32_t base_id = ops[2];
				auto *var = maybe_get_backing_variable(base_id);
				if (var && atomic_image_vars.count(var->self))
				{
					if (global_var_ids.find(base_id) != global_var_ids.end())
						added_arg_ids.insert(base_id);
				}
				break;
			}

			case OpExtInst:
			{
				uint32_t extension_set = ops[2];
				if (get<SPIRExtension>(extension_set).ext == SPIRExtension::GLSL)
				{
					auto op_450 = static_cast<GLSLstd450>(ops[3]);
					switch (op_450)
					{
					case GLSLstd450InterpolateAtCentroid:
					case GLSLstd450InterpolateAtSample:
					case GLSLstd450InterpolateAtOffset:
					{
						// For these, we really need the stage-in block. It is theoretically possible to pass the
						// interpolant object, but a) doing so would require us to create an entirely new variable
						// with Interpolant type, and b) if we have a struct or array, handling all the members and
						// elements could get unwieldy fast.
						added_arg_ids.insert(stage_in_var_id);
						break;
					}

					case GLSLstd450Modf:
					case GLSLstd450Frexp:
					{
						uint32_t base_id = ops[5];
						if (global_var_ids.find(base_id) != global_var_ids.end())
							added_arg_ids.insert(base_id);
						break;
					}

					default:
						break;
					}
				}
				break;
			}

			case OpGroupNonUniformInverseBallot:
			{
				added_arg_ids.insert(builtin_subgroup_invocation_id_id);
				break;
			}

			case OpGroupNonUniformBallotFindLSB:
			case OpGroupNonUniformBallotFindMSB:
			{
				added_arg_ids.insert(builtin_subgroup_size_id);
				break;
			}

			case OpGroupNonUniformBallotBitCount:
			{
				auto operation = static_cast<GroupOperation>(ops[3]);
				switch (operation)
				{
				case GroupOperationReduce:
					added_arg_ids.insert(builtin_subgroup_size_id);
					break;
				case GroupOperationInclusiveScan:
				case GroupOperationExclusiveScan:
					added_arg_ids.insert(builtin_subgroup_invocation_id_id);
					break;
				default:
					break;
				}
				break;
			}

			case OpDemoteToHelperInvocation:
				if (needs_manual_helper_invocation_updates() &&
				    (active_input_builtins.get(BuiltInHelperInvocation) || needs_helper_invocation))
					added_arg_ids.insert(builtin_helper_invocation_id);
				break;

			case OpIsHelperInvocationEXT:
				if (needs_manual_helper_invocation_updates())
					added_arg_ids.insert(builtin_helper_invocation_id);
				break;

			case OpRayQueryInitializeKHR:
			case OpRayQueryProceedKHR:
			case OpRayQueryTerminateKHR:
			case OpRayQueryGenerateIntersectionKHR:
			case OpRayQueryConfirmIntersectionKHR:
			{
				// Ray query accesses memory directly, need check pass down object if using Private storage class.
				uint32_t base_id = ops[0];
				if (global_var_ids.find(base_id) != global_var_ids.end())
					added_arg_ids.insert(base_id);
				break;
			}

			case OpRayQueryGetRayTMinKHR:
			case OpRayQueryGetRayFlagsKHR:
			case OpRayQueryGetWorldRayOriginKHR:
			case OpRayQueryGetWorldRayDirectionKHR:
			case OpRayQueryGetIntersectionCandidateAABBOpaqueKHR:
			case OpRayQueryGetIntersectionTypeKHR:
			case OpRayQueryGetIntersectionTKHR:
			case OpRayQueryGetIntersectionInstanceCustomIndexKHR:
			case OpRayQueryGetIntersectionInstanceIdKHR:
			case OpRayQueryGetIntersectionInstanceShaderBindingTableRecordOffsetKHR:
			case OpRayQueryGetIntersectionGeometryIndexKHR:
			case OpRayQueryGetIntersectionPrimitiveIndexKHR:
			case OpRayQueryGetIntersectionBarycentricsKHR:
			case OpRayQueryGetIntersectionFrontFaceKHR:
			case OpRayQueryGetIntersectionObjectRayDirectionKHR:
			case OpRayQueryGetIntersectionObjectRayOriginKHR:
			case OpRayQueryGetIntersectionObjectToWorldKHR:
			case OpRayQueryGetIntersectionWorldToObjectKHR:
			{
				// Ray query accesses memory directly, need check pass down object if using Private storage class.
				uint32_t base_id = ops[2];
				if (global_var_ids.find(base_id) != global_var_ids.end())
					added_arg_ids.insert(base_id);
				break;
			}

			default:
				break;
			}

			if (needs_manual_helper_invocation_updates() && b.terminator == SPIRBlock::Kill &&
			    (active_input_builtins.get(BuiltInHelperInvocation) || needs_helper_invocation))
				added_arg_ids.insert(builtin_helper_invocation_id);

			// TODO: Add all other operations which can affect memory.
			// We should consider a more unified system here to reduce boiler-plate.
			// This kind of analysis is done in several places ...
		}
	}

	function_global_vars[func_id] = added_arg_ids;

	// Add the global variables as arguments to the function
	if (func_id != ir.default_entry_point)
	{
		bool control_point_added_in = false;
		bool control_point_added_out = false;
		bool patch_added_in = false;
		bool patch_added_out = false;

		for (uint32_t arg_id : added_arg_ids)
		{
			auto &var = get<SPIRVariable>(arg_id);
			uint32_t type_id = var.basetype;
			auto *p_type = &get<SPIRType>(type_id);
			BuiltIn bi_type = BuiltIn(get_decoration(arg_id, DecorationBuiltIn));

			bool is_patch = has_decoration(arg_id, DecorationPatch) || is_patch_block(*p_type);
			bool is_block = has_decoration(p_type->self, DecorationBlock);
			bool is_control_point_storage =
			    !is_patch && ((is_tessellation_shader() && var.storage == StorageClassInput) ||
			                  (is_tesc_shader() && var.storage == StorageClassOutput));
			bool is_patch_block_storage = is_patch && is_block && var.storage == StorageClassOutput;
			bool is_builtin = is_builtin_variable(var);
			bool variable_is_stage_io =
					!is_builtin || bi_type == BuiltInPosition || bi_type == BuiltInPointSize ||
					bi_type == BuiltInClipDistance || bi_type == BuiltInCullDistance ||
					p_type->basetype == SPIRType::Struct;
			bool is_redirected_to_global_stage_io = (is_control_point_storage || is_patch_block_storage) &&
			                                        variable_is_stage_io;

			// If output is masked it is not considered part of the global stage IO interface.
			if (is_redirected_to_global_stage_io && var.storage == StorageClassOutput)
				is_redirected_to_global_stage_io = !is_stage_output_variable_masked(var);

			if (is_redirected_to_global_stage_io)
			{
				// Tessellation control shaders see inputs and per-point outputs as arrays.
				// Similarly, tessellation evaluation shaders see per-point inputs as arrays.
				// We collected them into a structure; we must pass the array of this
				// structure to the function.
				std::string name;
				if (is_patch)
					name = var.storage == StorageClassInput ? patch_stage_in_var_name : patch_stage_out_var_name;
				else
					name = var.storage == StorageClassInput ? "gl_in" : "gl_out";

				if (var.storage == StorageClassOutput && has_decoration(p_type->self, DecorationBlock))
				{
					// If we're redirecting a block, we might still need to access the original block
					// variable if we're masking some members.
					for (uint32_t mbr_idx = 0; mbr_idx < uint32_t(p_type->member_types.size()); mbr_idx++)
					{
						if (is_stage_output_block_member_masked(var, mbr_idx, true))
						{
							func.add_parameter(var.basetype, var.self, true);
							break;
						}
					}
				}

				if (var.storage == StorageClassInput)
				{
					auto &added_in = is_patch ? patch_added_in : control_point_added_in;
					if (added_in)
						continue;
					arg_id = is_patch ? patch_stage_in_var_id : stage_in_ptr_var_id;
					added_in = true;
				}
				else if (var.storage == StorageClassOutput)
				{
					auto &added_out = is_patch ? patch_added_out : control_point_added_out;
					if (added_out)
						continue;
					arg_id = is_patch ? patch_stage_out_var_id : stage_out_ptr_var_id;
					added_out = true;
				}

				type_id = get<SPIRVariable>(arg_id).basetype;
				uint32_t next_id = ir.increase_bound_by(1);
				func.add_parameter(type_id, next_id, true);
				set<SPIRVariable>(next_id, type_id, StorageClassFunction, 0, arg_id);

				set_name(next_id, name);
				if (is_tese_shader() && msl_options.raw_buffer_tese_input && var.storage == StorageClassInput)
					set_decoration(next_id, DecorationNonWritable);
			}
			else if (is_builtin && has_decoration(p_type->self, DecorationBlock))
			{
				// Get the pointee type
				type_id = get_pointee_type_id(type_id);
				p_type = &get<SPIRType>(type_id);

				uint32_t mbr_idx = 0;
				for (auto &mbr_type_id : p_type->member_types)
				{
					BuiltIn builtin = BuiltInMax;
					is_builtin = is_member_builtin(*p_type, mbr_idx, &builtin);
					if (is_builtin && has_active_builtin(builtin, var.storage))
					{
						// Add a arg variable with the same type and decorations as the member
						uint32_t next_ids = ir.increase_bound_by(2);
						uint32_t ptr_type_id = next_ids + 0;
						uint32_t var_id = next_ids + 1;

						// Make sure we have an actual pointer type,
						// so that we will get the appropriate address space when declaring these builtins.
						auto &ptr = set<SPIRType>(ptr_type_id, get<SPIRType>(mbr_type_id));
						ptr.self = mbr_type_id;
						ptr.storage = var.storage;
						ptr.pointer = true;
						ptr.pointer_depth++;
						ptr.parent_type = mbr_type_id;

						func.add_parameter(mbr_type_id, var_id, true);
						set<SPIRVariable>(var_id, ptr_type_id, StorageClassFunction);
						ir.meta[var_id].decoration = ir.meta[type_id].members[mbr_idx];
					}
					mbr_idx++;
				}
			}
			else
			{
				uint32_t next_id = ir.increase_bound_by(1);
				func.add_parameter(type_id, next_id, true);
				set<SPIRVariable>(next_id, type_id, StorageClassFunction, 0, arg_id);

				// Ensure the new variable has all the same meta info
				ir.meta[next_id] = ir.meta[arg_id];
			}
		}
	}
}

// For all variables that are some form of non-input-output interface block, mark that all the structs
// that are recursively contained within the type referenced by that variable should be packed tightly.
void CompilerMSL::mark_packable_structs()
{
	ir.for_each_typed_id<SPIRVariable>([&](uint32_t, SPIRVariable &var) {
		if (var.storage != StorageClassFunction && !is_hidden_variable(var))
		{
			auto &type = this->get<SPIRType>(var.basetype);
			if (type.pointer &&
			    (type.storage == StorageClassUniform || type.storage == StorageClassUniformConstant ||
			     type.storage == StorageClassPushConstant || type.storage == StorageClassStorageBuffer) &&
			    (has_decoration(type.self, DecorationBlock) || has_decoration(type.self, DecorationBufferBlock)))
				mark_as_packable(type);
		}

		if (var.storage == StorageClassWorkgroup)
		{
			auto *type = &this->get<SPIRType>(var.basetype);
			if (type->basetype == SPIRType::Struct)
				mark_as_workgroup_struct(*type);
		}
	});

	// Physical storage buffer pointers can appear outside of the context of a variable, if the address
	// is calculated from a ulong or uvec2 and cast to a pointer, so check if they need to be packed too.
	ir.for_each_typed_id<SPIRType>([&](uint32_t, SPIRType &type) {
		if (type.basetype == SPIRType::Struct && type.pointer && type.storage == StorageClassPhysicalStorageBuffer)
			mark_as_packable(type);
	});
}

// If the specified type is a struct, it and any nested structs
// are marked as packable with the SPIRVCrossDecorationBufferBlockRepacked decoration,
void CompilerMSL::mark_as_packable(SPIRType &type)
{
	// If this is not the base type (eg. it's a pointer or array), tunnel down
	if (type.parent_type)
	{
		mark_as_packable(get<SPIRType>(type.parent_type));
		return;
	}

	// Handle possible recursion when a struct contains a pointer to its own type nested somewhere.
	if (type.basetype == SPIRType::Struct && !has_extended_decoration(type.self, SPIRVCrossDecorationBufferBlockRepacked))
	{
		set_extended_decoration(type.self, SPIRVCrossDecorationBufferBlockRepacked);

		// Recurse
		uint32_t mbr_cnt = uint32_t(type.member_types.size());
		for (uint32_t mbr_idx = 0; mbr_idx < mbr_cnt; mbr_idx++)
		{
			uint32_t mbr_type_id = type.member_types[mbr_idx];
			auto &mbr_type = get<SPIRType>(mbr_type_id);
			mark_as_packable(mbr_type);
			if (mbr_type.type_alias)
			{
				auto &mbr_type_alias = get<SPIRType>(mbr_type.type_alias);
				mark_as_packable(mbr_type_alias);
			}
		}
	}
}

// If the specified type is a struct, it and any nested structs
// are marked as used with workgroup storage using the SPIRVCrossDecorationWorkgroupStruct decoration.
void CompilerMSL::mark_as_workgroup_struct(SPIRType &type)
{
	// If this is not the base type (eg. it's a pointer or array), tunnel down
	if (type.parent_type)
	{
		mark_as_workgroup_struct(get<SPIRType>(type.parent_type));
		return;
	}

	// Handle possible recursion when a struct contains a pointer to its own type nested somewhere.
	if (type.basetype == SPIRType::Struct && !has_extended_decoration(type.self, SPIRVCrossDecorationWorkgroupStruct))
	{
		set_extended_decoration(type.self, SPIRVCrossDecorationWorkgroupStruct);

		// Recurse
		uint32_t mbr_cnt = uint32_t(type.member_types.size());
		for (uint32_t mbr_idx = 0; mbr_idx < mbr_cnt; mbr_idx++)
		{
			uint32_t mbr_type_id = type.member_types[mbr_idx];
			auto &mbr_type = get<SPIRType>(mbr_type_id);
			mark_as_workgroup_struct(mbr_type);
			if (mbr_type.type_alias)
			{
				auto &mbr_type_alias = get<SPIRType>(mbr_type.type_alias);
				mark_as_workgroup_struct(mbr_type_alias);
			}
		}
	}
}

// If a shader input exists at the location, it is marked as being used by this shader
void CompilerMSL::mark_location_as_used_by_shader(uint32_t location, const SPIRType &type,
                                                  StorageClass storage, bool fallback)
{
	uint32_t count = type_to_location_count(type);
	switch (storage)
	{
	case StorageClassInput:
		for (uint32_t i = 0; i < count; i++)
		{
			location_inputs_in_use.insert(location + i);
			if (fallback)
				location_inputs_in_use_fallback.insert(location + i);
		}
		break;
	case StorageClassOutput:
		for (uint32_t i = 0; i < count; i++)
		{
			location_outputs_in_use.insert(location + i);
			if (fallback)
				location_outputs_in_use_fallback.insert(location + i);
		}
		break;
	default:
		return;
	}
}

uint32_t CompilerMSL::get_target_components_for_fragment_location(uint32_t location) const
{
	auto itr = fragment_output_components.find(location);
	if (itr == end(fragment_output_components))
		return 4;
	else
		return itr->second;
}

uint32_t CompilerMSL::build_extended_vector_type(uint32_t type_id, uint32_t components, SPIRType::BaseType basetype)
{
	uint32_t new_type_id = ir.increase_bound_by(1);
	auto &old_type = get<SPIRType>(type_id);
	auto *type = &set<SPIRType>(new_type_id, old_type);
	type->vecsize = components;
	if (basetype != SPIRType::Unknown)
		type->basetype = basetype;
	type->self = new_type_id;
	type->parent_type = type_id;
	type->array.clear();
	type->array_size_literal.clear();
	type->pointer = false;

	if (is_array(old_type))
	{
		uint32_t array_type_id = ir.increase_bound_by(1);
		type = &set<SPIRType>(array_type_id, *type);
		type->parent_type = new_type_id;
		type->array = old_type.array;
		type->array_size_literal = old_type.array_size_literal;
		new_type_id = array_type_id;
	}

	if (old_type.pointer)
	{
		uint32_t ptr_type_id = ir.increase_bound_by(1);
		type = &set<SPIRType>(ptr_type_id, *type);
		type->self = new_type_id;
		type->parent_type = new_type_id;
		type->storage = old_type.storage;
		type->pointer = true;
		type->pointer_depth++;
		new_type_id = ptr_type_id;
	}

	return new_type_id;
}

uint32_t CompilerMSL::build_msl_interpolant_type(uint32_t type_id, bool is_noperspective)
{
	uint32_t new_type_id = ir.increase_bound_by(1);
	SPIRType &type = set<SPIRType>(new_type_id, get<SPIRType>(type_id));
	type.basetype = SPIRType::Interpolant;
	type.parent_type = type_id;
	// In Metal, the pull-model interpolant type encodes perspective-vs-no-perspective in the type itself.
	// Add this decoration so we know which argument to pass to the template.
	if (is_noperspective)
		set_decoration(new_type_id, DecorationNoPerspective);
	return new_type_id;
}

bool CompilerMSL::add_component_variable_to_interface_block(spv::StorageClass storage, const std::string &ib_var_ref,
                                                            SPIRVariable &var,
                                                            const SPIRType &type,
                                                            InterfaceBlockMeta &meta)
{
	// Deal with Component decorations.
	const InterfaceBlockMeta::LocationMeta *location_meta = nullptr;
	uint32_t location = ~0u;
	if (has_decoration(var.self, DecorationLocation))
	{
		location = get_decoration(var.self, DecorationLocation);
		auto location_meta_itr = meta.location_meta.find(location);
		if (location_meta_itr != end(meta.location_meta))
			location_meta = &location_meta_itr->second;
	}

	// Check if we need to pad fragment output to match a certain number of components.
	if (location_meta)
	{
		bool pad_fragment_output = has_decoration(var.self, DecorationLocation) &&
		                           msl_options.pad_fragment_output_components &&
		                           get_entry_point().model == ExecutionModelFragment && storage == StorageClassOutput;

		auto &entry_func = get<SPIRFunction>(ir.default_entry_point);
		uint32_t start_component = get_decoration(var.self, DecorationComponent);
		uint32_t type_components = type.vecsize;
		uint32_t num_components = location_meta->num_components;

		if (pad_fragment_output)
		{
			uint32_t locn = get_decoration(var.self, DecorationLocation);
			num_components = max<uint32_t>(num_components, get_target_components_for_fragment_location(locn));
		}

		// We have already declared an IO block member as m_location_N.
		// Just emit an early-declared variable and fixup as needed.
		// Arrays need to be unrolled here since each location might need a different number of components.
		entry_func.add_local_variable(var.self);
		vars_needing_early_declaration.push_back(var.self);

		if (var.storage == StorageClassInput)
		{
			entry_func.fixup_hooks_in.push_back([=, &type, &var]() {
				if (!type.array.empty())
				{
					uint32_t array_size = to_array_size_literal(type);
					for (uint32_t loc_off = 0; loc_off < array_size; loc_off++)
					{
						statement(to_name(var.self), "[", loc_off, "]", " = ", ib_var_ref,
						          ".m_location_", location + loc_off,
						          vector_swizzle(type_components, start_component), ";");
					}
				}
				else
				{
					statement(to_name(var.self), " = ", ib_var_ref, ".m_location_", location,
					          vector_swizzle(type_components, start_component), ";");
				}
			});
		}
		else
		{
			entry_func.fixup_hooks_out.push_back([=, &type, &var]() {
				if (!type.array.empty())
				{
					uint32_t array_size = to_array_size_literal(type);
					for (uint32_t loc_off = 0; loc_off < array_size; loc_off++)
					{
						statement(ib_var_ref, ".m_location_", location + loc_off,
						          vector_swizzle(type_components, start_component), " = ",
						          to_name(var.self), "[", loc_off, "];");
					}
				}
				else
				{
					statement(ib_var_ref, ".m_location_", location,
					          vector_swizzle(type_components, start_component), " = ", to_name(var.self), ";");
				}
			});
		}
		return true;
	}
	else
		return false;
}

void CompilerMSL::add_plain_variable_to_interface_block(StorageClass storage, const string &ib_var_ref,
                                                        SPIRType &ib_type, SPIRVariable &var, InterfaceBlockMeta &meta)
{
	bool is_builtin = is_builtin_variable(var);
	BuiltIn builtin = BuiltIn(get_decoration(var.self, DecorationBuiltIn));
	bool is_flat = has_decoration(var.self, DecorationFlat);
	bool is_noperspective = has_decoration(var.self, DecorationNoPerspective);
	bool is_centroid = has_decoration(var.self, DecorationCentroid);
	bool is_sample = has_decoration(var.self, DecorationSample);

	// Add a reference to the variable type to the interface struct.
	uint32_t ib_mbr_idx = uint32_t(ib_type.member_types.size());
	uint32_t type_id = ensure_correct_builtin_type(var.basetype, builtin);
	var.basetype = type_id;

	type_id = get_pointee_type_id(var.basetype);
	if (meta.strip_array && is_array(get<SPIRType>(type_id)))
		type_id = get<SPIRType>(type_id).parent_type;
	auto &type = get<SPIRType>(type_id);
	uint32_t target_components = 0;
	uint32_t type_components = type.vecsize;

	bool padded_output = false;
	bool padded_input = false;
	uint32_t start_component = 0;

	auto &entry_func = get<SPIRFunction>(ir.default_entry_point);

	if (add_component_variable_to_interface_block(storage, ib_var_ref, var, type, meta))
		return;

	bool pad_fragment_output = has_decoration(var.self, DecorationLocation) &&
	                           msl_options.pad_fragment_output_components &&
	                           get_entry_point().model == ExecutionModelFragment && storage == StorageClassOutput;

	if (pad_fragment_output)
	{
		uint32_t locn = get_decoration(var.self, DecorationLocation);
		target_components = get_target_components_for_fragment_location(locn);
		if (type_components < target_components)
		{
			// Make a new type here.
			type_id = build_extended_vector_type(type_id, target_components);
			padded_output = true;
		}
	}

	if (storage == StorageClassInput && pull_model_inputs.count(var.self))
		ib_type.member_types.push_back(build_msl_interpolant_type(type_id, is_noperspective));
	else
		ib_type.member_types.push_back(type_id);

	// Give the member a name
	string mbr_name = ensure_valid_name(to_expression(var.self), "m");
	set_member_name(ib_type.self, ib_mbr_idx, mbr_name);

	// Update the original variable reference to include the structure reference
	string qual_var_name = ib_var_ref + "." + mbr_name;
	// If using pull-model interpolation, need to add a call to the correct interpolation method.
	if (storage == StorageClassInput && pull_model_inputs.count(var.self))
	{
		if (is_centroid)
			qual_var_name += ".interpolate_at_centroid()";
		else if (is_sample)
			qual_var_name += join(".interpolate_at_sample(", to_expression(builtin_sample_id_id), ")");
		else
			qual_var_name += ".interpolate_at_center()";
	}

	if (padded_output || padded_input)
	{
		entry_func.add_local_variable(var.self);
		vars_needing_early_declaration.push_back(var.self);

		if (padded_output)
		{
			entry_func.fixup_hooks_out.push_back([=, &var]() {
				statement(qual_var_name, vector_swizzle(type_components, start_component), " = ", to_name(var.self),
				          ";");
			});
		}
		else
		{
			entry_func.fixup_hooks_in.push_back([=, &var]() {
				statement(to_name(var.self), " = ", qual_var_name, vector_swizzle(type_components, start_component),
				          ";");
			});
		}
	}
	else if (!meta.strip_array)
		ir.meta[var.self].decoration.qualified_alias = qual_var_name;

	if (var.storage == StorageClassOutput && var.initializer != ID(0))
	{
		if (padded_output || padded_input)
		{
			entry_func.fixup_hooks_in.push_back(
			    [=, &var]() { statement(to_name(var.self), " = ", to_expression(var.initializer), ";"); });
		}
		else
		{
			if (meta.strip_array)
			{
				entry_func.fixup_hooks_in.push_back([=, &var]() {
					uint32_t index = get_extended_decoration(var.self, SPIRVCrossDecorationInterfaceMemberIndex);
					auto invocation = to_tesc_invocation_id();
					statement(to_expression(stage_out_ptr_var_id), "[",
					          invocation, "].",
					          to_member_name(ib_type, index), " = ", to_expression(var.initializer), "[",
					          invocation, "];");
				});
			}
			else
			{
				entry_func.fixup_hooks_in.push_back([=, &var]() {
					statement(qual_var_name, " = ", to_expression(var.initializer), ";");
				});
			}
		}
	}

	// Copy the variable location from the original variable to the member
	if (get_decoration_bitset(var.self).get(DecorationLocation))
	{
		uint32_t locn = get_decoration(var.self, DecorationLocation);
		uint32_t comp = get_decoration(var.self, DecorationComponent);
		if (storage == StorageClassInput)
		{
			type_id = ensure_correct_input_type(var.basetype, locn, comp, 0, meta.strip_array);
			var.basetype = type_id;

			type_id = get_pointee_type_id(type_id);
			if (meta.strip_array && is_array(get<SPIRType>(type_id)))
				type_id = get<SPIRType>(type_id).parent_type;
			if (pull_model_inputs.count(var.self))
				ib_type.member_types[ib_mbr_idx] = build_msl_interpolant_type(type_id, is_noperspective);
			else
				ib_type.member_types[ib_mbr_idx] = type_id;
		}
		set_member_decoration(ib_type.self, ib_mbr_idx, DecorationLocation, locn);
		if (comp)
			set_member_decoration(ib_type.self, ib_mbr_idx, DecorationComponent, comp);
		mark_location_as_used_by_shader(locn, get<SPIRType>(type_id), storage);
	}
	else if (is_builtin && is_tessellation_shader() && storage == StorageClassInput && inputs_by_builtin.count(builtin))
	{
		uint32_t locn = inputs_by_builtin[builtin].location;
		set_member_decoration(ib_type.self, ib_mbr_idx, DecorationLocation, locn);
		mark_location_as_used_by_shader(locn, type, storage);
	}
	else if (is_builtin && capture_output_to_buffer && storage == StorageClassOutput && outputs_by_builtin.count(builtin))
	{
		uint32_t locn = outputs_by_builtin[builtin].location;
		set_member_decoration(ib_type.self, ib_mbr_idx, DecorationLocation, locn);
		mark_location_as_used_by_shader(locn, type, storage);
	}

	if (get_decoration_bitset(var.self).get(DecorationComponent))
	{
		uint32_t component = get_decoration(var.self, DecorationComponent);
		set_member_decoration(ib_type.self, ib_mbr_idx, DecorationComponent, component);
	}

	if (get_decoration_bitset(var.self).get(DecorationIndex))
	{
		uint32_t index = get_decoration(var.self, DecorationIndex);
		set_member_decoration(ib_type.self, ib_mbr_idx, DecorationIndex, index);
	}

	// Mark the member as builtin if needed
	if (is_builtin)
	{
		set_member_decoration(ib_type.self, ib_mbr_idx, DecorationBuiltIn, builtin);
		if (builtin == BuiltInPosition && storage == StorageClassOutput)
			qual_pos_var_name = qual_var_name;
	}

	// Copy interpolation decorations if needed
	if (storage != StorageClassInput || !pull_model_inputs.count(var.self))
	{
		if (is_flat)
			set_member_decoration(ib_type.self, ib_mbr_idx, DecorationFlat);
		if (is_noperspective)
			set_member_decoration(ib_type.self, ib_mbr_idx, DecorationNoPerspective);
		if (is_centroid)
			set_member_decoration(ib_type.self, ib_mbr_idx, DecorationCentroid);
		if (is_sample)
			set_member_decoration(ib_type.self, ib_mbr_idx, DecorationSample);
	}

	set_extended_member_decoration(ib_type.self, ib_mbr_idx, SPIRVCrossDecorationInterfaceOrigID, var.self);
}

void CompilerMSL::add_composite_variable_to_interface_block(StorageClass storage, const string &ib_var_ref,
                                                            SPIRType &ib_type, SPIRVariable &var,
                                                            InterfaceBlockMeta &meta)
{
	auto &entry_func = get<SPIRFunction>(ir.default_entry_point);
	auto &var_type = meta.strip_array ? get_variable_element_type(var) : get_variable_data_type(var);
	uint32_t elem_cnt = 0;

	if (add_component_variable_to_interface_block(storage, ib_var_ref, var, var_type, meta))
		return;

	if (is_matrix(var_type))
	{
		if (is_array(var_type))
			SPIRV_CROSS_THROW("MSL cannot emit arrays-of-matrices in input and output variables.");

		elem_cnt = var_type.columns;
	}
	else if (is_array(var_type))
	{
		if (var_type.array.size() != 1)
			SPIRV_CROSS_THROW("MSL cannot emit arrays-of-arrays in input and output variables.");

		elem_cnt = to_array_size_literal(var_type);
	}

	bool is_builtin = is_builtin_variable(var);
	BuiltIn builtin = BuiltIn(get_decoration(var.self, DecorationBuiltIn));
	bool is_flat = has_decoration(var.self, DecorationFlat);
	bool is_noperspective = has_decoration(var.self, DecorationNoPerspective);
	bool is_centroid = has_decoration(var.self, DecorationCentroid);
	bool is_sample = has_decoration(var.self, DecorationSample);

	auto *usable_type = &var_type;
	if (usable_type->pointer)
		usable_type = &get<SPIRType>(usable_type->parent_type);
	while (is_array(*usable_type) || is_matrix(*usable_type))
		usable_type = &get<SPIRType>(usable_type->parent_type);

	// If a builtin, force it to have the proper name.
	if (is_builtin)
		set_name(var.self, builtin_to_glsl(builtin, StorageClassFunction));

	bool flatten_from_ib_var = false;
	string flatten_from_ib_mbr_name;

	if (storage == StorageClassOutput && is_builtin && builtin == BuiltInClipDistance)
	{
		// Also declare [[clip_distance]] attribute here.
		uint32_t clip_array_mbr_idx = uint32_t(ib_type.member_types.size());
		ib_type.member_types.push_back(get_variable_data_type_id(var));
		set_member_decoration(ib_type.self, clip_array_mbr_idx, DecorationBuiltIn, BuiltInClipDistance);

		flatten_from_ib_mbr_name = builtin_to_glsl(BuiltInClipDistance, StorageClassOutput);
		set_member_name(ib_type.self, clip_array_mbr_idx, flatten_from_ib_mbr_name);

		// When we flatten, we flatten directly from the "out" struct,
		// not from a function variable.
		flatten_from_ib_var = true;

		if (!msl_options.enable_clip_distance_user_varying)
			return;
	}
	else if (!meta.strip_array)
	{
		// Only flatten/unflatten IO composites for non-tessellation cases where arrays are not stripped.
		entry_func.add_local_variable(var.self);
		// We need to declare the variable early and at entry-point scope.
		vars_needing_early_declaration.push_back(var.self);
	}

	for (uint32_t i = 0; i < elem_cnt; i++)
	{
		// Add a reference to the variable type to the interface struct.
		uint32_t ib_mbr_idx = uint32_t(ib_type.member_types.size());

		uint32_t target_components = 0;
		bool padded_output = false;
		uint32_t type_id = usable_type->self;

		// Check if we need to pad fragment output to match a certain number of components.
		if (get_decoration_bitset(var.self).get(DecorationLocation) && msl_options.pad_fragment_output_components &&
		    get_entry_point().model == ExecutionModelFragment && storage == StorageClassOutput)
		{
			uint32_t locn = get_decoration(var.self, DecorationLocation) + i;
			target_components = get_target_components_for_fragment_location(locn);
			if (usable_type->vecsize < target_components)
			{
				// Make a new type here.
				type_id = build_extended_vector_type(usable_type->self, target_components);
				padded_output = true;
			}
		}

		if (storage == StorageClassInput && pull_model_inputs.count(var.self))
			ib_type.member_types.push_back(build_msl_interpolant_type(get_pointee_type_id(type_id), is_noperspective));
		else
			ib_type.member_types.push_back(get_pointee_type_id(type_id));

		// Give the member a name
		string mbr_name = ensure_valid_name(join(to_expression(var.self), "_", i), "m");
		set_member_name(ib_type.self, ib_mbr_idx, mbr_name);

		// There is no qualified alias since we need to flatten the internal array on return.
		if (get_decoration_bitset(var.self).get(DecorationLocation))
		{
			uint32_t locn = get_decoration(var.self, DecorationLocation) + i;
			uint32_t comp = get_decoration(var.self, DecorationComponent);
			if (storage == StorageClassInput)
			{
				var.basetype = ensure_correct_input_type(var.basetype, locn, comp, 0, meta.strip_array);
				uint32_t mbr_type_id = ensure_correct_input_type(usable_type->self, locn, comp, 0, meta.strip_array);
				if (storage == StorageClassInput && pull_model_inputs.count(var.self))
					ib_type.member_types[ib_mbr_idx] = build_msl_interpolant_type(mbr_type_id, is_noperspective);
				else
					ib_type.member_types[ib_mbr_idx] = mbr_type_id;
			}
			set_member_decoration(ib_type.self, ib_mbr_idx, DecorationLocation, locn);
			if (comp)
				set_member_decoration(ib_type.self, ib_mbr_idx, DecorationComponent, comp);
			mark_location_as_used_by_shader(locn, *usable_type, storage);
		}
		else if (is_builtin && is_tessellation_shader() && storage == StorageClassInput && inputs_by_builtin.count(builtin))
		{
			uint32_t locn = inputs_by_builtin[builtin].location + i;
			set_member_decoration(ib_type.self, ib_mbr_idx, DecorationLocation, locn);
			mark_location_as_used_by_shader(locn, *usable_type, storage);
		}
		else if (is_builtin && capture_output_to_buffer && storage == StorageClassOutput && outputs_by_builtin.count(builtin))
		{
			uint32_t locn = outputs_by_builtin[builtin].location + i;
			set_member_decoration(ib_type.self, ib_mbr_idx, DecorationLocation, locn);
			mark_location_as_used_by_shader(locn, *usable_type, storage);
		}
		else if (is_builtin && (builtin == BuiltInClipDistance || builtin == BuiltInCullDistance))
		{
			// Declare the Clip/CullDistance as [[user(clip/cullN)]].
			set_member_decoration(ib_type.self, ib_mbr_idx, DecorationBuiltIn, builtin);
			set_member_decoration(ib_type.self, ib_mbr_idx, DecorationIndex, i);
		}

		if (get_decoration_bitset(var.self).get(DecorationIndex))
		{
			uint32_t index = get_decoration(var.self, DecorationIndex);
			set_member_decoration(ib_type.self, ib_mbr_idx, DecorationIndex, index);
		}

		if (storage != StorageClassInput || !pull_model_inputs.count(var.self))
		{
			// Copy interpolation decorations if needed
			if (is_flat)
				set_member_decoration(ib_type.self, ib_mbr_idx, DecorationFlat);
			if (is_noperspective)
				set_member_decoration(ib_type.self, ib_mbr_idx, DecorationNoPerspective);
			if (is_centroid)
				set_member_decoration(ib_type.self, ib_mbr_idx, DecorationCentroid);
			if (is_sample)
				set_member_decoration(ib_type.self, ib_mbr_idx, DecorationSample);
		}

		set_extended_member_decoration(ib_type.self, ib_mbr_idx, SPIRVCrossDecorationInterfaceOrigID, var.self);

		// Only flatten/unflatten IO composites for non-tessellation cases where arrays are not stripped.
		if (!meta.strip_array)
		{
			switch (storage)
			{
			case StorageClassInput:
				entry_func.fixup_hooks_in.push_back([=, &var]() {
					if (pull_model_inputs.count(var.self))
					{
						string lerp_call;
						if (is_centroid)
							lerp_call = ".interpolate_at_centroid()";
						else if (is_sample)
							lerp_call = join(".interpolate_at_sample(", to_expression(builtin_sample_id_id), ")");
						else
							lerp_call = ".interpolate_at_center()";
						statement(to_name(var.self), "[", i, "] = ", ib_var_ref, ".", mbr_name, lerp_call, ";");
					}
					else
					{
						statement(to_name(var.self), "[", i, "] = ", ib_var_ref, ".", mbr_name, ";");
					}
				});
				break;

			case StorageClassOutput:
				entry_func.fixup_hooks_out.push_back([=, &var]() {
					if (padded_output)
					{
						auto &padded_type = this->get<SPIRType>(type_id);
						statement(
						    ib_var_ref, ".", mbr_name, " = ",
						    remap_swizzle(padded_type, usable_type->vecsize, join(to_name(var.self), "[", i, "]")),
						    ";");