Browse Source

Updated spirv-cross.

Бранимир Караџић 3 years ago
parent
commit
6bbd3aafce

+ 36 - 20
3rdparty/spirv-cross/main.cpp

@@ -77,7 +77,7 @@ struct CLICallbacks
 struct CLIParser
 struct CLIParser
 {
 {
 	CLIParser(CLICallbacks cbs_, int argc_, char *argv_[])
 	CLIParser(CLICallbacks cbs_, int argc_, char *argv_[])
-	    : cbs(move(cbs_))
+	    : cbs(std::move(cbs_))
 	    , argc(argc_)
 	    , argc(argc_)
 	    , argv(argv_)
 	    , argv(argv_)
 	{
 	{
@@ -669,6 +669,8 @@ struct CLIArguments
 	bool emit_line_directives = false;
 	bool emit_line_directives = false;
 	bool enable_storage_image_qualifier_deduction = true;
 	bool enable_storage_image_qualifier_deduction = true;
 	bool force_zero_initialized_variables = false;
 	bool force_zero_initialized_variables = false;
+	bool relax_nan_checks = false;
+	uint32_t force_recompile_max_debug_iterations = 3;
 	SmallVector<uint32_t> msl_discrete_descriptor_sets;
 	SmallVector<uint32_t> msl_discrete_descriptor_sets;
 	SmallVector<uint32_t> msl_device_argument_buffers;
 	SmallVector<uint32_t> msl_device_argument_buffers;
 	SmallVector<pair<uint32_t, uint32_t>> msl_dynamic_buffers;
 	SmallVector<pair<uint32_t, uint32_t>> msl_dynamic_buffers;
@@ -790,6 +792,9 @@ static void print_help_hlsl()
 	// clang-format off
 	// clang-format off
 	fprintf(stderr, "\nHLSL options:\n"
 	fprintf(stderr, "\nHLSL options:\n"
 	                "\t[--shader-model]:\n\t\tEnables a specific shader model, e.g. --shader-model 50 for SM 5.0.\n"
 	                "\t[--shader-model]:\n\t\tEnables a specific shader model, e.g. --shader-model 50 for SM 5.0.\n"
+	                "\t[--flatten-ubo]:\n\t\tEmit UBOs as plain uniform arrays.\n"
+	                "\t\tE.g.: uniform MyUBO { vec4 a; float b, c, d, e; }; will be emitted as uniform float4 MyUBO[2];\n"
+	                "\t\tCaveat: You cannot mix and match floating-point and integer in the same UBO with this option.\n"
 	                "\t[--hlsl-enable-compat]:\n\t\tAllow point size and point coord to be used, even if they won't work as expected.\n"
 	                "\t[--hlsl-enable-compat]:\n\t\tAllow point size and point coord to be used, even if they won't work as expected.\n"
 	                "\t\tPointSize is ignored, and PointCoord returns (0.5, 0.5).\n"
 	                "\t\tPointSize is ignored, and PointCoord returns (0.5, 0.5).\n"
 	                "\t[--hlsl-support-nonzero-basevertex-baseinstance]:\n\t\tSupport base vertex and base instance by emitting a special cbuffer declared as:\n"
 	                "\t[--hlsl-support-nonzero-basevertex-baseinstance]:\n\t\tSupport base vertex and base instance by emitting a special cbuffer declared as:\n"
@@ -915,6 +920,7 @@ static void print_help_common()
 	                "\t[--mask-stage-output-builtin <Position|PointSize|ClipDistance|CullDistance>]:\n"
 	                "\t[--mask-stage-output-builtin <Position|PointSize|ClipDistance|CullDistance>]:\n"
 	                "\t\tIf a stage output variable with matching builtin is active, "
 	                "\t\tIf a stage output variable with matching builtin is active, "
 	                "optimize away the variable if it can affect cross-stage linking correctness.\n"
 	                "optimize away the variable if it can affect cross-stage linking correctness.\n"
+	                "\t[--relax-nan-checks]:\n\t\tRelax NaN checks for N{Clamp,Min,Max} and ordered vs. unordered compare instructions.\n"
 	);
 	);
 	// clang-format on
 	// clang-format on
 }
 }
@@ -932,6 +938,8 @@ static void print_help_obscure()
 	                "\t\tdo not attempt to analyze usage, and always emit read/write state.\n"
 	                "\t\tdo not attempt to analyze usage, and always emit read/write state.\n"
 	                "\t[--flatten-multidimensional-arrays]:\n\t\tDo not support multi-dimensional arrays and flatten them to one dimension.\n"
 	                "\t[--flatten-multidimensional-arrays]:\n\t\tDo not support multi-dimensional arrays and flatten them to one dimension.\n"
 	                "\t[--cpp-interface-name <name>]:\n\t\tEmit a specific class name in C++ codegen.\n"
 	                "\t[--cpp-interface-name <name>]:\n\t\tEmit a specific class name in C++ codegen.\n"
+	                "\t[--force-recompile-max-debug-iterations <count>]:\n\t\tAllow compilation loop to run for N loops.\n"
+	                "\t\tCan be used to triage workarounds, but should not be used as a crutch, since it masks an implementation bug.\n"
 	);
 	);
 	// clang-format on
 	// clang-format on
 }
 }
@@ -1083,7 +1091,7 @@ static HLSLBindingFlags hlsl_resource_type_to_flag(const std::string &arg)
 
 
 static string compile_iteration(const CLIArguments &args, std::vector<uint32_t> spirv_file)
 static string compile_iteration(const CLIArguments &args, std::vector<uint32_t> spirv_file)
 {
 {
-	Parser spirv_parser(move(spirv_file));
+	Parser spirv_parser(std::move(spirv_file));
 	spirv_parser.parse();
 	spirv_parser.parse();
 
 
 	unique_ptr<CompilerGLSL> compiler;
 	unique_ptr<CompilerGLSL> compiler;
@@ -1092,13 +1100,13 @@ static string compile_iteration(const CLIArguments &args, std::vector<uint32_t>
 
 
 	if (args.cpp)
 	if (args.cpp)
 	{
 	{
-		compiler.reset(new CompilerCPP(move(spirv_parser.get_parsed_ir())));
+		compiler.reset(new CompilerCPP(std::move(spirv_parser.get_parsed_ir())));
 		if (args.cpp_interface_name)
 		if (args.cpp_interface_name)
 			static_cast<CompilerCPP *>(compiler.get())->set_interface_name(args.cpp_interface_name);
 			static_cast<CompilerCPP *>(compiler.get())->set_interface_name(args.cpp_interface_name);
 	}
 	}
 	else if (args.msl)
 	else if (args.msl)
 	{
 	{
-		compiler.reset(new CompilerMSL(move(spirv_parser.get_parsed_ir())));
+		compiler.reset(new CompilerMSL(std::move(spirv_parser.get_parsed_ir())));
 
 
 		auto *msl_comp = static_cast<CompilerMSL *>(compiler.get());
 		auto *msl_comp = static_cast<CompilerMSL *>(compiler.get());
 		auto msl_opts = msl_comp->get_msl_options();
 		auto msl_opts = msl_comp->get_msl_options();
@@ -1156,13 +1164,13 @@ static string compile_iteration(const CLIArguments &args, std::vector<uint32_t>
 			msl_comp->set_combined_sampler_suffix(args.msl_combined_sampler_suffix);
 			msl_comp->set_combined_sampler_suffix(args.msl_combined_sampler_suffix);
 	}
 	}
 	else if (args.hlsl)
 	else if (args.hlsl)
-		compiler.reset(new CompilerHLSL(move(spirv_parser.get_parsed_ir())));
+		compiler.reset(new CompilerHLSL(std::move(spirv_parser.get_parsed_ir())));
 	else
 	else
 	{
 	{
 		combined_image_samplers = !args.vulkan_semantics;
 		combined_image_samplers = !args.vulkan_semantics;
 		if (!args.vulkan_semantics || args.vulkan_glsl_disable_ext_samplerless_texture_functions)
 		if (!args.vulkan_semantics || args.vulkan_glsl_disable_ext_samplerless_texture_functions)
 			build_dummy_sampler = true;
 			build_dummy_sampler = true;
-		compiler.reset(new CompilerGLSL(move(spirv_parser.get_parsed_ir())));
+		compiler.reset(new CompilerGLSL(std::move(spirv_parser.get_parsed_ir())));
 	}
 	}
 
 
 	if (!args.variable_type_remaps.empty())
 	if (!args.variable_type_remaps.empty())
@@ -1173,7 +1181,7 @@ static string compile_iteration(const CLIArguments &args, std::vector<uint32_t>
 					out = remap.new_variable_type;
 					out = remap.new_variable_type;
 		};
 		};
 
 
-		compiler->set_variable_type_remap_callback(move(remap_cb));
+		compiler->set_variable_type_remap_callback(std::move(remap_cb));
 	}
 	}
 
 
 	for (auto &masked : args.masked_stage_outputs)
 	for (auto &masked : args.masked_stage_outputs)
@@ -1286,6 +1294,8 @@ static string compile_iteration(const CLIArguments &args, std::vector<uint32_t>
 	opts.emit_line_directives = args.emit_line_directives;
 	opts.emit_line_directives = args.emit_line_directives;
 	opts.enable_storage_image_qualifier_deduction = args.enable_storage_image_qualifier_deduction;
 	opts.enable_storage_image_qualifier_deduction = args.enable_storage_image_qualifier_deduction;
 	opts.force_zero_initialized_variables = args.force_zero_initialized_variables;
 	opts.force_zero_initialized_variables = args.force_zero_initialized_variables;
+	opts.relax_nan_checks = args.relax_nan_checks;
+	opts.force_recompile_max_debug_iterations = args.force_recompile_max_debug_iterations;
 	compiler->set_common_options(opts);
 	compiler->set_common_options(opts);
 
 
 	for (auto &fetch : args.glsl_ext_framebuffer_fetch)
 	for (auto &fetch : args.glsl_ext_framebuffer_fetch)
@@ -1345,7 +1355,7 @@ static string compile_iteration(const CLIArguments &args, std::vector<uint32_t>
 	{
 	{
 		auto active = compiler->get_active_interface_variables();
 		auto active = compiler->get_active_interface_variables();
 		res = compiler->get_shader_resources(active);
 		res = compiler->get_shader_resources(active);
-		compiler->set_enabled_interface_variables(move(active));
+		compiler->set_enabled_interface_variables(std::move(active));
 	}
 	}
 	else
 	else
 		res = compiler->get_shader_resources();
 		res = compiler->get_shader_resources();
@@ -1360,7 +1370,7 @@ static string compile_iteration(const CLIArguments &args, std::vector<uint32_t>
 
 
 	auto pls_inputs = remap_pls(args.pls_in, res.stage_inputs, &res.subpass_inputs);
 	auto pls_inputs = remap_pls(args.pls_in, res.stage_inputs, &res.subpass_inputs);
 	auto pls_outputs = remap_pls(args.pls_out, res.stage_outputs, nullptr);
 	auto pls_outputs = remap_pls(args.pls_out, res.stage_outputs, nullptr);
-	compiler->remap_pixel_local_storage(move(pls_inputs), move(pls_outputs));
+	compiler->remap_pixel_local_storage(std::move(pls_inputs), std::move(pls_outputs));
 
 
 	for (auto &ext : args.extensions)
 	for (auto &ext : args.extensions)
 		compiler->require_extension(ext);
 		compiler->require_extension(ext);
@@ -1589,7 +1599,7 @@ static int main_inner(int argc, char *argv[])
 		auto old_name = parser.next_string();
 		auto old_name = parser.next_string();
 		auto new_name = parser.next_string();
 		auto new_name = parser.next_string();
 		auto model = stage_to_execution_model(parser.next_string());
 		auto model = stage_to_execution_model(parser.next_string());
-		args.entry_point_rename.push_back({ old_name, new_name, move(model) });
+		args.entry_point_rename.push_back({ old_name, new_name, std::move(model) });
 	});
 	});
 	cbs.add("--entry", [&args](CLIParser &parser) { args.entry = parser.next_string(); });
 	cbs.add("--entry", [&args](CLIParser &parser) { args.entry = parser.next_string(); });
 	cbs.add("--stage", [&args](CLIParser &parser) { args.entry_stage = parser.next_string(); });
 	cbs.add("--stage", [&args](CLIParser &parser) { args.entry_stage = parser.next_string(); });
@@ -1598,20 +1608,20 @@ static int main_inner(int argc, char *argv[])
 		HLSLVertexAttributeRemap remap;
 		HLSLVertexAttributeRemap remap;
 		remap.location = parser.next_uint();
 		remap.location = parser.next_uint();
 		remap.semantic = parser.next_string();
 		remap.semantic = parser.next_string();
-		args.hlsl_attr_remap.push_back(move(remap));
+		args.hlsl_attr_remap.push_back(std::move(remap));
 	});
 	});
 
 
 	cbs.add("--remap", [&args](CLIParser &parser) {
 	cbs.add("--remap", [&args](CLIParser &parser) {
 		string src = parser.next_string();
 		string src = parser.next_string();
 		string dst = parser.next_string();
 		string dst = parser.next_string();
 		uint32_t components = parser.next_uint();
 		uint32_t components = parser.next_uint();
-		args.remaps.push_back({ move(src), move(dst), components });
+		args.remaps.push_back({ std::move(src), std::move(dst), components });
 	});
 	});
 
 
 	cbs.add("--remap-variable-type", [&args](CLIParser &parser) {
 	cbs.add("--remap-variable-type", [&args](CLIParser &parser) {
 		string var_name = parser.next_string();
 		string var_name = parser.next_string();
 		string new_type = parser.next_string();
 		string new_type = parser.next_string();
-		args.variable_type_remaps.push_back({ move(var_name), move(new_type) });
+		args.variable_type_remaps.push_back({ std::move(var_name), std::move(new_type) });
 	});
 	});
 
 
 	cbs.add("--rename-interface-variable", [&args](CLIParser &parser) {
 	cbs.add("--rename-interface-variable", [&args](CLIParser &parser) {
@@ -1624,18 +1634,18 @@ static int main_inner(int argc, char *argv[])
 
 
 		uint32_t loc = parser.next_uint();
 		uint32_t loc = parser.next_uint();
 		string var_name = parser.next_string();
 		string var_name = parser.next_string();
-		args.interface_variable_renames.push_back({ cls, loc, move(var_name) });
+		args.interface_variable_renames.push_back({ cls, loc, std::move(var_name) });
 	});
 	});
 
 
 	cbs.add("--pls-in", [&args](CLIParser &parser) {
 	cbs.add("--pls-in", [&args](CLIParser &parser) {
 		auto fmt = pls_format(parser.next_string());
 		auto fmt = pls_format(parser.next_string());
 		auto name = parser.next_string();
 		auto name = parser.next_string();
-		args.pls_in.push_back({ move(fmt), move(name) });
+		args.pls_in.push_back({ std::move(fmt), std::move(name) });
 	});
 	});
 	cbs.add("--pls-out", [&args](CLIParser &parser) {
 	cbs.add("--pls-out", [&args](CLIParser &parser) {
 		auto fmt = pls_format(parser.next_string());
 		auto fmt = pls_format(parser.next_string());
 		auto name = parser.next_string();
 		auto name = parser.next_string();
-		args.pls_out.push_back({ move(fmt), move(name) });
+		args.pls_out.push_back({ std::move(fmt), std::move(name) });
 	});
 	});
 	cbs.add("--shader-model", [&args](CLIParser &parser) {
 	cbs.add("--shader-model", [&args](CLIParser &parser) {
 		args.shader_model = parser.next_uint();
 		args.shader_model = parser.next_uint();
@@ -1678,11 +1688,17 @@ static int main_inner(int argc, char *argv[])
 		args.masked_stage_builtins.push_back(masked_builtin);
 		args.masked_stage_builtins.push_back(masked_builtin);
 	});
 	});
 
 
+	cbs.add("--force-recompile-max-debug-iterations", [&](CLIParser &parser) {
+		args.force_recompile_max_debug_iterations = parser.next_uint();
+	});
+
+	cbs.add("--relax-nan-checks", [&](CLIParser &) { args.relax_nan_checks = true; });
+
 	cbs.default_handler = [&args](const char *value) { args.input = value; };
 	cbs.default_handler = [&args](const char *value) { args.input = value; };
 	cbs.add("-", [&args](CLIParser &) { args.input = "-"; });
 	cbs.add("-", [&args](CLIParser &) { args.input = "-"; });
 	cbs.error_handler = [] { print_help(); };
 	cbs.error_handler = [] { print_help(); };
 
 
-	CLIParser parser{ move(cbs), argc - 1, argv + 1 };
+	CLIParser parser{ std::move(cbs), argc - 1, argv + 1 };
 	if (!parser.parse())
 	if (!parser.parse())
 		return EXIT_FAILURE;
 		return EXIT_FAILURE;
 	else if (parser.ended_state)
 	else if (parser.ended_state)
@@ -1702,10 +1718,10 @@ static int main_inner(int argc, char *argv[])
 	// Special case reflection because it has little to do with the path followed by code-outputting compilers
 	// Special case reflection because it has little to do with the path followed by code-outputting compilers
 	if (!args.reflect.empty())
 	if (!args.reflect.empty())
 	{
 	{
-		Parser spirv_parser(move(spirv_file));
+		Parser spirv_parser(std::move(spirv_file));
 		spirv_parser.parse();
 		spirv_parser.parse();
 
 
-		CompilerReflection compiler(move(spirv_parser.get_parsed_ir()));
+		CompilerReflection compiler(std::move(spirv_parser.get_parsed_ir()));
 		compiler.set_format(args.reflect);
 		compiler.set_format(args.reflect);
 		auto json = compiler.compile();
 		auto json = compiler.compile();
 		if (args.output)
 		if (args.output)
@@ -1718,7 +1734,7 @@ static int main_inner(int argc, char *argv[])
 	string compiled_output;
 	string compiled_output;
 
 
 	if (args.iterations == 1)
 	if (args.iterations == 1)
-		compiled_output = compile_iteration(args, move(spirv_file));
+		compiled_output = compile_iteration(args, std::move(spirv_file));
 	else
 	else
 	{
 	{
 		for (unsigned i = 0; i < args.iterations; i++)
 		for (unsigned i = 0; i < args.iterations; i++)

+ 1 - 2
3rdparty/spirv-cross/spirv_common.hpp

@@ -690,7 +690,7 @@ struct SPIRExpression : IVariant
 
 
 	// Only created by the backend target to avoid creating tons of temporaries.
 	// Only created by the backend target to avoid creating tons of temporaries.
 	SPIRExpression(std::string expr, TypeID expression_type_, bool immutable_)
 	SPIRExpression(std::string expr, TypeID expression_type_, bool immutable_)
-	    : expression(move(expr))
+	    : expression(std::move(expr))
 	    , expression_type(expression_type_)
 	    , expression_type(expression_type_)
 	    , immutable(immutable_)
 	    , immutable(immutable_)
 	{
 	{
@@ -1072,7 +1072,6 @@ struct SPIRVariable : IVariant
 
 
 	// Temporaries which can remain forwarded as long as this variable is not modified.
 	// Temporaries which can remain forwarded as long as this variable is not modified.
 	SmallVector<ID> dependees;
 	SmallVector<ID> dependees;
-	bool forwardable = true;
 
 
 	bool deferred_declaration = false;
 	bool deferred_declaration = false;
 	bool phi_variable = false;
 	bool phi_variable = false;

+ 15 - 10
3rdparty/spirv-cross/spirv_cross.cpp

@@ -36,16 +36,16 @@ using namespace SPIRV_CROSS_NAMESPACE;
 
 
 Compiler::Compiler(vector<uint32_t> ir_)
 Compiler::Compiler(vector<uint32_t> ir_)
 {
 {
-	Parser parser(move(ir_));
+	Parser parser(std::move(ir_));
 	parser.parse();
 	parser.parse();
-	set_ir(move(parser.get_parsed_ir()));
+	set_ir(std::move(parser.get_parsed_ir()));
 }
 }
 
 
 Compiler::Compiler(const uint32_t *ir_, size_t word_count)
 Compiler::Compiler(const uint32_t *ir_, size_t word_count)
 {
 {
 	Parser parser(ir_, word_count);
 	Parser parser(ir_, word_count);
 	parser.parse();
 	parser.parse();
-	set_ir(move(parser.get_parsed_ir()));
+	set_ir(std::move(parser.get_parsed_ir()));
 }
 }
 
 
 Compiler::Compiler(const ParsedIR &ir_)
 Compiler::Compiler(const ParsedIR &ir_)
@@ -55,12 +55,12 @@ Compiler::Compiler(const ParsedIR &ir_)
 
 
 Compiler::Compiler(ParsedIR &&ir_)
 Compiler::Compiler(ParsedIR &&ir_)
 {
 {
-	set_ir(move(ir_));
+	set_ir(std::move(ir_));
 }
 }
 
 
 void Compiler::set_ir(ParsedIR &&ir_)
 void Compiler::set_ir(ParsedIR &&ir_)
 {
 {
-	ir = move(ir_);
+	ir = std::move(ir_);
 	parse_fixup();
 	parse_fixup();
 }
 }
 
 
@@ -852,7 +852,7 @@ unordered_set<VariableID> Compiler::get_active_interface_variables() const
 
 
 void Compiler::set_enabled_interface_variables(std::unordered_set<VariableID> active_variables)
 void Compiler::set_enabled_interface_variables(std::unordered_set<VariableID> active_variables)
 {
 {
-	active_interface_variables = move(active_variables);
+	active_interface_variables = std::move(active_variables);
 	check_active_interface_variables = true;
 	check_active_interface_variables = true;
 }
 }
 
 
@@ -1677,6 +1677,11 @@ const SmallVector<SPIRBlock::Case> &Compiler::get_case_list(const SPIRBlock &blo
 		const auto &type = get<SPIRType>(var->basetype);
 		const auto &type = get<SPIRType>(var->basetype);
 		width = type.width;
 		width = type.width;
 	}
 	}
+	else if (const auto *undef = maybe_get<SPIRUndef>(block.condition))
+	{
+		const auto &type = get<SPIRType>(undef->basetype);
+		width = type.width;
+	}
 	else
 	else
 	{
 	{
 		auto search = ir.load_type_width.find(block.condition);
 		auto search = ir.load_type_width.find(block.condition);
@@ -2508,7 +2513,7 @@ void Compiler::CombinedImageSamplerHandler::push_remap_parameters(const SPIRFunc
 	unordered_map<uint32_t, uint32_t> remapping;
 	unordered_map<uint32_t, uint32_t> remapping;
 	for (uint32_t i = 0; i < length; i++)
 	for (uint32_t i = 0; i < length; i++)
 		remapping[func.arguments[i].id] = remap_parameter(args[i]);
 		remapping[func.arguments[i].id] = remap_parameter(args[i]);
-	parameter_remapping.push(move(remapping));
+	parameter_remapping.push(std::move(remapping));
 }
 }
 
 
 void Compiler::CombinedImageSamplerHandler::pop_remap_parameters()
 void Compiler::CombinedImageSamplerHandler::pop_remap_parameters()
@@ -4362,7 +4367,7 @@ void Compiler::analyze_image_and_sampler_usage()
 	handler.dependency_hierarchy.clear();
 	handler.dependency_hierarchy.clear();
 	traverse_all_reachable_opcodes(get<SPIRFunction>(ir.default_entry_point), handler);
 	traverse_all_reachable_opcodes(get<SPIRFunction>(ir.default_entry_point), handler);
 
 
-	comparison_ids = move(handler.comparison_ids);
+	comparison_ids = std::move(handler.comparison_ids);
 	need_subpass_input = handler.need_subpass_input;
 	need_subpass_input = handler.need_subpass_input;
 
 
 	// Forward information from separate images and samplers into combined image samplers.
 	// Forward information from separate images and samplers into combined image samplers.
@@ -4415,7 +4420,7 @@ void Compiler::build_function_control_flow_graphs_and_analyze()
 	CFGBuilder handler(*this);
 	CFGBuilder handler(*this);
 	handler.function_cfgs[ir.default_entry_point].reset(new CFG(*this, get<SPIRFunction>(ir.default_entry_point)));
 	handler.function_cfgs[ir.default_entry_point].reset(new CFG(*this, get<SPIRFunction>(ir.default_entry_point)));
 	traverse_all_reachable_opcodes(get<SPIRFunction>(ir.default_entry_point), handler);
 	traverse_all_reachable_opcodes(get<SPIRFunction>(ir.default_entry_point), handler);
-	function_cfgs = move(handler.function_cfgs);
+	function_cfgs = std::move(handler.function_cfgs);
 	bool single_function = function_cfgs.size() <= 1;
 	bool single_function = function_cfgs.size() <= 1;
 
 
 	for (auto &f : function_cfgs)
 	for (auto &f : function_cfgs)
@@ -5024,7 +5029,7 @@ void Compiler::analyze_non_block_pointer_types()
 	for (auto type : handler.non_block_types)
 	for (auto type : handler.non_block_types)
 		physical_storage_non_block_pointer_types.push_back(type);
 		physical_storage_non_block_pointer_types.push_back(type);
 	sort(begin(physical_storage_non_block_pointer_types), end(physical_storage_non_block_pointer_types));
 	sort(begin(physical_storage_non_block_pointer_types), end(physical_storage_non_block_pointer_types));
-	physical_storage_type_to_alignment = move(handler.physical_block_type_meta);
+	physical_storage_type_to_alignment = std::move(handler.physical_block_type_meta);
 }
 }
 
 
 bool Compiler::InterlockedResourceAccessPrepassHandler::handle(Op op, const uint32_t *, uint32_t)
 bool Compiler::InterlockedResourceAccessPrepassHandler::handle(Op op, const uint32_t *, uint32_t)

+ 10 - 7
3rdparty/spirv-cross/spirv_cross_c.cpp

@@ -251,7 +251,7 @@ spvc_result spvc_context_parse_spirv(spvc_context context, const SpvId *spirv, s
 		pir->context = context;
 		pir->context = context;
 		Parser parser(spirv, word_count);
 		Parser parser(spirv, word_count);
 		parser.parse();
 		parser.parse();
-		pir->parsed = move(parser.get_parsed_ir());
+		pir->parsed = std::move(parser.get_parsed_ir());
 		*parsed_ir = pir.get();
 		*parsed_ir = pir.get();
 		context->allocations.push_back(std::move(pir));
 		context->allocations.push_back(std::move(pir));
 	}
 	}
@@ -283,7 +283,7 @@ spvc_result spvc_context_create_compiler(spvc_context context, spvc_backend back
 		{
 		{
 		case SPVC_BACKEND_NONE:
 		case SPVC_BACKEND_NONE:
 			if (mode == SPVC_CAPTURE_MODE_TAKE_OWNERSHIP)
 			if (mode == SPVC_CAPTURE_MODE_TAKE_OWNERSHIP)
-				comp->compiler.reset(new Compiler(move(parsed_ir->parsed)));
+				comp->compiler.reset(new Compiler(std::move(parsed_ir->parsed)));
 			else if (mode == SPVC_CAPTURE_MODE_COPY)
 			else if (mode == SPVC_CAPTURE_MODE_COPY)
 				comp->compiler.reset(new Compiler(parsed_ir->parsed));
 				comp->compiler.reset(new Compiler(parsed_ir->parsed));
 			break;
 			break;
@@ -291,7 +291,7 @@ spvc_result spvc_context_create_compiler(spvc_context context, spvc_backend back
 #if SPIRV_CROSS_C_API_GLSL
 #if SPIRV_CROSS_C_API_GLSL
 		case SPVC_BACKEND_GLSL:
 		case SPVC_BACKEND_GLSL:
 			if (mode == SPVC_CAPTURE_MODE_TAKE_OWNERSHIP)
 			if (mode == SPVC_CAPTURE_MODE_TAKE_OWNERSHIP)
-				comp->compiler.reset(new CompilerGLSL(move(parsed_ir->parsed)));
+				comp->compiler.reset(new CompilerGLSL(std::move(parsed_ir->parsed)));
 			else if (mode == SPVC_CAPTURE_MODE_COPY)
 			else if (mode == SPVC_CAPTURE_MODE_COPY)
 				comp->compiler.reset(new CompilerGLSL(parsed_ir->parsed));
 				comp->compiler.reset(new CompilerGLSL(parsed_ir->parsed));
 			break;
 			break;
@@ -300,7 +300,7 @@ spvc_result spvc_context_create_compiler(spvc_context context, spvc_backend back
 #if SPIRV_CROSS_C_API_HLSL
 #if SPIRV_CROSS_C_API_HLSL
 		case SPVC_BACKEND_HLSL:
 		case SPVC_BACKEND_HLSL:
 			if (mode == SPVC_CAPTURE_MODE_TAKE_OWNERSHIP)
 			if (mode == SPVC_CAPTURE_MODE_TAKE_OWNERSHIP)
-				comp->compiler.reset(new CompilerHLSL(move(parsed_ir->parsed)));
+				comp->compiler.reset(new CompilerHLSL(std::move(parsed_ir->parsed)));
 			else if (mode == SPVC_CAPTURE_MODE_COPY)
 			else if (mode == SPVC_CAPTURE_MODE_COPY)
 				comp->compiler.reset(new CompilerHLSL(parsed_ir->parsed));
 				comp->compiler.reset(new CompilerHLSL(parsed_ir->parsed));
 			break;
 			break;
@@ -309,7 +309,7 @@ spvc_result spvc_context_create_compiler(spvc_context context, spvc_backend back
 #if SPIRV_CROSS_C_API_MSL
 #if SPIRV_CROSS_C_API_MSL
 		case SPVC_BACKEND_MSL:
 		case SPVC_BACKEND_MSL:
 			if (mode == SPVC_CAPTURE_MODE_TAKE_OWNERSHIP)
 			if (mode == SPVC_CAPTURE_MODE_TAKE_OWNERSHIP)
-				comp->compiler.reset(new CompilerMSL(move(parsed_ir->parsed)));
+				comp->compiler.reset(new CompilerMSL(std::move(parsed_ir->parsed)));
 			else if (mode == SPVC_CAPTURE_MODE_COPY)
 			else if (mode == SPVC_CAPTURE_MODE_COPY)
 				comp->compiler.reset(new CompilerMSL(parsed_ir->parsed));
 				comp->compiler.reset(new CompilerMSL(parsed_ir->parsed));
 			break;
 			break;
@@ -318,7 +318,7 @@ spvc_result spvc_context_create_compiler(spvc_context context, spvc_backend back
 #if SPIRV_CROSS_C_API_CPP
 #if SPIRV_CROSS_C_API_CPP
 		case SPVC_BACKEND_CPP:
 		case SPVC_BACKEND_CPP:
 			if (mode == SPVC_CAPTURE_MODE_TAKE_OWNERSHIP)
 			if (mode == SPVC_CAPTURE_MODE_TAKE_OWNERSHIP)
-				comp->compiler.reset(new CompilerCPP(move(parsed_ir->parsed)));
+				comp->compiler.reset(new CompilerCPP(std::move(parsed_ir->parsed)));
 			else if (mode == SPVC_CAPTURE_MODE_COPY)
 			else if (mode == SPVC_CAPTURE_MODE_COPY)
 				comp->compiler.reset(new CompilerCPP(parsed_ir->parsed));
 				comp->compiler.reset(new CompilerCPP(parsed_ir->parsed));
 			break;
 			break;
@@ -327,7 +327,7 @@ spvc_result spvc_context_create_compiler(spvc_context context, spvc_backend back
 #if SPIRV_CROSS_C_API_REFLECT
 #if SPIRV_CROSS_C_API_REFLECT
 		case SPVC_BACKEND_JSON:
 		case SPVC_BACKEND_JSON:
 			if (mode == SPVC_CAPTURE_MODE_TAKE_OWNERSHIP)
 			if (mode == SPVC_CAPTURE_MODE_TAKE_OWNERSHIP)
-				comp->compiler.reset(new CompilerReflection(move(parsed_ir->parsed)));
+				comp->compiler.reset(new CompilerReflection(std::move(parsed_ir->parsed)));
 			else if (mode == SPVC_CAPTURE_MODE_COPY)
 			else if (mode == SPVC_CAPTURE_MODE_COPY)
 				comp->compiler.reset(new CompilerReflection(parsed_ir->parsed));
 				comp->compiler.reset(new CompilerReflection(parsed_ir->parsed));
 			break;
 			break;
@@ -475,6 +475,9 @@ spvc_result spvc_compiler_options_set_uint(spvc_compiler_options options, spvc_c
 	case SPVC_COMPILER_OPTION_GLSL_OVR_MULTIVIEW_VIEW_COUNT:
 	case SPVC_COMPILER_OPTION_GLSL_OVR_MULTIVIEW_VIEW_COUNT:
 		options->glsl.ovr_multiview_view_count = value;
 		options->glsl.ovr_multiview_view_count = value;
 		break;
 		break;
+	case SPVC_COMPILER_OPTION_RELAX_NAN_CHECKS:
+		options->glsl.relax_nan_checks = value != 0;
+		break;
 #endif
 #endif
 
 
 #if SPIRV_CROSS_C_API_HLSL
 #if SPIRV_CROSS_C_API_HLSL

+ 3 - 1
3rdparty/spirv-cross/spirv_cross_c.h

@@ -40,7 +40,7 @@ extern "C" {
 /* Bumped if ABI or API breaks backwards compatibility. */
 /* Bumped if ABI or API breaks backwards compatibility. */
 #define SPVC_C_API_VERSION_MAJOR 0
 #define SPVC_C_API_VERSION_MAJOR 0
 /* Bumped if APIs or enumerations are added in a backwards compatible way. */
 /* Bumped if APIs or enumerations are added in a backwards compatible way. */
-#define SPVC_C_API_VERSION_MINOR 48
+#define SPVC_C_API_VERSION_MINOR 49
 /* Bumped if internal implementation details change. */
 /* Bumped if internal implementation details change. */
 #define SPVC_C_API_VERSION_PATCH 0
 #define SPVC_C_API_VERSION_PATCH 0
 
 
@@ -677,6 +677,8 @@ typedef enum spvc_compiler_option
 
 
 	SPVC_COMPILER_OPTION_GLSL_OVR_MULTIVIEW_VIEW_COUNT = 77 | SPVC_COMPILER_OPTION_GLSL_BIT,
 	SPVC_COMPILER_OPTION_GLSL_OVR_MULTIVIEW_VIEW_COUNT = 77 | SPVC_COMPILER_OPTION_GLSL_BIT,
 
 
+	SPVC_COMPILER_OPTION_RELAX_NAN_CHECKS = 78 | SPVC_COMPILER_OPTION_COMMON_BIT,
+
 	SPVC_COMPILER_OPTION_INT_MAX = 0x7fffffff
 	SPVC_COMPILER_OPTION_INT_MAX = 0x7fffffff
 } spvc_compiler_option;
 } spvc_compiler_option;
 
 

+ 14 - 14
3rdparty/spirv-cross/spirv_cross_parsed_ir.cpp

@@ -54,26 +54,26 @@ ParsedIR::ParsedIR()
 // Should have been default-implemented, but need this on MSVC 2013.
 // Should have been default-implemented, but need this on MSVC 2013.
 ParsedIR::ParsedIR(ParsedIR &&other) SPIRV_CROSS_NOEXCEPT
 ParsedIR::ParsedIR(ParsedIR &&other) SPIRV_CROSS_NOEXCEPT
 {
 {
-	*this = move(other);
+	*this = std::move(other);
 }
 }
 
 
 ParsedIR &ParsedIR::operator=(ParsedIR &&other) SPIRV_CROSS_NOEXCEPT
 ParsedIR &ParsedIR::operator=(ParsedIR &&other) SPIRV_CROSS_NOEXCEPT
 {
 {
 	if (this != &other)
 	if (this != &other)
 	{
 	{
-		pool_group = move(other.pool_group);
-		spirv = move(other.spirv);
-		meta = move(other.meta);
+		pool_group = std::move(other.pool_group);
+		spirv = std::move(other.spirv);
+		meta = std::move(other.meta);
 		for (int i = 0; i < TypeCount; i++)
 		for (int i = 0; i < TypeCount; i++)
-			ids_for_type[i] = move(other.ids_for_type[i]);
-		ids_for_constant_or_type = move(other.ids_for_constant_or_type);
-		ids_for_constant_or_variable = move(other.ids_for_constant_or_variable);
-		declared_capabilities = move(other.declared_capabilities);
-		declared_extensions = move(other.declared_extensions);
-		block_meta = move(other.block_meta);
-		continue_block_to_loop_header = move(other.continue_block_to_loop_header);
-		entry_points = move(other.entry_points);
-		ids = move(other.ids);
+			ids_for_type[i] = std::move(other.ids_for_type[i]);
+		ids_for_constant_or_type = std::move(other.ids_for_constant_or_type);
+		ids_for_constant_or_variable = std::move(other.ids_for_constant_or_variable);
+		declared_capabilities = std::move(other.declared_capabilities);
+		declared_extensions = std::move(other.declared_extensions);
+		block_meta = std::move(other.block_meta);
+		continue_block_to_loop_header = std::move(other.continue_block_to_loop_header);
+		entry_points = std::move(other.entry_points);
+		ids = std::move(other.ids);
 		addressing_model = other.addressing_model;
 		addressing_model = other.addressing_model;
 		memory_model = other.memory_model;
 		memory_model = other.memory_model;
 
 
@@ -999,7 +999,7 @@ ParsedIR::LoopLock::LoopLock(uint32_t *lock_)
 
 
 ParsedIR::LoopLock::LoopLock(LoopLock &&other) SPIRV_CROSS_NOEXCEPT
 ParsedIR::LoopLock::LoopLock(LoopLock &&other) SPIRV_CROSS_NOEXCEPT
 {
 {
-	*this = move(other);
+	*this = std::move(other);
 }
 }
 
 
 ParsedIR::LoopLock &ParsedIR::LoopLock::operator=(LoopLock &&other) SPIRV_CROSS_NOEXCEPT
 ParsedIR::LoopLock &ParsedIR::LoopLock::operator=(LoopLock &&other) SPIRV_CROSS_NOEXCEPT

+ 129 - 47
3rdparty/spirv-cross/spirv_glsl.cpp

@@ -306,8 +306,8 @@ void CompilerGLSL::reset(uint32_t iteration_count)
 	// It is highly context-sensitive when we need to force recompilation,
 	// It is highly context-sensitive when we need to force recompilation,
 	// and it is not practical with the current architecture
 	// and it is not practical with the current architecture
 	// to resolve everything up front.
 	// to resolve everything up front.
-	if (iteration_count >= 3 && !is_force_recompile_forward_progress)
-		SPIRV_CROSS_THROW("Over 3 compilation loops detected and no forward progress was made. Must be a bug!");
+	if (iteration_count >= options.force_recompile_max_debug_iterations && !is_force_recompile_forward_progress)
+		SPIRV_CROSS_THROW("Maximum compilation loops detected and no forward progress was made. Must be a SPIRV-Cross bug!");
 
 
 	// We do some speculative optimizations which should pretty much always work out,
 	// We do some speculative optimizations which should pretty much always work out,
 	// but just in case the SPIR-V is rather weird, recompile until it's happy.
 	// but just in case the SPIR-V is rather weird, recompile until it's happy.
@@ -4298,10 +4298,8 @@ string CompilerGLSL::to_func_call_arg(const SPIRFunction::Parameter &, uint32_t
 	return to_expression(name_id);
 	return to_expression(name_id);
 }
 }
 
 
-void CompilerGLSL::handle_invalid_expression(uint32_t id)
+void CompilerGLSL::force_temporary_and_recompile(uint32_t id)
 {
 {
-	// We tried to read an invalidated expression.
-	// This means we need another pass at compilation, but next time, force temporary variables so that they cannot be invalidated.
 	auto res = forced_temporaries.insert(id);
 	auto res = forced_temporaries.insert(id);
 
 
 	// Forcing new temporaries guarantees forward progress.
 	// Forcing new temporaries guarantees forward progress.
@@ -4311,6 +4309,14 @@ void CompilerGLSL::handle_invalid_expression(uint32_t id)
 		force_recompile();
 		force_recompile();
 }
 }
 
 
+void CompilerGLSL::handle_invalid_expression(uint32_t id)
+{
+	// We tried to read an invalidated expression.
+	// This means we need another pass at compilation, but next time,
+	// force temporary variables so that they cannot be invalidated.
+	force_temporary_and_recompile(id);
+}
+
 // Converts the format of the current expression from packed to unpacked,
 // Converts the format of the current expression from packed to unpacked,
 // by wrapping the expression in a constructor of the appropriate type.
 // by wrapping the expression in a constructor of the appropriate type.
 // GLSL does not support packed formats, so simply return the expression.
 // GLSL does not support packed formats, so simply return the expression.
@@ -7237,18 +7243,11 @@ string CompilerGLSL::to_function_args(const TextureFunctionArguments &args, bool
 			forward = forward && should_forward(args.lod);
 			forward = forward && should_forward(args.lod);
 			farg_str += ", ";
 			farg_str += ", ";
 
 
-			auto &lod_expr_type = expression_type(args.lod);
-
 			// Lod expression for TexelFetch in GLSL must be int, and only int.
 			// Lod expression for TexelFetch in GLSL must be int, and only int.
-			if (args.base.is_fetch && imgtype.image.dim != DimBuffer && !imgtype.image.ms &&
-			    lod_expr_type.basetype != SPIRType::Int)
-			{
-				farg_str += join("int(", to_expression(args.lod), ")");
-			}
+			if (args.base.is_fetch && imgtype.image.dim != DimBuffer && !imgtype.image.ms)
+				farg_str += bitcast_expression(SPIRType::Int, args.lod);
 			else
 			else
-			{
 				farg_str += to_expression(args.lod);
 				farg_str += to_expression(args.lod);
-			}
 		}
 		}
 	}
 	}
 	else if (args.base.is_fetch && imgtype.image.dim != DimBuffer && !imgtype.image.ms)
 	else if (args.base.is_fetch && imgtype.image.dim != DimBuffer && !imgtype.image.ms)
@@ -7261,19 +7260,19 @@ string CompilerGLSL::to_function_args(const TextureFunctionArguments &args, bool
 	{
 	{
 		forward = forward && should_forward(args.coffset);
 		forward = forward && should_forward(args.coffset);
 		farg_str += ", ";
 		farg_str += ", ";
-		farg_str += to_expression(args.coffset);
+		farg_str += bitcast_expression(SPIRType::Int, args.coffset);
 	}
 	}
 	else if (args.offset)
 	else if (args.offset)
 	{
 	{
 		forward = forward && should_forward(args.offset);
 		forward = forward && should_forward(args.offset);
 		farg_str += ", ";
 		farg_str += ", ";
-		farg_str += to_expression(args.offset);
+		farg_str += bitcast_expression(SPIRType::Int, args.offset);
 	}
 	}
 
 
 	if (args.sample)
 	if (args.sample)
 	{
 	{
 		farg_str += ", ";
 		farg_str += ", ";
-		farg_str += to_expression(args.sample);
+		farg_str += bitcast_expression(SPIRType::Int, args.sample);
 	}
 	}
 
 
 	if (args.min_lod)
 	if (args.min_lod)
@@ -7300,11 +7299,7 @@ string CompilerGLSL::to_function_args(const TextureFunctionArguments &args, bool
 	{
 	{
 		forward = forward && should_forward(args.component);
 		forward = forward && should_forward(args.component);
 		farg_str += ", ";
 		farg_str += ", ";
-		auto &component_type = expression_type(args.component);
-		if (component_type.basetype == SPIRType::Int)
-			farg_str += to_expression(args.component);
-		else
-			farg_str += join("int(", to_expression(args.component), ")");
+		farg_str += bitcast_expression(SPIRType::Int, args.component);
 	}
 	}
 
 
 	*p_forward = forward;
 	*p_forward = forward;
@@ -7312,6 +7307,63 @@ string CompilerGLSL::to_function_args(const TextureFunctionArguments &args, bool
 	return farg_str;
 	return farg_str;
 }
 }
 
 
+Op CompilerGLSL::get_remapped_spirv_op(Op op) const
+{
+	if (options.relax_nan_checks)
+	{
+		switch (op)
+		{
+		case OpFUnordLessThan:
+			op = OpFOrdLessThan;
+			break;
+		case OpFUnordLessThanEqual:
+			op = OpFOrdLessThanEqual;
+			break;
+		case OpFUnordGreaterThan:
+			op = OpFOrdGreaterThan;
+			break;
+		case OpFUnordGreaterThanEqual:
+			op = OpFOrdGreaterThanEqual;
+			break;
+		case OpFUnordEqual:
+			op = OpFOrdEqual;
+			break;
+		case OpFOrdNotEqual:
+			op = OpFUnordNotEqual;
+			break;
+
+		default:
+			break;
+		}
+	}
+
+	return op;
+}
+
+GLSLstd450 CompilerGLSL::get_remapped_glsl_op(GLSLstd450 std450_op) const
+{
+	// Relax to non-NaN aware opcodes.
+	if (options.relax_nan_checks)
+	{
+		switch (std450_op)
+		{
+		case GLSLstd450NClamp:
+			std450_op = GLSLstd450FClamp;
+			break;
+		case GLSLstd450NMin:
+			std450_op = GLSLstd450FMin;
+			break;
+		case GLSLstd450NMax:
+			std450_op = GLSLstd450FMax;
+			break;
+		default:
+			break;
+		}
+	}
+
+	return std450_op;
+}
+
 void CompilerGLSL::emit_glsl_op(uint32_t result_type, uint32_t id, uint32_t eop, const uint32_t *args, uint32_t length)
 void CompilerGLSL::emit_glsl_op(uint32_t result_type, uint32_t id, uint32_t eop, const uint32_t *args, uint32_t length)
 {
 {
 	auto op = static_cast<GLSLstd450>(eop);
 	auto op = static_cast<GLSLstd450>(eop);
@@ -7324,6 +7376,8 @@ void CompilerGLSL::emit_glsl_op(uint32_t result_type, uint32_t id, uint32_t eop,
 	auto int_type = to_signed_basetype(integer_width);
 	auto int_type = to_signed_basetype(integer_width);
 	auto uint_type = to_unsigned_basetype(integer_width);
 	auto uint_type = to_unsigned_basetype(integer_width);
 
 
+	op = get_remapped_glsl_op(op);
+
 	switch (op)
 	switch (op)
 	{
 	{
 	// FP fiddling
 	// FP fiddling
@@ -9189,8 +9243,13 @@ std::string CompilerGLSL::flattened_access_chain_struct(uint32_t base, const uin
 {
 {
 	std::string expr;
 	std::string expr;
 
 
-	expr += type_to_glsl_constructor(target_type);
-	expr += "(";
+	if (backend.can_declare_struct_inline)
+	{
+		expr += type_to_glsl_constructor(target_type);
+		expr += "(";
+	}
+	else
+		expr += "{";
 
 
 	for (uint32_t i = 0; i < uint32_t(target_type.member_types.size()); ++i)
 	for (uint32_t i = 0; i < uint32_t(target_type.member_types.size()); ++i)
 	{
 	{
@@ -9220,7 +9279,7 @@ std::string CompilerGLSL::flattened_access_chain_struct(uint32_t base, const uin
 			expr += tmp;
 			expr += tmp;
 	}
 	}
 
 
-	expr += ")";
+	expr += backend.can_declare_struct_inline ? ")" : "}";
 
 
 	return expr;
 	return expr;
 }
 }
@@ -9526,8 +9585,11 @@ bool CompilerGLSL::should_forward(uint32_t id) const
 	// This is important because otherwise we'll get local sampler copies (highp sampler2D foo = bar) that are invalid in OpenGL GLSL
 	// This is important because otherwise we'll get local sampler copies (highp sampler2D foo = bar) that are invalid in OpenGL GLSL
 
 
 	auto *var = maybe_get<SPIRVariable>(id);
 	auto *var = maybe_get<SPIRVariable>(id);
-	if (var && var->forwardable)
-		return true;
+	if (var)
+	{
+		// Never forward volatile variables, e.g. SPIR-V 1.6 IsHelperInvocation.
+		return !has_decoration(id, DecorationVolatile);
+	}
 
 
 	// For debugging emit temporary variables for all expressions
 	// For debugging emit temporary variables for all expressions
 	if (options.force_temporary)
 	if (options.force_temporary)
@@ -9540,6 +9602,12 @@ bool CompilerGLSL::should_forward(uint32_t id) const
 	if (expr && expr->expression_dependencies.size() >= max_expression_dependencies)
 	if (expr && expr->expression_dependencies.size() >= max_expression_dependencies)
 		return false;
 		return false;
 
 
+	if (expr && expr->loaded_from && has_decoration(expr->loaded_from, DecorationVolatile))
+	{
+		// Never forward volatile variables.
+		return false;
+	}
+
 	// Immutable expression can always be forwarded.
 	// Immutable expression can always be forwarded.
 	if (is_immutable(id))
 	if (is_immutable(id))
 		return true;
 		return true;
@@ -9596,9 +9664,8 @@ void CompilerGLSL::track_expression_read(uint32_t id)
 			//if (v == 2)
 			//if (v == 2)
 			//    fprintf(stderr, "ID %u was forced to temporary due to more than 1 expression use!\n", id);
 			//    fprintf(stderr, "ID %u was forced to temporary due to more than 1 expression use!\n", id);
 
 
-			forced_temporaries.insert(id);
 			// Force a recompile after this pass to avoid forwarding this variable.
 			// Force a recompile after this pass to avoid forwarding this variable.
-			force_recompile();
+			force_temporary_and_recompile(id);
 		}
 		}
 	}
 	}
 }
 }
@@ -9952,9 +10019,8 @@ void CompilerGLSL::disallow_forwarding_in_expression_chain(const SPIRExpression
 	if (expression_is_forwarded(expr.self) && !expression_suppresses_usage_tracking(expr.self) &&
 	if (expression_is_forwarded(expr.self) && !expression_suppresses_usage_tracking(expr.self) &&
 	    forced_invariant_temporaries.count(expr.self) == 0)
 	    forced_invariant_temporaries.count(expr.self) == 0)
 	{
 	{
-		forced_temporaries.insert(expr.self);
+		force_temporary_and_recompile(expr.self);
 		forced_invariant_temporaries.insert(expr.self);
 		forced_invariant_temporaries.insert(expr.self);
-		force_recompile();
 
 
 		for (auto &dependent : expr.expression_dependencies)
 		for (auto &dependent : expr.expression_dependencies)
 			disallow_forwarding_in_expression_chain(get<SPIRExpression>(dependent));
 			disallow_forwarding_in_expression_chain(get<SPIRExpression>(dependent));
@@ -10096,6 +10162,8 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction)
 	auto int_type = to_signed_basetype(integer_width);
 	auto int_type = to_signed_basetype(integer_width);
 	auto uint_type = to_unsigned_basetype(integer_width);
 	auto uint_type = to_unsigned_basetype(integer_width);
 
 
+	opcode = get_remapped_spirv_op(opcode);
+
 	switch (opcode)
 	switch (opcode)
 	{
 	{
 	// Dealing with memory
 	// Dealing with memory
@@ -10239,9 +10307,19 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction)
 		// If an expression is mutable and forwardable, we speculate that it is immutable.
 		// If an expression is mutable and forwardable, we speculate that it is immutable.
 		AccessChainMeta meta;
 		AccessChainMeta meta;
 		bool ptr_chain = opcode == OpPtrAccessChain;
 		bool ptr_chain = opcode == OpPtrAccessChain;
-		auto e = access_chain(ops[2], &ops[3], length - 3, get<SPIRType>(ops[0]), &meta, ptr_chain);
+		auto &target_type = get<SPIRType>(ops[0]);
+		auto e = access_chain(ops[2], &ops[3], length - 3, target_type, &meta, ptr_chain);
 
 
-		auto &expr = set<SPIRExpression>(ops[1], move(e), ops[0], should_forward(ops[2]));
+		// If the base is flattened UBO of struct type, the expression has to be a composite.
+		// In that case, backends which do not support inline syntax need it to be bound to a temporary.
+		// Otherwise, invalid expressions like ({UBO[0].xyz, UBO[0].w, UBO[1]}).member are emitted.
+		bool requires_temporary = false;
+		if (flattened_buffer_blocks.count(ops[2]) && target_type.basetype == SPIRType::Struct)
+			requires_temporary = !backend.can_declare_struct_inline;
+
+		auto &expr = requires_temporary ?
+                         emit_op(ops[0], ops[1], std::move(e), false) :
+                         set<SPIRExpression>(ops[1], std::move(e), ops[0], should_forward(ops[2]));
 
 
 		auto *backing_variable = maybe_get_backing_variable(ops[2]);
 		auto *backing_variable = maybe_get_backing_variable(ops[2]);
 		expr.loaded_from = backing_variable ? backing_variable->self : ID(ops[2]);
 		expr.loaded_from = backing_variable ? backing_variable->self : ID(ops[2]);
@@ -11212,7 +11290,11 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction)
 
 
 	case OpLogicalNotEqual:
 	case OpLogicalNotEqual:
 	case OpFOrdNotEqual:
 	case OpFOrdNotEqual:
+	case OpFUnordNotEqual:
 	{
 	{
+		// GLSL is fuzzy on what to do with ordered vs unordered not equal.
+		// glslang started emitting UnorderedNotEqual some time ago to harmonize with IEEE,
+		// but this means we have no easy way of implementing ordered not equal.
 		if (expression_type(ops[2]).vecsize > 1)
 		if (expression_type(ops[2]).vecsize > 1)
 			GLSL_BFOP(notEqual);
 			GLSL_BFOP(notEqual);
 		else
 		else
@@ -12016,7 +12098,7 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction)
 			pure = false;
 			pure = false;
 		}
 		}
 
 
-		if (var && var->forwardable)
+		if (var)
 		{
 		{
 			bool forward = forced_temporaries.find(id) == end(forced_temporaries);
 			bool forward = forced_temporaries.find(id) == end(forced_temporaries);
 			auto &e = emit_op(result_type, id, imgexpr, forward);
 			auto &e = emit_op(result_type, id, imgexpr, forward);
@@ -12532,7 +12614,6 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction)
 		break;
 		break;
 
 
 	case OpFUnordEqual:
 	case OpFUnordEqual:
-	case OpFUnordNotEqual:
 	case OpFUnordLessThan:
 	case OpFUnordLessThan:
 	case OpFUnordGreaterThan:
 	case OpFUnordGreaterThan:
 	case OpFUnordLessThanEqual:
 	case OpFUnordLessThanEqual:
@@ -12555,10 +12636,6 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction)
 				comp_op = "notEqual";
 				comp_op = "notEqual";
 				break;
 				break;
 
 
-			case OpFUnordNotEqual:
-				comp_op = "equal";
-				break;
-
 			case OpFUnordLessThan:
 			case OpFUnordLessThan:
 				comp_op = "greaterThanEqual";
 				comp_op = "greaterThanEqual";
 				break;
 				break;
@@ -12591,10 +12668,6 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction)
 				comp_op = " != ";
 				comp_op = " != ";
 				break;
 				break;
 
 
-			case OpFUnordNotEqual:
-				comp_op = " == ";
-				break;
-
 			case OpFUnordLessThan:
 			case OpFUnordLessThan:
 				comp_op = " >= ";
 				comp_op = " >= ";
 				break;
 				break;
@@ -12786,6 +12859,8 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction)
 		if (!options.vulkan_semantics)
 		if (!options.vulkan_semantics)
 			SPIRV_CROSS_THROW("GL_EXT_demote_to_helper_invocation is only supported in Vulkan GLSL.");
 			SPIRV_CROSS_THROW("GL_EXT_demote_to_helper_invocation is only supported in Vulkan GLSL.");
 		require_extension_internal("GL_EXT_demote_to_helper_invocation");
 		require_extension_internal("GL_EXT_demote_to_helper_invocation");
+		// Helper lane state with demote is volatile by nature.
+		// Do not forward this.
 		emit_op(ops[0], ops[1], "helperInvocationEXT()", false);
 		emit_op(ops[0], ops[1], "helperInvocationEXT()", false);
 		break;
 		break;
 
 
@@ -14637,9 +14712,16 @@ void CompilerGLSL::emit_hoisted_temporaries(SmallVector<pair<TypeID, ID>> &tempo
 
 
 	for (auto &tmp : temporaries)
 	for (auto &tmp : temporaries)
 	{
 	{
+		auto &type = get<SPIRType>(tmp.first);
+
+		// There are some rare scenarios where we are asked to declare pointer types as hoisted temporaries.
+		// This should be ignored unless we're doing actual variable pointers and backend supports it.
+		// Access chains cannot normally be lowered to temporaries in GLSL and HLSL.
+		if (type.pointer && !backend.native_pointers)
+			continue;
+
 		add_local_variable_name(tmp.second);
 		add_local_variable_name(tmp.second);
 		auto &flags = ir.meta[tmp.second].decoration.decoration_flags;
 		auto &flags = ir.meta[tmp.second].decoration.decoration_flags;
-		auto &type = get<SPIRType>(tmp.first);
 
 
 		// Not all targets support pointer literals, so don't bother with that case.
 		// Not all targets support pointer literals, so don't bother with that case.
 		string initializer;
 		string initializer;
@@ -15502,7 +15584,7 @@ void CompilerGLSL::unroll_array_from_complex_load(uint32_t target_id, uint32_t s
 			statement(new_expr, "[i] = ", expr, "[i];");
 			statement(new_expr, "[i] = ", expr, "[i];");
 		end_scope();
 		end_scope();
 
 
-		expr = move(new_expr);
+		expr = std::move(new_expr);
 	}
 	}
 }
 }
 
 
@@ -15837,7 +15919,7 @@ void CompilerGLSL::emit_copy_logical_type(uint32_t lhs_id, uint32_t lhs_type_id,
 		rhs_id = id + 1;
 		rhs_id = id + 1;
 
 
 		{
 		{
-			auto &lhs_expr = set<SPIRExpression>(lhs_id, move(lhs), lhs_type_id, true);
+			auto &lhs_expr = set<SPIRExpression>(lhs_id, std::move(lhs), lhs_type_id, true);
 			lhs_expr.need_transpose = lhs_meta.need_transpose;
 			lhs_expr.need_transpose = lhs_meta.need_transpose;
 
 
 			if (lhs_meta.storage_is_packed)
 			if (lhs_meta.storage_is_packed)
@@ -15850,7 +15932,7 @@ void CompilerGLSL::emit_copy_logical_type(uint32_t lhs_id, uint32_t lhs_type_id,
 		}
 		}
 
 
 		{
 		{
-			auto &rhs_expr = set<SPIRExpression>(rhs_id, move(rhs), rhs_type_id, true);
+			auto &rhs_expr = set<SPIRExpression>(rhs_id, std::move(rhs), rhs_type_id, true);
 			rhs_expr.need_transpose = rhs_meta.need_transpose;
 			rhs_expr.need_transpose = rhs_meta.need_transpose;
 
 
 			if (rhs_meta.storage_is_packed)
 			if (rhs_meta.storage_is_packed)

+ 18 - 0
3rdparty/spirv-cross/spirv_glsl.hpp

@@ -83,6 +83,11 @@ public:
 
 
 		// Debug option to always emit temporary variables for all expressions.
 		// Debug option to always emit temporary variables for all expressions.
 		bool force_temporary = false;
 		bool force_temporary = false;
+		// Debug option, can be increased in an attempt to workaround SPIRV-Cross bugs temporarily.
+		// If this limit has to be increased, it points to an implementation bug.
+		// In certain scenarios, the maximum number of debug iterations may increase beyond this limit
+		// as long as we can prove we're making certain kinds of forward progress.
+		uint32_t force_recompile_max_debug_iterations = 3;
 
 
 		// If true, Vulkan GLSL features are used instead of GL-compatible features.
 		// If true, Vulkan GLSL features are used instead of GL-compatible features.
 		// Mostly useful for debugging SPIR-V files.
 		// Mostly useful for debugging SPIR-V files.
@@ -133,6 +138,13 @@ public:
 		// what happens on legacy GLSL targets for blocks and structs.
 		// what happens on legacy GLSL targets for blocks and structs.
 		bool force_flattened_io_blocks = false;
 		bool force_flattened_io_blocks = false;
 
 
+		// For opcodes where we have to perform explicit additional nan checks, very ugly code is generated.
+		// If we opt-in, ignore these requirements.
+		// In opcodes like NClamp/NMin/NMax and FP compare, ignore NaN behavior.
+		// Use FClamp/FMin/FMax semantics for clamps and lets implementation choose ordered or unordered
+		// compares.
+		bool relax_nan_checks = false;
+
 		// If non-zero, controls layout(num_views = N) in; in GL_OVR_multiview2.
 		// If non-zero, controls layout(num_views = N) in; in GL_OVR_multiview2.
 		uint32_t ovr_multiview_view_count = 0;
 		uint32_t ovr_multiview_view_count = 0;
 
 
@@ -357,6 +369,11 @@ protected:
 
 
 	virtual void emit_instruction(const Instruction &instr);
 	virtual void emit_instruction(const Instruction &instr);
 	void emit_block_instructions(SPIRBlock &block);
 	void emit_block_instructions(SPIRBlock &block);
+
+	// For relax_nan_checks.
+	GLSLstd450 get_remapped_glsl_op(GLSLstd450 std450_op) const;
+	spv::Op get_remapped_spirv_op(spv::Op op) const;
+
 	virtual void emit_glsl_op(uint32_t result_type, uint32_t result_id, uint32_t op, const uint32_t *args,
 	virtual void emit_glsl_op(uint32_t result_type, uint32_t result_id, uint32_t op, const uint32_t *args,
 	                          uint32_t count);
 	                          uint32_t count);
 	virtual void emit_spv_amd_shader_ballot_op(uint32_t result_type, uint32_t result_id, uint32_t op,
 	virtual void emit_spv_amd_shader_ballot_op(uint32_t result_type, uint32_t result_id, uint32_t op,
@@ -881,6 +898,7 @@ protected:
 
 
 	void check_function_call_constraints(const uint32_t *args, uint32_t length);
 	void check_function_call_constraints(const uint32_t *args, uint32_t length);
 	void handle_invalid_expression(uint32_t id);
 	void handle_invalid_expression(uint32_t id);
+	void force_temporary_and_recompile(uint32_t id);
 	void find_static_extensions();
 	void find_static_extensions();
 
 
 	std::string emit_for_loop_initializers(const SPIRBlock &block);
 	std::string emit_for_loop_initializers(const SPIRBlock &block);

+ 40 - 9
3rdparty/spirv-cross/spirv_hlsl.cpp

@@ -728,6 +728,11 @@ void CompilerHLSL::emit_builtin_inputs_in_struct()
 			// Handled specially.
 			// Handled specially.
 			break;
 			break;
 
 
+		case BuiltInHelperInvocation:
+			if (hlsl_options.shader_model < 50 || get_entry_point().model != ExecutionModelFragment)
+				SPIRV_CROSS_THROW("Helper Invocation input is only supported in PS 5.0 or higher.");
+			break;
+
 		case BuiltInClipDistance:
 		case BuiltInClipDistance:
 			// HLSL is a bit weird here, use SV_ClipDistance0, SV_ClipDistance1 and so on with vectors.
 			// HLSL is a bit weird here, use SV_ClipDistance0, SV_ClipDistance1 and so on with vectors.
 			for (uint32_t clip = 0; clip < clip_distance_count; clip += 4)
 			for (uint32_t clip = 0; clip < clip_distance_count; clip += 4)
@@ -984,6 +989,8 @@ std::string CompilerHLSL::builtin_to_glsl(spv::BuiltIn builtin, spv::StorageClas
 		return "WaveGetLaneIndex()";
 		return "WaveGetLaneIndex()";
 	case BuiltInSubgroupSize:
 	case BuiltInSubgroupSize:
 		return "WaveGetLaneCount()";
 		return "WaveGetLaneCount()";
+	case BuiltInHelperInvocation:
+		return "IsHelperLane()";
 
 
 	default:
 	default:
 		return CompilerGLSL::builtin_to_glsl(builtin, storage);
 		return CompilerGLSL::builtin_to_glsl(builtin, storage);
@@ -1103,6 +1110,11 @@ void CompilerHLSL::emit_builtin_variables()
 			type = "uint4";
 			type = "uint4";
 			break;
 			break;
 
 
+		case BuiltInHelperInvocation:
+			if (hlsl_options.shader_model < 50)
+				SPIRV_CROSS_THROW("Need SM 5.0 for Helper Invocation.");
+			break;
+
 		case BuiltInClipDistance:
 		case BuiltInClipDistance:
 			array_size = clip_distance_count;
 			array_size = clip_distance_count;
 			type = "float";
 			type = "float";
@@ -1335,7 +1347,8 @@ void CompilerHLSL::emit_resources()
 		}
 		}
 	});
 	});
 
 
-	if (execution.model == ExecutionModelVertex && hlsl_options.shader_model <= 30)
+	if (execution.model == ExecutionModelVertex && hlsl_options.shader_model <= 30 &&
+	    active_output_builtins.get(BuiltInPosition))
 	{
 	{
 		statement("uniform float4 gl_HalfPixel;");
 		statement("uniform float4 gl_HalfPixel;");
 		emitted = true;
 		emitted = true;
@@ -2102,7 +2115,11 @@ void CompilerHLSL::emit_buffer_block(const SPIRVariable &var)
 
 
 	bool is_uav = var.storage == StorageClassStorageBuffer || has_decoration(type.self, DecorationBufferBlock);
 	bool is_uav = var.storage == StorageClassStorageBuffer || has_decoration(type.self, DecorationBufferBlock);
 
 
-	if (is_uav)
+	if (flattened_buffer_blocks.count(var.self))
+	{
+		emit_buffer_block_flattened(var);
+	}
+	else if (is_uav)
 	{
 	{
 		Bitset flags = ir.get_buffer_block_flags(var);
 		Bitset flags = ir.get_buffer_block_flags(var);
 		bool is_readonly = flags.get(DecorationNonWritable) && !is_hlsl_force_storage_buffer_as_uav(var.self);
 		bool is_readonly = flags.get(DecorationNonWritable) && !is_hlsl_force_storage_buffer_as_uav(var.self);
@@ -2207,7 +2224,11 @@ void CompilerHLSL::emit_buffer_block(const SPIRVariable &var)
 
 
 void CompilerHLSL::emit_push_constant_block(const SPIRVariable &var)
 void CompilerHLSL::emit_push_constant_block(const SPIRVariable &var)
 {
 {
-	if (root_constants_layout.empty())
+	if (flattened_buffer_blocks.count(var.self))
+	{
+		emit_buffer_block_flattened(var);
+	}
+	else if (root_constants_layout.empty())
 	{
 	{
 		emit_buffer_block(var);
 		emit_buffer_block(var);
 	}
 	}
@@ -2365,7 +2386,7 @@ void CompilerHLSL::emit_function_prototype(SPIRFunction &func, const Bitset &ret
 		out_argument += " ";
 		out_argument += " ";
 		out_argument += "spvReturnValue";
 		out_argument += "spvReturnValue";
 		out_argument += type_to_array_glsl(type);
 		out_argument += type_to_array_glsl(type);
-		arglist.push_back(move(out_argument));
+		arglist.push_back(std::move(out_argument));
 	}
 	}
 
 
 	for (auto &arg : func.arguments)
 	for (auto &arg : func.arguments)
@@ -2512,6 +2533,7 @@ void CompilerHLSL::emit_hlsl_entry_point()
 		case BuiltInPointCoord:
 		case BuiltInPointCoord:
 		case BuiltInSubgroupSize:
 		case BuiltInSubgroupSize:
 		case BuiltInSubgroupLocalInvocationId:
 		case BuiltInSubgroupLocalInvocationId:
+		case BuiltInHelperInvocation:
 			break;
 			break;
 
 
 		case BuiltInSubgroupEqMask:
 		case BuiltInSubgroupEqMask:
@@ -2731,7 +2753,7 @@ void CompilerHLSL::emit_hlsl_entry_point()
 
 
 void CompilerHLSL::emit_fixup()
 void CompilerHLSL::emit_fixup()
 {
 {
-	if (is_vertex_like_shader())
+	if (is_vertex_like_shader() && active_output_builtins.get(BuiltInPosition))
 	{
 	{
 		// Do various mangling on the gl_Position.
 		// Do various mangling on the gl_Position.
 		if (hlsl_options.shader_model <= 30)
 		if (hlsl_options.shader_model <= 30)
@@ -3515,6 +3537,8 @@ void CompilerHLSL::emit_glsl_op(uint32_t result_type, uint32_t id, uint32_t eop,
 	auto int_type = to_signed_basetype(integer_width);
 	auto int_type = to_signed_basetype(integer_width);
 	auto uint_type = to_unsigned_basetype(integer_width);
 	auto uint_type = to_unsigned_basetype(integer_width);
 
 
+	op = get_remapped_glsl_op(op);
+
 	switch (op)
 	switch (op)
 	{
 	{
 	case GLSLstd450InverseSqrt:
 	case GLSLstd450InverseSqrt:
@@ -3987,7 +4011,7 @@ void CompilerHLSL::read_access_chain(string *expr, const string &lhs, const SPIR
 	if (lhs.empty())
 	if (lhs.empty())
 	{
 	{
 		assert(expr);
 		assert(expr);
-		*expr = move(load_expr);
+		*expr = std::move(load_expr);
 	}
 	}
 	else
 	else
 		statement(lhs, " = ", load_expr, ";");
 		statement(lhs, " = ", load_expr, ";");
@@ -4783,6 +4807,8 @@ void CompilerHLSL::emit_instruction(const Instruction &instruction)
 	auto int_type = to_signed_basetype(integer_width);
 	auto int_type = to_signed_basetype(integer_width);
 	auto uint_type = to_unsigned_basetype(integer_width);
 	auto uint_type = to_unsigned_basetype(integer_width);
 
 
+	opcode = get_remapped_spirv_op(opcode);
+
 	switch (opcode)
 	switch (opcode)
 	{
 	{
 	case OpAccessChain:
 	case OpAccessChain:
@@ -5333,7 +5359,7 @@ void CompilerHLSL::emit_instruction(const Instruction &instruction)
 				                        image_format_to_components(get<SPIRType>(var->basetype).image.format), imgexpr);
 				                        image_format_to_components(get<SPIRType>(var->basetype).image.format), imgexpr);
 		}
 		}
 
 
-		if (var && var->forwardable)
+		if (var)
 		{
 		{
 			bool forward = forced_temporaries.find(id) == end(forced_temporaries);
 			bool forward = forced_temporaries.find(id) == end(forced_temporaries);
 			auto &e = emit_op(result_type, id, imgexpr, forward);
 			auto &e = emit_op(result_type, id, imgexpr, forward);
@@ -5577,7 +5603,12 @@ void CompilerHLSL::emit_instruction(const Instruction &instruction)
 	}
 	}
 
 
 	case OpIsHelperInvocationEXT:
 	case OpIsHelperInvocationEXT:
-		SPIRV_CROSS_THROW("helperInvocationEXT() is not supported in HLSL.");
+		if (hlsl_options.shader_model < 50 || get_entry_point().model != ExecutionModelFragment)
+			SPIRV_CROSS_THROW("Helper Invocation input is only supported in PS 5.0 or higher.");
+		// Helper lane state with demote is volatile by nature.
+		// Do not forward this.
+		emit_op(ops[0], ops[1], "IsHelperLane()", false);
+		break;
 
 
 	case OpBeginInvocationInterlockEXT:
 	case OpBeginInvocationInterlockEXT:
 	case OpEndInvocationInterlockEXT:
 	case OpEndInvocationInterlockEXT:
@@ -5664,7 +5695,7 @@ void CompilerHLSL::require_texture_query_variant(uint32_t var_id)
 
 
 void CompilerHLSL::set_root_constant_layouts(std::vector<RootConstants> layout)
 void CompilerHLSL::set_root_constant_layouts(std::vector<RootConstants> layout)
 {
 {
-	root_constants_layout = move(layout);
+	root_constants_layout = std::move(layout);
 }
 }
 
 
 void CompilerHLSL::add_vertex_attribute_remap(const HLSLVertexAttributeRemap &vertex_attributes)
 void CompilerHLSL::add_vertex_attribute_remap(const HLSLVertexAttributeRemap &vertex_attributes)

+ 252 - 160
3rdparty/spirv-cross/spirv_msl.cpp

@@ -37,7 +37,7 @@ static const uint32_t k_unknown_component = ~0u;
 static const char *force_inline = "static inline __attribute__((always_inline))";
 static const char *force_inline = "static inline __attribute__((always_inline))";
 
 
 CompilerMSL::CompilerMSL(std::vector<uint32_t> spirv_)
 CompilerMSL::CompilerMSL(std::vector<uint32_t> spirv_)
-    : CompilerGLSL(move(spirv_))
+    : CompilerGLSL(std::move(spirv_))
 {
 {
 }
 }
 
 
@@ -1548,6 +1548,14 @@ void CompilerMSL::extract_global_variables_from_functions()
 	// Uniforms
 	// Uniforms
 	unordered_set<uint32_t> global_var_ids;
 	unordered_set<uint32_t> global_var_ids;
 	ir.for_each_typed_id<SPIRVariable>([&](uint32_t, SPIRVariable &var) {
 	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) &&
+		    BuiltIn(get_decoration(var.self, DecorationBuiltIn)) == BuiltInHelperInvocation)
+		{
+			return;
+		}
+
 		if (var.storage == StorageClassInput || var.storage == StorageClassOutput ||
 		if (var.storage == StorageClassInput || var.storage == StorageClassOutput ||
 		    var.storage == StorageClassUniform || var.storage == StorageClassUniformConstant ||
 		    var.storage == StorageClassUniform || var.storage == StorageClassUniformConstant ||
 		    var.storage == StorageClassPushConstant || var.storage == StorageClassStorageBuffer)
 		    var.storage == StorageClassPushConstant || var.storage == StorageClassStorageBuffer)
@@ -2508,12 +2516,15 @@ void CompilerMSL::add_composite_variable_to_interface_block(StorageClass storage
 	}
 	}
 }
 }
 
 
-void CompilerMSL::add_composite_member_variable_to_interface_block(StorageClass storage, const string &ib_var_ref,
-                                                                   SPIRType &ib_type, SPIRVariable &var,
-                                                                   uint32_t mbr_idx, InterfaceBlockMeta &meta)
+void CompilerMSL::add_composite_member_variable_to_interface_block(StorageClass storage,
+                                                                   const string &ib_var_ref, SPIRType &ib_type,
+                                                                   SPIRVariable &var, SPIRType &var_type,
+                                                                   uint32_t mbr_idx, InterfaceBlockMeta &meta,
+                                                                   const string &mbr_name_qual,
+                                                                   const string &var_chain_qual,
+                                                                   uint32_t &location, uint32_t &var_mbr_idx)
 {
 {
 	auto &entry_func = get<SPIRFunction>(ir.default_entry_point);
 	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);
 
 
 	BuiltIn builtin = BuiltInMax;
 	BuiltIn builtin = BuiltInMax;
 	bool is_builtin = is_member_builtin(var_type, mbr_idx, &builtin);
 	bool is_builtin = is_member_builtin(var_type, mbr_idx, &builtin);
@@ -2528,8 +2539,8 @@ void CompilerMSL::add_composite_member_variable_to_interface_block(StorageClass
 
 
 	uint32_t mbr_type_id = var_type.member_types[mbr_idx];
 	uint32_t mbr_type_id = var_type.member_types[mbr_idx];
 	auto &mbr_type = get<SPIRType>(mbr_type_id);
 	auto &mbr_type = get<SPIRType>(mbr_type_id);
-	uint32_t elem_cnt = 0;
 
 
+	uint32_t elem_cnt = 1;
 	if (is_matrix(mbr_type))
 	if (is_matrix(mbr_type))
 	{
 	{
 		if (is_array(mbr_type))
 		if (is_array(mbr_type))
@@ -2572,6 +2583,27 @@ void CompilerMSL::add_composite_member_variable_to_interface_block(StorageClass
 			return;
 			return;
 	}
 	}
 
 
+	// Recursively handle nested structures.
+	if (mbr_type.basetype == SPIRType::Struct)
+	{
+		for (uint32_t i = 0; i < elem_cnt; i++)
+		{
+			string mbr_name = append_member_name(mbr_name_qual, var_type, mbr_idx) + (elem_cnt == 1 ? "" : join("_", i));
+			string var_chain = join(var_chain_qual, ".", to_member_name(var_type, mbr_idx), (elem_cnt == 1 ? "" : join("[", i, "]")));
+			uint32_t sub_mbr_cnt = uint32_t(mbr_type.member_types.size());
+			for (uint32_t sub_mbr_idx = 0; sub_mbr_idx < sub_mbr_cnt; sub_mbr_idx++)
+			{
+				add_composite_member_variable_to_interface_block(storage, ib_var_ref, ib_type,
+																 var, mbr_type, sub_mbr_idx,
+																 meta, mbr_name, var_chain,
+																 location, var_mbr_idx);
+				// FIXME: Recursive structs and tessellation breaks here.
+				var_mbr_idx++;
+			}
+		}
+		return;
+	}
+
 	for (uint32_t i = 0; i < elem_cnt; i++)
 	for (uint32_t i = 0; i < elem_cnt; i++)
 	{
 	{
 		// Add a reference to the variable type to the interface struct.
 		// Add a reference to the variable type to the interface struct.
@@ -2582,26 +2614,40 @@ void CompilerMSL::add_composite_member_variable_to_interface_block(StorageClass
 			ib_type.member_types.push_back(usable_type->self);
 			ib_type.member_types.push_back(usable_type->self);
 
 
 		// Give the member a name
 		// Give the member a name
-		string mbr_name = ensure_valid_name(join(to_qualified_member_name(var_type, mbr_idx), "_", i), "m");
+		string mbr_name = ensure_valid_name(append_member_name(mbr_name_qual, var_type, mbr_idx) + (elem_cnt == 1 ? "" : join("_", i)), "m");
 		set_member_name(ib_type.self, ib_mbr_idx, mbr_name);
 		set_member_name(ib_type.self, ib_mbr_idx, mbr_name);
 
 
-		if (has_member_decoration(var_type.self, mbr_idx, DecorationLocation))
+		// Once we determine the location of the first member within nested structures,
+		// from a var of the topmost structure, the remaining flattened members of
+		// the nested structures will have consecutive location values. At this point,
+		// we've recursively tunnelled into structs, arrays, and matrices, and are
+		// down to a single location for each member now.
+		if (!is_builtin && location != UINT32_MAX)
 		{
 		{
-			uint32_t locn = get_member_decoration(var_type.self, mbr_idx, DecorationLocation) + i;
-			set_member_decoration(ib_type.self, ib_mbr_idx, DecorationLocation, locn);
-			mark_location_as_used_by_shader(locn, *usable_type, storage);
+			set_member_decoration(ib_type.self, ib_mbr_idx, DecorationLocation, location);
+			mark_location_as_used_by_shader(location, *usable_type, storage);
+			location++;
+		}
+		else if (has_member_decoration(var_type.self, mbr_idx, DecorationLocation))
+		{
+			location = get_member_decoration(var_type.self, mbr_idx, DecorationLocation) + i;
+			set_member_decoration(ib_type.self, ib_mbr_idx, DecorationLocation, location);
+			mark_location_as_used_by_shader(location, *usable_type, storage);
+			location++;
 		}
 		}
 		else if (has_decoration(var.self, DecorationLocation))
 		else if (has_decoration(var.self, DecorationLocation))
 		{
 		{
-			uint32_t locn = get_accumulated_member_location(var, mbr_idx, meta.strip_array) + i;
-			set_member_decoration(ib_type.self, ib_mbr_idx, DecorationLocation, locn);
-			mark_location_as_used_by_shader(locn, *usable_type, storage);
+			location = get_accumulated_member_location(var, mbr_idx, meta.strip_array) + i;
+			set_member_decoration(ib_type.self, ib_mbr_idx, DecorationLocation, location);
+			mark_location_as_used_by_shader(location, *usable_type, storage);
+			location++;
 		}
 		}
 		else if (is_builtin && is_tessellation_shader() && storage == StorageClassInput && inputs_by_builtin.count(builtin))
 		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);
+			location = inputs_by_builtin[builtin].location + i;
+			set_member_decoration(ib_type.self, ib_mbr_idx, DecorationLocation, location);
+			mark_location_as_used_by_shader(location, *usable_type, storage);
+			location++;
 		}
 		}
 		else if (is_builtin && (builtin == BuiltInClipDistance || builtin == BuiltInCullDistance))
 		else if (is_builtin && (builtin == BuiltInClipDistance || builtin == BuiltInCullDistance))
 		{
 		{
@@ -2611,7 +2657,7 @@ void CompilerMSL::add_composite_member_variable_to_interface_block(StorageClass
 		}
 		}
 
 
 		if (has_member_decoration(var_type.self, mbr_idx, DecorationComponent))
 		if (has_member_decoration(var_type.self, mbr_idx, DecorationComponent))
-			SPIRV_CROSS_THROW("DecorationComponent on matrices and arrays make little sense.");
+			SPIRV_CROSS_THROW("DecorationComponent on matrices and arrays is not supported.");
 
 
 		if (storage != StorageClassInput || !pull_model_inputs.count(var.self))
 		if (storage != StorageClassInput || !pull_model_inputs.count(var.self))
 		{
 		{
@@ -2627,47 +2673,36 @@ void CompilerMSL::add_composite_member_variable_to_interface_block(StorageClass
 		}
 		}
 
 
 		set_extended_member_decoration(ib_type.self, ib_mbr_idx, SPIRVCrossDecorationInterfaceOrigID, var.self);
 		set_extended_member_decoration(ib_type.self, ib_mbr_idx, SPIRVCrossDecorationInterfaceOrigID, var.self);
-		set_extended_member_decoration(ib_type.self, ib_mbr_idx, SPIRVCrossDecorationInterfaceMemberIndex, mbr_idx);
+		set_extended_member_decoration(ib_type.self, ib_mbr_idx, SPIRVCrossDecorationInterfaceMemberIndex, var_mbr_idx);
 
 
 		// Unflatten or flatten from [[stage_in]] or [[stage_out]] as appropriate.
 		// Unflatten or flatten from [[stage_in]] or [[stage_out]] as appropriate.
 		if (!meta.strip_array && meta.allow_local_declaration)
 		if (!meta.strip_array && meta.allow_local_declaration)
 		{
 		{
+			string var_chain = join(var_chain_qual, ".", to_member_name(var_type, mbr_idx), (elem_cnt == 1 ? "" : join("[", i, "]")));
 			switch (storage)
 			switch (storage)
 			{
 			{
 			case StorageClassInput:
 			case StorageClassInput:
-				entry_func.fixup_hooks_in.push_back([=, &var, &var_type]() {
+				entry_func.fixup_hooks_in.push_back([=, &var]() {
+					string lerp_call;
 					if (pull_model_inputs.count(var.self))
 					if (pull_model_inputs.count(var.self))
 					{
 					{
-						string lerp_call;
 						if (is_centroid)
 						if (is_centroid)
 							lerp_call = ".interpolate_at_centroid()";
 							lerp_call = ".interpolate_at_centroid()";
 						else if (is_sample)
 						else if (is_sample)
 							lerp_call = join(".interpolate_at_sample(", to_expression(builtin_sample_id_id), ")");
 							lerp_call = join(".interpolate_at_sample(", to_expression(builtin_sample_id_id), ")");
 						else
 						else
 							lerp_call = ".interpolate_at_center()";
 							lerp_call = ".interpolate_at_center()";
-						statement(to_name(var.self), ".", to_member_name(var_type, mbr_idx), "[", i, "] = ", ib_var_ref,
-						          ".", mbr_name, lerp_call, ";");
-					}
-					else
-					{
-						statement(to_name(var.self), ".", to_member_name(var_type, mbr_idx), "[", i, "] = ", ib_var_ref,
-						          ".", mbr_name, ";");
 					}
 					}
+					statement(var_chain, " = ", ib_var_ref, ".", mbr_name, lerp_call, ";");
 				});
 				});
 				break;
 				break;
 
 
 			case StorageClassOutput:
 			case StorageClassOutput:
-				entry_func.fixup_hooks_out.push_back([=, &var, &var_type]() {
+				entry_func.fixup_hooks_out.push_back([=]() {
 					if (flatten_from_ib_var)
 					if (flatten_from_ib_var)
-					{
-						statement(ib_var_ref, ".", mbr_name, " = ", ib_var_ref, ".", flatten_from_ib_mbr_name, "[", i,
-						          "];");
-					}
+						statement(ib_var_ref, ".", mbr_name, " = ", ib_var_ref, ".", flatten_from_ib_mbr_name, "[", i, "];");
 					else
 					else
-					{
-						statement(ib_var_ref, ".", mbr_name, " = ", to_name(var.self), ".",
-						          to_member_name(var_type, mbr_idx), "[", i, "];");
-					}
+						statement(ib_var_ref, ".", mbr_name, " = ", var_chain, ";");
 				});
 				});
 				break;
 				break;
 
 
@@ -2678,11 +2713,14 @@ void CompilerMSL::add_composite_member_variable_to_interface_block(StorageClass
 	}
 	}
 }
 }
 
 
-void CompilerMSL::add_plain_member_variable_to_interface_block(StorageClass storage, const string &ib_var_ref,
-                                                               SPIRType &ib_type, SPIRVariable &var, uint32_t mbr_idx,
-                                                               InterfaceBlockMeta &meta)
+void CompilerMSL::add_plain_member_variable_to_interface_block(StorageClass storage,
+                                                               const string &ib_var_ref, SPIRType &ib_type,
+                                                               SPIRVariable &var, SPIRType &var_type,
+                                                               uint32_t mbr_idx, InterfaceBlockMeta &meta,
+                                                               const string &mbr_name_qual,
+                                                               const string &var_chain_qual,
+                                                               uint32_t &location, uint32_t &var_mbr_idx)
 {
 {
-	auto &var_type = meta.strip_array ? get_variable_element_type(var) : get_variable_data_type(var);
 	auto &entry_func = get<SPIRFunction>(ir.default_entry_point);
 	auto &entry_func = get<SPIRFunction>(ir.default_entry_point);
 
 
 	BuiltIn builtin = BuiltInMax;
 	BuiltIn builtin = BuiltInMax;
@@ -2707,7 +2745,7 @@ void CompilerMSL::add_plain_member_variable_to_interface_block(StorageClass stor
 		ib_type.member_types.push_back(mbr_type_id);
 		ib_type.member_types.push_back(mbr_type_id);
 
 
 	// Give the member a name
 	// Give the member a name
-	string mbr_name = ensure_valid_name(to_qualified_member_name(var_type, mbr_idx), "m");
+	string mbr_name = ensure_valid_name(append_member_name(mbr_name_qual, var_type, mbr_idx), "m");
 	set_member_name(ib_type.self, ib_mbr_idx, mbr_name);
 	set_member_name(ib_type.self, ib_mbr_idx, mbr_name);
 
 
 	// Update the original variable reference to include the structure reference
 	// Update the original variable reference to include the structure reference
@@ -2724,7 +2762,7 @@ void CompilerMSL::add_plain_member_variable_to_interface_block(StorageClass stor
 	}
 	}
 
 
 	bool flatten_stage_out = false;
 	bool flatten_stage_out = false;
-
+	string var_chain = var_chain_qual + "." + to_member_name(var_type, mbr_idx);
 	if (is_builtin && !meta.strip_array)
 	if (is_builtin && !meta.strip_array)
 	{
 	{
 		// For the builtin gl_PerVertex, we cannot treat it as a block anyways,
 		// For the builtin gl_PerVertex, we cannot treat it as a block anyways,
@@ -2737,15 +2775,15 @@ void CompilerMSL::add_plain_member_variable_to_interface_block(StorageClass stor
 		switch (storage)
 		switch (storage)
 		{
 		{
 		case StorageClassInput:
 		case StorageClassInput:
-			entry_func.fixup_hooks_in.push_back([=, &var, &var_type]() {
-				statement(to_name(var.self), ".", to_member_name(var_type, mbr_idx), " = ", qual_var_name, ";");
+			entry_func.fixup_hooks_in.push_back([=]() {
+				statement(var_chain, " = ", qual_var_name, ";");
 			});
 			});
 			break;
 			break;
 
 
 		case StorageClassOutput:
 		case StorageClassOutput:
 			flatten_stage_out = true;
 			flatten_stage_out = true;
-			entry_func.fixup_hooks_out.push_back([=, &var, &var_type]() {
-				statement(qual_var_name, " = ", to_name(var.self), ".", to_member_name(var_type, mbr_idx), ";");
+			entry_func.fixup_hooks_out.push_back([=]() {
+				statement(qual_var_name, " = ", var_chain, ";");
 			});
 			});
 			break;
 			break;
 
 
@@ -2754,48 +2792,56 @@ void CompilerMSL::add_plain_member_variable_to_interface_block(StorageClass stor
 		}
 		}
 	}
 	}
 
 
-	// Copy the variable location from the original variable to the member
-	if (has_member_decoration(var_type.self, mbr_idx, DecorationLocation))
+	// Once we determine the location of the first member within nested structures,
+	// from a var of the topmost structure, the remaining flattened members of
+	// the nested structures will have consecutive location values. At this point,
+	// we've recursively tunnelled into structs, arrays, and matrices, and are
+	// down to a single location for each member now.
+	if (!is_builtin && location != UINT32_MAX)
 	{
 	{
-		uint32_t locn = get_member_decoration(var_type.self, mbr_idx, DecorationLocation);
+		set_member_decoration(ib_type.self, ib_mbr_idx, DecorationLocation, location);
+		mark_location_as_used_by_shader(location, get<SPIRType>(mbr_type_id), storage);
+		location++;
+	}
+	else if (has_member_decoration(var_type.self, mbr_idx, DecorationLocation))
+	{
+		location = get_member_decoration(var_type.self, mbr_idx, DecorationLocation);
 		uint32_t comp = get_member_decoration(var_type.self, mbr_idx, DecorationComponent);
 		uint32_t comp = get_member_decoration(var_type.self, mbr_idx, DecorationComponent);
 		if (storage == StorageClassInput)
 		if (storage == StorageClassInput)
 		{
 		{
-			mbr_type_id = ensure_correct_input_type(mbr_type_id, locn, comp, 0, meta.strip_array);
+			mbr_type_id = ensure_correct_input_type(mbr_type_id, location, comp, 0, meta.strip_array);
 			var_type.member_types[mbr_idx] = mbr_type_id;
 			var_type.member_types[mbr_idx] = mbr_type_id;
 			if (storage == StorageClassInput && pull_model_inputs.count(var.self))
 			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);
 				ib_type.member_types[ib_mbr_idx] = build_msl_interpolant_type(mbr_type_id, is_noperspective);
 			else
 			else
 				ib_type.member_types[ib_mbr_idx] = mbr_type_id;
 				ib_type.member_types[ib_mbr_idx] = mbr_type_id;
 		}
 		}
-		set_member_decoration(ib_type.self, ib_mbr_idx, DecorationLocation, locn);
-		mark_location_as_used_by_shader(locn, get<SPIRType>(mbr_type_id), storage);
+		set_member_decoration(ib_type.self, ib_mbr_idx, DecorationLocation, location);
+		mark_location_as_used_by_shader(location, get<SPIRType>(mbr_type_id), storage);
+		location++;
 	}
 	}
 	else if (has_decoration(var.self, DecorationLocation))
 	else if (has_decoration(var.self, DecorationLocation))
 	{
 	{
-		// The block itself might have a location and in this case, all members of the block
-		// receive incrementing locations.
-		uint32_t locn = get_accumulated_member_location(var, mbr_idx, meta.strip_array);
+		location = get_accumulated_member_location(var, mbr_idx, meta.strip_array);
 		if (storage == StorageClassInput)
 		if (storage == StorageClassInput)
 		{
 		{
-			mbr_type_id = ensure_correct_input_type(mbr_type_id, locn, 0, 0, meta.strip_array);
+			mbr_type_id = ensure_correct_input_type(mbr_type_id, location, 0, 0, meta.strip_array);
 			var_type.member_types[mbr_idx] = mbr_type_id;
 			var_type.member_types[mbr_idx] = mbr_type_id;
 			if (storage == StorageClassInput && pull_model_inputs.count(var.self))
 			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);
 				ib_type.member_types[ib_mbr_idx] = build_msl_interpolant_type(mbr_type_id, is_noperspective);
 			else
 			else
 				ib_type.member_types[ib_mbr_idx] = mbr_type_id;
 				ib_type.member_types[ib_mbr_idx] = mbr_type_id;
 		}
 		}
-		set_member_decoration(ib_type.self, ib_mbr_idx, DecorationLocation, locn);
-		mark_location_as_used_by_shader(locn, get<SPIRType>(mbr_type_id), storage);
+		set_member_decoration(ib_type.self, ib_mbr_idx, DecorationLocation, location);
+		mark_location_as_used_by_shader(location, get<SPIRType>(mbr_type_id), storage);
+		location++;
 	}
 	}
 	else if (is_builtin && is_tessellation_shader() && storage == StorageClassInput && inputs_by_builtin.count(builtin))
 	else if (is_builtin && is_tessellation_shader() && storage == StorageClassInput && inputs_by_builtin.count(builtin))
 	{
 	{
-		uint32_t locn = 0;
-		auto builtin_itr = inputs_by_builtin.find(builtin);
-		if (builtin_itr != end(inputs_by_builtin))
-			locn = builtin_itr->second.location;
-		set_member_decoration(ib_type.self, ib_mbr_idx, DecorationLocation, locn);
-		mark_location_as_used_by_shader(locn, get<SPIRType>(mbr_type_id), storage);
+		location = inputs_by_builtin[builtin].location;
+		set_member_decoration(ib_type.self, ib_mbr_idx, DecorationLocation, location);
+		mark_location_as_used_by_shader(location, get<SPIRType>(mbr_type_id), storage);
+		location++;
 	}
 	}
 
 
 	// Copy the component location, if present.
 	// Copy the component location, if present.
@@ -2854,7 +2900,7 @@ void CompilerMSL::add_plain_member_variable_to_interface_block(StorageClass stor
 	}
 	}
 
 
 	set_extended_member_decoration(ib_type.self, ib_mbr_idx, SPIRVCrossDecorationInterfaceOrigID, var.self);
 	set_extended_member_decoration(ib_type.self, ib_mbr_idx, SPIRVCrossDecorationInterfaceOrigID, var.self);
-	set_extended_member_decoration(ib_type.self, ib_mbr_idx, SPIRVCrossDecorationInterfaceMemberIndex, mbr_idx);
+	set_extended_member_decoration(ib_type.self, ib_mbr_idx, SPIRVCrossDecorationInterfaceMemberIndex, var_mbr_idx);
 }
 }
 
 
 // In Metal, the tessellation levels are stored as tightly packed half-precision floating point values.
 // In Metal, the tessellation levels are stored as tightly packed half-precision floating point values.
@@ -3114,67 +3160,102 @@ void CompilerMSL::add_variable_to_interface_block(StorageClass storage, const st
 		else
 		else
 		{
 		{
 			bool masked_block = false;
 			bool masked_block = false;
+			uint32_t location = UINT32_MAX;
+			uint32_t var_mbr_idx = 0;
+			uint32_t elem_cnt = 1;
+			if (is_matrix(var_type))
+			{
+				if (is_array(var_type))
+					SPIRV_CROSS_THROW("MSL cannot emit arrays-of-matrices in input and output variables.");
 
 
-			// Flatten the struct members into the interface struct
-			for (uint32_t mbr_idx = 0; mbr_idx < uint32_t(var_type.member_types.size()); mbr_idx++)
+				elem_cnt = var_type.columns;
+			}
+			else if (is_array(var_type))
 			{
 			{
-				builtin = BuiltInMax;
-				is_builtin = is_member_builtin(var_type, mbr_idx, &builtin);
-				auto &mbr_type = get<SPIRType>(var_type.member_types[mbr_idx]);
+				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);
+			}
 
 
-				if (storage == StorageClassOutput && is_stage_output_block_member_masked(var, mbr_idx, meta.strip_array))
+			for (uint32_t elem_idx = 0; elem_idx < elem_cnt; elem_idx++)
+			{
+				// Flatten the struct members into the interface struct
+				for (uint32_t mbr_idx = 0; mbr_idx < uint32_t(var_type.member_types.size()); mbr_idx++)
 				{
 				{
-					if (is_block)
-						masked_block = true;
+					builtin = BuiltInMax;
+					is_builtin = is_member_builtin(var_type, mbr_idx, &builtin);
+					auto &mbr_type = get<SPIRType>(var_type.member_types[mbr_idx]);
 
 
-					// Non-builtin block output variables are just ignored, since they will still access
-					// the block variable as-is. They're just not flattened.
-					if (is_builtin && !meta.strip_array)
+					if (storage == StorageClassOutput && is_stage_output_block_member_masked(var, mbr_idx, meta.strip_array))
 					{
 					{
-						// Emit a fake variable instead.
-						uint32_t ids = ir.increase_bound_by(2);
-						uint32_t ptr_type_id = ids + 0;
-						uint32_t var_id = ids + 1;
-
-						auto ptr_type = mbr_type;
-						ptr_type.pointer = true;
-						ptr_type.pointer_depth++;
-						ptr_type.parent_type = var_type.member_types[mbr_idx];
-						ptr_type.storage = StorageClassOutput;
-
-						uint32_t initializer = 0;
-						if (var.initializer)
-							if (auto *c = maybe_get<SPIRConstant>(var.initializer))
-								initializer = c->subconstants[mbr_idx];
-
-						set<SPIRType>(ptr_type_id, ptr_type);
-						set<SPIRVariable>(var_id, ptr_type_id, StorageClassOutput, initializer);
-						entry_func.add_local_variable(var_id);
-						vars_needing_early_declaration.push_back(var_id);
-						set_name(var_id, builtin_to_glsl(builtin, StorageClassOutput));
-						set_decoration(var_id, DecorationBuiltIn, builtin);
-					}
-				}
-				else if (!is_builtin || has_active_builtin(builtin, storage))
-				{
-					bool is_composite_type = is_matrix(mbr_type) || is_array(mbr_type);
-					bool attribute_load_store =
-					    storage == StorageClassInput && get_execution_model() != ExecutionModelFragment;
-					bool storage_is_stage_io = variable_storage_requires_stage_io(storage);
+						location = UINT32_MAX;		// Skip this member and resolve location again on next var member
 
 
-					// Clip/CullDistance always need to be declared as user attributes.
-					if (builtin == BuiltInClipDistance || builtin == BuiltInCullDistance)
-						is_builtin = false;
+						if (is_block)
+							masked_block = true;
 
 
-					if ((!is_builtin || attribute_load_store) && storage_is_stage_io && is_composite_type)
-					{
-						add_composite_member_variable_to_interface_block(storage, ib_var_ref, ib_type, var, mbr_idx,
-						                                                 meta);
+						// Non-builtin block output variables are just ignored, since they will still access
+						// the block variable as-is. They're just not flattened.
+						if (is_builtin && !meta.strip_array)
+						{
+							// Emit a fake variable instead.
+							uint32_t ids = ir.increase_bound_by(2);
+							uint32_t ptr_type_id = ids + 0;
+							uint32_t var_id = ids + 1;
+
+							auto ptr_type = mbr_type;
+							ptr_type.pointer = true;
+							ptr_type.pointer_depth++;
+							ptr_type.parent_type = var_type.member_types[mbr_idx];
+							ptr_type.storage = StorageClassOutput;
+
+							uint32_t initializer = 0;
+							if (var.initializer)
+								if (auto *c = maybe_get<SPIRConstant>(var.initializer))
+									initializer = c->subconstants[mbr_idx];
+
+							set<SPIRType>(ptr_type_id, ptr_type);
+							set<SPIRVariable>(var_id, ptr_type_id, StorageClassOutput, initializer);
+							entry_func.add_local_variable(var_id);
+							vars_needing_early_declaration.push_back(var_id);
+							set_name(var_id, builtin_to_glsl(builtin, StorageClassOutput));
+							set_decoration(var_id, DecorationBuiltIn, builtin);
+						}
 					}
 					}
-					else
+					else if (!is_builtin || has_active_builtin(builtin, storage))
 					{
 					{
-						add_plain_member_variable_to_interface_block(storage, ib_var_ref, ib_type, var, mbr_idx, meta);
+						bool is_composite_type = is_matrix(mbr_type) || is_array(mbr_type) || mbr_type.basetype == SPIRType::Struct;
+						bool attribute_load_store =
+						storage == StorageClassInput && get_execution_model() != ExecutionModelFragment;
+						bool storage_is_stage_io = variable_storage_requires_stage_io(storage);
+
+						// Clip/CullDistance always need to be declared as user attributes.
+						if (builtin == BuiltInClipDistance || builtin == BuiltInCullDistance)
+							is_builtin = false;
+
+						string mbr_name_qual = to_name(var_type.self);
+						string var_chain_qual = to_name(var.self);
+						if (elem_cnt > 1) {
+							mbr_name_qual += join("_", elem_idx);
+							var_chain_qual += join("[", elem_idx, "]");
+						}
+
+						if ((!is_builtin || attribute_load_store) && storage_is_stage_io && is_composite_type)
+						{
+							add_composite_member_variable_to_interface_block(storage, ib_var_ref, ib_type,
+							                                                 var, var_type, mbr_idx, meta,
+							                                                 mbr_name_qual, var_chain_qual,
+							                                                 location, var_mbr_idx);
+						}
+						else
+						{
+							add_plain_member_variable_to_interface_block(storage, ib_var_ref, ib_type,
+							                                             var, var_type, mbr_idx, meta,
+							                                             mbr_name_qual, var_chain_qual,
+							                                             location, var_mbr_idx);
+						}
 					}
 					}
+					var_mbr_idx++;
 				}
 				}
 			}
 			}
 
 
@@ -5370,7 +5451,7 @@ void CompilerMSL::emit_custom_functions()
 			statement("template<typename T>");
 			statement("template<typename T>");
 			statement("inline T spvSubgroupBroadcast(T value, ushort lane)");
 			statement("inline T spvSubgroupBroadcast(T value, ushort lane)");
 			begin_scope();
 			begin_scope();
-			if (msl_options.is_ios() && !msl_options.ios_use_simdgroup_functions)
+			if (msl_options.use_quadgroup_operation())
 				statement("return quad_broadcast(value, lane);");
 				statement("return quad_broadcast(value, lane);");
 			else
 			else
 				statement("return simd_broadcast(value, lane);");
 				statement("return simd_broadcast(value, lane);");
@@ -5379,7 +5460,7 @@ void CompilerMSL::emit_custom_functions()
 			statement("template<>");
 			statement("template<>");
 			statement("inline bool spvSubgroupBroadcast(bool value, ushort lane)");
 			statement("inline bool spvSubgroupBroadcast(bool value, ushort lane)");
 			begin_scope();
 			begin_scope();
-			if (msl_options.is_ios() && !msl_options.ios_use_simdgroup_functions)
+			if (msl_options.use_quadgroup_operation())
 				statement("return !!quad_broadcast((ushort)value, lane);");
 				statement("return !!quad_broadcast((ushort)value, lane);");
 			else
 			else
 				statement("return !!simd_broadcast((ushort)value, lane);");
 				statement("return !!simd_broadcast((ushort)value, lane);");
@@ -5388,7 +5469,7 @@ void CompilerMSL::emit_custom_functions()
 			statement("template<uint N>");
 			statement("template<uint N>");
 			statement("inline vec<bool, N> spvSubgroupBroadcast(vec<bool, N> value, ushort lane)");
 			statement("inline vec<bool, N> spvSubgroupBroadcast(vec<bool, N> value, ushort lane)");
 			begin_scope();
 			begin_scope();
-			if (msl_options.is_ios() && !msl_options.ios_use_simdgroup_functions)
+			if (msl_options.use_quadgroup_operation())
 				statement("return (vec<bool, N>)quad_broadcast((vec<ushort, N>)value, lane);");
 				statement("return (vec<bool, N>)quad_broadcast((vec<ushort, N>)value, lane);");
 			else
 			else
 				statement("return (vec<bool, N>)simd_broadcast((vec<ushort, N>)value, lane);");
 				statement("return (vec<bool, N>)simd_broadcast((vec<ushort, N>)value, lane);");
@@ -5400,7 +5481,7 @@ void CompilerMSL::emit_custom_functions()
 			statement("template<typename T>");
 			statement("template<typename T>");
 			statement("inline T spvSubgroupBroadcastFirst(T value)");
 			statement("inline T spvSubgroupBroadcastFirst(T value)");
 			begin_scope();
 			begin_scope();
-			if (msl_options.is_ios() && !msl_options.ios_use_simdgroup_functions)
+			if (msl_options.use_quadgroup_operation())
 				statement("return quad_broadcast_first(value);");
 				statement("return quad_broadcast_first(value);");
 			else
 			else
 				statement("return simd_broadcast_first(value);");
 				statement("return simd_broadcast_first(value);");
@@ -5409,7 +5490,7 @@ void CompilerMSL::emit_custom_functions()
 			statement("template<>");
 			statement("template<>");
 			statement("inline bool spvSubgroupBroadcastFirst(bool value)");
 			statement("inline bool spvSubgroupBroadcastFirst(bool value)");
 			begin_scope();
 			begin_scope();
-			if (msl_options.is_ios() && !msl_options.ios_use_simdgroup_functions)
+			if (msl_options.use_quadgroup_operation())
 				statement("return !!quad_broadcast_first((ushort)value);");
 				statement("return !!quad_broadcast_first((ushort)value);");
 			else
 			else
 				statement("return !!simd_broadcast_first((ushort)value);");
 				statement("return !!simd_broadcast_first((ushort)value);");
@@ -5418,7 +5499,7 @@ void CompilerMSL::emit_custom_functions()
 			statement("template<uint N>");
 			statement("template<uint N>");
 			statement("inline vec<bool, N> spvSubgroupBroadcastFirst(vec<bool, N> value)");
 			statement("inline vec<bool, N> spvSubgroupBroadcastFirst(vec<bool, N> value)");
 			begin_scope();
 			begin_scope();
-			if (msl_options.is_ios() && !msl_options.ios_use_simdgroup_functions)
+			if (msl_options.use_quadgroup_operation())
 				statement("return (vec<bool, N>)quad_broadcast_first((vec<ushort, N>)value);");
 				statement("return (vec<bool, N>)quad_broadcast_first((vec<ushort, N>)value);");
 			else
 			else
 				statement("return (vec<bool, N>)simd_broadcast_first((vec<ushort, N>)value);");
 				statement("return (vec<bool, N>)simd_broadcast_first((vec<ushort, N>)value);");
@@ -5429,7 +5510,7 @@ void CompilerMSL::emit_custom_functions()
 		case SPVFuncImplSubgroupBallot:
 		case SPVFuncImplSubgroupBallot:
 			statement("inline uint4 spvSubgroupBallot(bool value)");
 			statement("inline uint4 spvSubgroupBallot(bool value)");
 			begin_scope();
 			begin_scope();
-			if (msl_options.is_ios() && !msl_options.ios_use_simdgroup_functions)
+			if (msl_options.use_quadgroup_operation())
 			{
 			{
 				statement("return uint4((quad_vote::vote_t)quad_ballot(value), 0, 0, 0);");
 				statement("return uint4((quad_vote::vote_t)quad_ballot(value), 0, 0, 0);");
 			}
 			}
@@ -5557,7 +5638,7 @@ void CompilerMSL::emit_custom_functions()
 			statement("template<typename T>");
 			statement("template<typename T>");
 			statement("inline bool spvSubgroupAllEqual(T value)");
 			statement("inline bool spvSubgroupAllEqual(T value)");
 			begin_scope();
 			begin_scope();
-			if (msl_options.is_ios() && !msl_options.ios_use_simdgroup_functions)
+			if (msl_options.use_quadgroup_operation())
 				statement("return quad_all(all(value == quad_broadcast_first(value)));");
 				statement("return quad_all(all(value == quad_broadcast_first(value)));");
 			else
 			else
 				statement("return simd_all(all(value == simd_broadcast_first(value)));");
 				statement("return simd_all(all(value == simd_broadcast_first(value)));");
@@ -5566,7 +5647,7 @@ void CompilerMSL::emit_custom_functions()
 			statement("template<>");
 			statement("template<>");
 			statement("inline bool spvSubgroupAllEqual(bool value)");
 			statement("inline bool spvSubgroupAllEqual(bool value)");
 			begin_scope();
 			begin_scope();
-			if (msl_options.is_ios() && !msl_options.ios_use_simdgroup_functions)
+			if (msl_options.use_quadgroup_operation())
 				statement("return quad_all(value) || !quad_any(value);");
 				statement("return quad_all(value) || !quad_any(value);");
 			else
 			else
 				statement("return simd_all(value) || !simd_any(value);");
 				statement("return simd_all(value) || !simd_any(value);");
@@ -5575,7 +5656,7 @@ void CompilerMSL::emit_custom_functions()
 			statement("template<uint N>");
 			statement("template<uint N>");
 			statement("inline bool spvSubgroupAllEqual(vec<bool, N> value)");
 			statement("inline bool spvSubgroupAllEqual(vec<bool, N> value)");
 			begin_scope();
 			begin_scope();
-			if (msl_options.is_ios() && !msl_options.ios_use_simdgroup_functions)
+			if (msl_options.use_quadgroup_operation())
 				statement("return quad_all(all(value == (vec<bool, N>)quad_broadcast_first((vec<ushort, N>)value)));");
 				statement("return quad_all(all(value == (vec<bool, N>)quad_broadcast_first((vec<ushort, N>)value)));");
 			else
 			else
 				statement("return simd_all(all(value == (vec<bool, N>)simd_broadcast_first((vec<ushort, N>)value)));");
 				statement("return simd_all(all(value == (vec<bool, N>)simd_broadcast_first((vec<ushort, N>)value)));");
@@ -5587,7 +5668,7 @@ void CompilerMSL::emit_custom_functions()
 			statement("template<typename T>");
 			statement("template<typename T>");
 			statement("inline T spvSubgroupShuffle(T value, ushort lane)");
 			statement("inline T spvSubgroupShuffle(T value, ushort lane)");
 			begin_scope();
 			begin_scope();
-			if (msl_options.is_ios() && !msl_options.ios_use_simdgroup_functions)
+			if (msl_options.use_quadgroup_operation())
 				statement("return quad_shuffle(value, lane);");
 				statement("return quad_shuffle(value, lane);");
 			else
 			else
 				statement("return simd_shuffle(value, lane);");
 				statement("return simd_shuffle(value, lane);");
@@ -5596,7 +5677,7 @@ void CompilerMSL::emit_custom_functions()
 			statement("template<>");
 			statement("template<>");
 			statement("inline bool spvSubgroupShuffle(bool value, ushort lane)");
 			statement("inline bool spvSubgroupShuffle(bool value, ushort lane)");
 			begin_scope();
 			begin_scope();
-			if (msl_options.is_ios() && !msl_options.ios_use_simdgroup_functions)
+			if (msl_options.use_quadgroup_operation())
 				statement("return !!quad_shuffle((ushort)value, lane);");
 				statement("return !!quad_shuffle((ushort)value, lane);");
 			else
 			else
 				statement("return !!simd_shuffle((ushort)value, lane);");
 				statement("return !!simd_shuffle((ushort)value, lane);");
@@ -5605,7 +5686,7 @@ void CompilerMSL::emit_custom_functions()
 			statement("template<uint N>");
 			statement("template<uint N>");
 			statement("inline vec<bool, N> spvSubgroupShuffle(vec<bool, N> value, ushort lane)");
 			statement("inline vec<bool, N> spvSubgroupShuffle(vec<bool, N> value, ushort lane)");
 			begin_scope();
 			begin_scope();
-			if (msl_options.is_ios() && !msl_options.ios_use_simdgroup_functions)
+			if (msl_options.use_quadgroup_operation())
 				statement("return (vec<bool, N>)quad_shuffle((vec<ushort, N>)value, lane);");
 				statement("return (vec<bool, N>)quad_shuffle((vec<ushort, N>)value, lane);");
 			else
 			else
 				statement("return (vec<bool, N>)simd_shuffle((vec<ushort, N>)value, lane);");
 				statement("return (vec<bool, N>)simd_shuffle((vec<ushort, N>)value, lane);");
@@ -5617,7 +5698,7 @@ void CompilerMSL::emit_custom_functions()
 			statement("template<typename T>");
 			statement("template<typename T>");
 			statement("inline T spvSubgroupShuffleXor(T value, ushort mask)");
 			statement("inline T spvSubgroupShuffleXor(T value, ushort mask)");
 			begin_scope();
 			begin_scope();
-			if (msl_options.is_ios() && !msl_options.ios_use_simdgroup_functions)
+			if (msl_options.use_quadgroup_operation())
 				statement("return quad_shuffle_xor(value, mask);");
 				statement("return quad_shuffle_xor(value, mask);");
 			else
 			else
 				statement("return simd_shuffle_xor(value, mask);");
 				statement("return simd_shuffle_xor(value, mask);");
@@ -5626,7 +5707,7 @@ void CompilerMSL::emit_custom_functions()
 			statement("template<>");
 			statement("template<>");
 			statement("inline bool spvSubgroupShuffleXor(bool value, ushort mask)");
 			statement("inline bool spvSubgroupShuffleXor(bool value, ushort mask)");
 			begin_scope();
 			begin_scope();
-			if (msl_options.is_ios() && !msl_options.ios_use_simdgroup_functions)
+			if (msl_options.use_quadgroup_operation())
 				statement("return !!quad_shuffle_xor((ushort)value, mask);");
 				statement("return !!quad_shuffle_xor((ushort)value, mask);");
 			else
 			else
 				statement("return !!simd_shuffle_xor((ushort)value, mask);");
 				statement("return !!simd_shuffle_xor((ushort)value, mask);");
@@ -5635,7 +5716,7 @@ void CompilerMSL::emit_custom_functions()
 			statement("template<uint N>");
 			statement("template<uint N>");
 			statement("inline vec<bool, N> spvSubgroupShuffleXor(vec<bool, N> value, ushort mask)");
 			statement("inline vec<bool, N> spvSubgroupShuffleXor(vec<bool, N> value, ushort mask)");
 			begin_scope();
 			begin_scope();
-			if (msl_options.is_ios() && !msl_options.ios_use_simdgroup_functions)
+			if (msl_options.use_quadgroup_operation())
 				statement("return (vec<bool, N>)quad_shuffle_xor((vec<ushort, N>)value, mask);");
 				statement("return (vec<bool, N>)quad_shuffle_xor((vec<ushort, N>)value, mask);");
 			else
 			else
 				statement("return (vec<bool, N>)simd_shuffle_xor((vec<ushort, N>)value, mask);");
 				statement("return (vec<bool, N>)simd_shuffle_xor((vec<ushort, N>)value, mask);");
@@ -5647,7 +5728,7 @@ void CompilerMSL::emit_custom_functions()
 			statement("template<typename T>");
 			statement("template<typename T>");
 			statement("inline T spvSubgroupShuffleUp(T value, ushort delta)");
 			statement("inline T spvSubgroupShuffleUp(T value, ushort delta)");
 			begin_scope();
 			begin_scope();
-			if (msl_options.is_ios() && !msl_options.ios_use_simdgroup_functions)
+			if (msl_options.use_quadgroup_operation())
 				statement("return quad_shuffle_up(value, delta);");
 				statement("return quad_shuffle_up(value, delta);");
 			else
 			else
 				statement("return simd_shuffle_up(value, delta);");
 				statement("return simd_shuffle_up(value, delta);");
@@ -5656,7 +5737,7 @@ void CompilerMSL::emit_custom_functions()
 			statement("template<>");
 			statement("template<>");
 			statement("inline bool spvSubgroupShuffleUp(bool value, ushort delta)");
 			statement("inline bool spvSubgroupShuffleUp(bool value, ushort delta)");
 			begin_scope();
 			begin_scope();
-			if (msl_options.is_ios() && !msl_options.ios_use_simdgroup_functions)
+			if (msl_options.use_quadgroup_operation())
 				statement("return !!quad_shuffle_up((ushort)value, delta);");
 				statement("return !!quad_shuffle_up((ushort)value, delta);");
 			else
 			else
 				statement("return !!simd_shuffle_up((ushort)value, delta);");
 				statement("return !!simd_shuffle_up((ushort)value, delta);");
@@ -5665,7 +5746,7 @@ void CompilerMSL::emit_custom_functions()
 			statement("template<uint N>");
 			statement("template<uint N>");
 			statement("inline vec<bool, N> spvSubgroupShuffleUp(vec<bool, N> value, ushort delta)");
 			statement("inline vec<bool, N> spvSubgroupShuffleUp(vec<bool, N> value, ushort delta)");
 			begin_scope();
 			begin_scope();
-			if (msl_options.is_ios() && !msl_options.ios_use_simdgroup_functions)
+			if (msl_options.use_quadgroup_operation())
 				statement("return (vec<bool, N>)quad_shuffle_up((vec<ushort, N>)value, delta);");
 				statement("return (vec<bool, N>)quad_shuffle_up((vec<ushort, N>)value, delta);");
 			else
 			else
 				statement("return (vec<bool, N>)simd_shuffle_up((vec<ushort, N>)value, delta);");
 				statement("return (vec<bool, N>)simd_shuffle_up((vec<ushort, N>)value, delta);");
@@ -5677,7 +5758,7 @@ void CompilerMSL::emit_custom_functions()
 			statement("template<typename T>");
 			statement("template<typename T>");
 			statement("inline T spvSubgroupShuffleDown(T value, ushort delta)");
 			statement("inline T spvSubgroupShuffleDown(T value, ushort delta)");
 			begin_scope();
 			begin_scope();
-			if (msl_options.is_ios() && !msl_options.ios_use_simdgroup_functions)
+			if (msl_options.use_quadgroup_operation())
 				statement("return quad_shuffle_down(value, delta);");
 				statement("return quad_shuffle_down(value, delta);");
 			else
 			else
 				statement("return simd_shuffle_down(value, delta);");
 				statement("return simd_shuffle_down(value, delta);");
@@ -5686,7 +5767,7 @@ void CompilerMSL::emit_custom_functions()
 			statement("template<>");
 			statement("template<>");
 			statement("inline bool spvSubgroupShuffleDown(bool value, ushort delta)");
 			statement("inline bool spvSubgroupShuffleDown(bool value, ushort delta)");
 			begin_scope();
 			begin_scope();
-			if (msl_options.is_ios() && !msl_options.ios_use_simdgroup_functions)
+			if (msl_options.use_quadgroup_operation())
 				statement("return !!quad_shuffle_down((ushort)value, delta);");
 				statement("return !!quad_shuffle_down((ushort)value, delta);");
 			else
 			else
 				statement("return !!simd_shuffle_down((ushort)value, delta);");
 				statement("return !!simd_shuffle_down((ushort)value, delta);");
@@ -5695,7 +5776,7 @@ void CompilerMSL::emit_custom_functions()
 			statement("template<uint N>");
 			statement("template<uint N>");
 			statement("inline vec<bool, N> spvSubgroupShuffleDown(vec<bool, N> value, ushort delta)");
 			statement("inline vec<bool, N> spvSubgroupShuffleDown(vec<bool, N> value, ushort delta)");
 			begin_scope();
 			begin_scope();
-			if (msl_options.is_ios() && !msl_options.ios_use_simdgroup_functions)
+			if (msl_options.use_quadgroup_operation())
 				statement("return (vec<bool, N>)quad_shuffle_down((vec<ushort, N>)value, delta);");
 				statement("return (vec<bool, N>)quad_shuffle_down((vec<ushort, N>)value, delta);");
 			else
 			else
 				statement("return (vec<bool, N>)simd_shuffle_down((vec<ushort, N>)value, delta);");
 				statement("return (vec<bool, N>)simd_shuffle_down((vec<ushort, N>)value, delta);");
@@ -7368,7 +7449,7 @@ bool CompilerMSL::emit_tessellation_access_chain(const uint32_t *ops, uint32_t l
 		    expr_type->vecsize > result_ptr_type.vecsize)
 		    expr_type->vecsize > result_ptr_type.vecsize)
 			e += vector_swizzle(result_ptr_type.vecsize, 0);
 			e += vector_swizzle(result_ptr_type.vecsize, 0);
 
 
-		auto &expr = set<SPIRExpression>(ops[1], move(e), ops[0], should_forward(ops[2]));
+		auto &expr = set<SPIRExpression>(ops[1], std::move(e), ops[0], should_forward(ops[2]));
 		expr.loaded_from = var->self;
 		expr.loaded_from = var->self;
 		expr.need_transpose = meta.need_transpose;
 		expr.need_transpose = meta.need_transpose;
 		expr.access_chain = true;
 		expr.access_chain = true;
@@ -7559,6 +7640,8 @@ void CompilerMSL::emit_instruction(const Instruction &instruction)
 	auto ops = stream(instruction);
 	auto ops = stream(instruction);
 	auto opcode = static_cast<Op>(instruction.op);
 	auto opcode = static_cast<Op>(instruction.op);
 
 
+	opcode = get_remapped_spirv_op(opcode);
+
 	// If we need to do implicit bitcasts, make sure we do it with the correct type.
 	// If we need to do implicit bitcasts, make sure we do it with the correct type.
 	uint32_t integer_width = get_integer_width_for_instruction(instruction);
 	uint32_t integer_width = get_integer_width_for_instruction(instruction);
 	auto int_type = to_signed_basetype(integer_width);
 	auto int_type = to_signed_basetype(integer_width);
@@ -7601,6 +7684,10 @@ void CompilerMSL::emit_instruction(const Instruction &instruction)
 
 
 	case OpLogicalNotEqual:
 	case OpLogicalNotEqual:
 	case OpFOrdNotEqual:
 	case OpFOrdNotEqual:
+		// TODO: Should probably negate the == result here.
+		// Typically OrdNotEqual comes from GLSL which itself does not really specify what
+		// happens with NaN.
+		// Consider fixing this if we run into real issues.
 		MSL_BOP(!=);
 		MSL_BOP(!=);
 		break;
 		break;
 
 
@@ -7657,7 +7744,9 @@ void CompilerMSL::emit_instruction(const Instruction &instruction)
 		break;
 		break;
 
 
 	case OpFUnordNotEqual:
 	case OpFUnordNotEqual:
-		MSL_UNORD_BOP(!=);
+		// not equal in MSL generates une opcodes to begin with.
+		// Since unordered not equal is how it works in C, just inherit that behavior.
+		MSL_BOP(!=);
 		break;
 		break;
 
 
 	case OpFUnordGreaterThan:
 	case OpFUnordGreaterThan:
@@ -8920,6 +9009,8 @@ void CompilerMSL::emit_glsl_op(uint32_t result_type, uint32_t id, uint32_t eop,
 	auto int_type = to_signed_basetype(integer_width);
 	auto int_type = to_signed_basetype(integer_width);
 	auto uint_type = to_unsigned_basetype(integer_width);
 	auto uint_type = to_unsigned_basetype(integer_width);
 
 
+	op = get_remapped_glsl_op(op);
+
 	switch (op)
 	switch (op)
 	{
 	{
 	case GLSLstd450Sinh:
 	case GLSLstd450Sinh:
@@ -9335,6 +9426,9 @@ void CompilerMSL::emit_function_prototype(SPIRFunction &func, const Bitset &)
 		else
 		else
 			decl += entry_point_args_classic(!func.arguments.empty());
 			decl += entry_point_args_classic(!func.arguments.empty());
 
 
+		// append entry point args to avoid conflicts in local variable names.
+		local_variable_names.insert(resource_names.begin(), resource_names.end());
+
 		// If entry point function has variables that require early declaration,
 		// If entry point function has variables that require early declaration,
 		// ensure they each have an empty initializer, creating one if needed.
 		// ensure they each have an empty initializer, creating one if needed.
 		// This is done at this late stage because the initialization expression
 		// This is done at this late stage because the initialization expression
@@ -10567,7 +10661,7 @@ string CompilerMSL::convert_row_major_matrix(string exp_str, const SPIRType &exp
 {
 {
 	if (!is_matrix(exp_type))
 	if (!is_matrix(exp_type))
 	{
 	{
-		return CompilerGLSL::convert_row_major_matrix(move(exp_str), exp_type, physical_type_id, is_packed);
+		return CompilerGLSL::convert_row_major_matrix(std::move(exp_str), exp_type, physical_type_id, is_packed);
 	}
 	}
 	else
 	else
 	{
 	{
@@ -12077,16 +12171,6 @@ void CompilerMSL::fix_up_shader_inputs_outputs()
 					});
 					});
 				}
 				}
 				break;
 				break;
-			case BuiltInHelperInvocation:
-				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.");
-
-				entry_func.fixup_hooks_in.push_back([=]() {
-					statement(builtin_type_decl(bi_type), " ", to_expression(var_id), " = simd_is_helper_thread();");
-				});
-				break;
 			case BuiltInInvocationId:
 			case BuiltInInvocationId:
 				// This is direct-mapped without multi-patch workgroups.
 				// This is direct-mapped without multi-patch workgroups.
 				if (get_execution_model() != ExecutionModelTessellationControl || !msl_options.multi_patch_workgroup)
 				if (get_execution_model() != ExecutionModelTessellationControl || !msl_options.multi_patch_workgroup)
@@ -12979,8 +13063,8 @@ string CompilerMSL::to_name(uint32_t id, bool allow_alias) const
 	return Compiler::to_name(id, allow_alias);
 	return Compiler::to_name(id, allow_alias);
 }
 }
 
 
-// Returns a name that combines the name of the struct with the name of the member, except for Builtins
-string CompilerMSL::to_qualified_member_name(const SPIRType &type, uint32_t index)
+// Appends the name of the member to the variable qualifier string, except for Builtins.
+string CompilerMSL::append_member_name(const string &qualifier, const SPIRType &type, uint32_t index)
 {
 {
 	// Don't qualify Builtin names because they are unique and are treated as such when building expressions
 	// Don't qualify Builtin names because they are unique and are treated as such when building expressions
 	BuiltIn builtin = BuiltInMax;
 	BuiltIn builtin = BuiltInMax;
@@ -12991,7 +13075,7 @@ string CompilerMSL::to_qualified_member_name(const SPIRType &type, uint32_t inde
 	string mbr_name = to_member_name(type, index);
 	string mbr_name = to_member_name(type, index);
 	size_t startPos = mbr_name.find_first_not_of("_");
 	size_t startPos = mbr_name.find_first_not_of("_");
 	mbr_name = (startPos != string::npos) ? mbr_name.substr(startPos) : "";
 	mbr_name = (startPos != string::npos) ? mbr_name.substr(startPos) : "";
-	return join(to_name(type.self), "_", mbr_name);
+	return join(qualifier, "_", mbr_name);
 }
 }
 
 
 // Ensures that the specified name is permanently usable by prepending a prefix
 // Ensures that the specified name is permanently usable by prepending a prefix
@@ -13969,7 +14053,7 @@ void CompilerMSL::emit_subgroup_op(const Instruction &i)
 	switch (op)
 	switch (op)
 	{
 	{
 	case OpGroupNonUniformElect:
 	case OpGroupNonUniformElect:
-		if (msl_options.is_ios() && !msl_options.ios_use_simdgroup_functions)
+		if (msl_options.use_quadgroup_operation())
 			emit_op(result_type, id, "quad_is_first()", false);
 			emit_op(result_type, id, "quad_is_first()", false);
 		else
 		else
 			emit_op(result_type, id, "simd_is_first()", false);
 			emit_op(result_type, id, "simd_is_first()", false);
@@ -14042,14 +14126,14 @@ void CompilerMSL::emit_subgroup_op(const Instruction &i)
 		break;
 		break;
 
 
 	case OpGroupNonUniformAll:
 	case OpGroupNonUniformAll:
-		if (msl_options.is_ios() && !msl_options.ios_use_simdgroup_functions)
+		if (msl_options.use_quadgroup_operation())
 			emit_unary_func_op(result_type, id, ops[3], "quad_all");
 			emit_unary_func_op(result_type, id, ops[3], "quad_all");
 		else
 		else
 			emit_unary_func_op(result_type, id, ops[3], "simd_all");
 			emit_unary_func_op(result_type, id, ops[3], "simd_all");
 		break;
 		break;
 
 
 	case OpGroupNonUniformAny:
 	case OpGroupNonUniformAny:
-		if (msl_options.is_ios() && !msl_options.ios_use_simdgroup_functions)
+		if (msl_options.use_quadgroup_operation())
 			emit_unary_func_op(result_type, id, ops[3], "quad_any");
 			emit_unary_func_op(result_type, id, ops[3], "quad_any");
 		else
 		else
 			emit_unary_func_op(result_type, id, ops[3], "simd_any");
 			emit_unary_func_op(result_type, id, ops[3], "simd_any");
@@ -14381,6 +14465,14 @@ string CompilerMSL::builtin_to_glsl(BuiltIn builtin, StorageClass storage)
 		}
 		}
 		break;
 		break;
 
 
+	case BuiltInHelperInvocation:
+		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.");
+		// In SPIR-V 1.6 with Volatile HelperInvocation, we cannot emit a fixup early.
+		return "simd_is_helper_thread()";
+
 	default:
 	default:
 		break;
 		break;
 	}
 	}
@@ -14547,7 +14639,7 @@ string CompilerMSL::builtin_qualifier(BuiltIn builtin)
 			SPIRV_CROSS_THROW("NumSubgroups is handled specially with emulation.");
 			SPIRV_CROSS_THROW("NumSubgroups is handled specially with emulation.");
 		if (!msl_options.supports_msl_version(2))
 		if (!msl_options.supports_msl_version(2))
 			SPIRV_CROSS_THROW("Subgroup builtins require Metal 2.0.");
 			SPIRV_CROSS_THROW("Subgroup builtins require Metal 2.0.");
-		return msl_options.is_ios() ? "quadgroups_per_threadgroup" : "simdgroups_per_threadgroup";
+		return msl_options.use_quadgroup_operation() ? "quadgroups_per_threadgroup" : "simdgroups_per_threadgroup";
 
 
 	case BuiltInSubgroupId:
 	case BuiltInSubgroupId:
 		if (msl_options.emulate_subgroups)
 		if (msl_options.emulate_subgroups)
@@ -14555,7 +14647,7 @@ string CompilerMSL::builtin_qualifier(BuiltIn builtin)
 			SPIRV_CROSS_THROW("SubgroupId is handled specially with emulation.");
 			SPIRV_CROSS_THROW("SubgroupId is handled specially with emulation.");
 		if (!msl_options.supports_msl_version(2))
 		if (!msl_options.supports_msl_version(2))
 			SPIRV_CROSS_THROW("Subgroup builtins require Metal 2.0.");
 			SPIRV_CROSS_THROW("Subgroup builtins require Metal 2.0.");
-		return msl_options.is_ios() ? "quadgroup_index_in_threadgroup" : "simdgroup_index_in_threadgroup";
+		return msl_options.use_quadgroup_operation() ? "quadgroup_index_in_threadgroup" : "simdgroup_index_in_threadgroup";
 
 
 	case BuiltInSubgroupLocalInvocationId:
 	case BuiltInSubgroupLocalInvocationId:
 		if (msl_options.emulate_subgroups)
 		if (msl_options.emulate_subgroups)
@@ -14574,7 +14666,7 @@ string CompilerMSL::builtin_qualifier(BuiltIn builtin)
 			// We are generating a Metal kernel function.
 			// We are generating a Metal kernel function.
 			if (!msl_options.supports_msl_version(2))
 			if (!msl_options.supports_msl_version(2))
 				SPIRV_CROSS_THROW("Subgroup builtins in kernel functions require Metal 2.0.");
 				SPIRV_CROSS_THROW("Subgroup builtins in kernel functions require Metal 2.0.");
-			return msl_options.is_ios() ? "thread_index_in_quadgroup" : "thread_index_in_simdgroup";
+			return msl_options.use_quadgroup_operation() ? "thread_index_in_quadgroup" : "thread_index_in_simdgroup";
 		}
 		}
 		else
 		else
 			SPIRV_CROSS_THROW("Subgroup builtins are not available in this type of function.");
 			SPIRV_CROSS_THROW("Subgroup builtins are not available in this type of function.");

+ 21 - 8
3rdparty/spirv-cross/spirv_msl.hpp

@@ -393,7 +393,7 @@ public:
 		// and will be addressed using the current ViewIndex.
 		// and will be addressed using the current ViewIndex.
 		bool arrayed_subpass_input = false;
 		bool arrayed_subpass_input = false;
 
 
-		// Whether to use SIMD-group or quadgroup functions to implement group nnon-uniform
+		// Whether to use SIMD-group or quadgroup functions to implement group non-uniform
 		// operations. Some GPUs on iOS do not support the SIMD-group functions, only the
 		// operations. Some GPUs on iOS do not support the SIMD-group functions, only the
 		// quadgroup functions.
 		// quadgroup functions.
 		bool ios_use_simdgroup_functions = false;
 		bool ios_use_simdgroup_functions = false;
@@ -445,6 +445,11 @@ public:
 			return platform == macOS;
 			return platform == macOS;
 		}
 		}
 
 
+		bool use_quadgroup_operation() const
+		{
+			return is_ios() && !ios_use_simdgroup_functions;
+		}
+
 		void set_msl_version(uint32_t major, uint32_t minor = 0, uint32_t patch = 0)
 		void set_msl_version(uint32_t major, uint32_t minor = 0, uint32_t patch = 0)
 		{
 		{
 			msl_version = make_msl_version(major, minor, patch);
 			msl_version = make_msl_version(major, minor, patch);
@@ -826,12 +831,20 @@ protected:
 	bool add_component_variable_to_interface_block(spv::StorageClass storage, const std::string &ib_var_ref,
 	bool add_component_variable_to_interface_block(spv::StorageClass storage, const std::string &ib_var_ref,
 	                                               SPIRVariable &var, const SPIRType &type,
 	                                               SPIRVariable &var, const SPIRType &type,
 	                                               InterfaceBlockMeta &meta);
 	                                               InterfaceBlockMeta &meta);
-	void add_plain_member_variable_to_interface_block(spv::StorageClass storage, const std::string &ib_var_ref,
-	                                                  SPIRType &ib_type, SPIRVariable &var, uint32_t index,
-	                                                  InterfaceBlockMeta &meta);
-	void add_composite_member_variable_to_interface_block(spv::StorageClass storage, const std::string &ib_var_ref,
-	                                                      SPIRType &ib_type, SPIRVariable &var, uint32_t index,
-	                                                      InterfaceBlockMeta &meta);
+	void add_plain_member_variable_to_interface_block(spv::StorageClass storage,
+	                                                  const std::string &ib_var_ref, SPIRType &ib_type,
+	                                                  SPIRVariable &var, SPIRType &var_type,
+	                                                  uint32_t mbr_idx, InterfaceBlockMeta &meta,
+	                                                  const std::string &mbr_name_qual,
+	                                                  const std::string &var_chain_qual,
+	                                                  uint32_t &location, uint32_t &var_mbr_idx);
+	void add_composite_member_variable_to_interface_block(spv::StorageClass storage,
+	                                                      const std::string &ib_var_ref, SPIRType &ib_type,
+	                                                      SPIRVariable &var, SPIRType &var_type,
+	                                                      uint32_t mbr_idx, InterfaceBlockMeta &meta,
+	                                                      const std::string &mbr_name_qual,
+	                                                      const std::string &var_chain_qual,
+	                                                      uint32_t &location, uint32_t &var_mbr_idx);
 	void add_tess_level_input_to_interface_block(const std::string &ib_var_ref, SPIRType &ib_type, SPIRVariable &var);
 	void add_tess_level_input_to_interface_block(const std::string &ib_var_ref, SPIRType &ib_type, SPIRVariable &var);
 
 
 	void fix_up_interface_member_indices(spv::StorageClass storage, uint32_t ib_type_id);
 	void fix_up_interface_member_indices(spv::StorageClass storage, uint32_t ib_type_id);
@@ -858,7 +871,7 @@ protected:
 	std::string entry_point_arg_stage_in();
 	std::string entry_point_arg_stage_in();
 	void entry_point_args_builtin(std::string &args);
 	void entry_point_args_builtin(std::string &args);
 	void entry_point_args_discrete_descriptors(std::string &args);
 	void entry_point_args_discrete_descriptors(std::string &args);
-	std::string to_qualified_member_name(const SPIRType &type, uint32_t index);
+	std::string append_member_name(const std::string &qualifier, const SPIRType &type, uint32_t index);
 	std::string ensure_valid_name(std::string name, std::string pfx);
 	std::string ensure_valid_name(std::string name, std::string pfx);
 	std::string to_sampler_expression(uint32_t id);
 	std::string to_sampler_expression(uint32_t id);
 	std::string to_swizzle_expression(uint32_t id);
 	std::string to_swizzle_expression(uint32_t id);

+ 12 - 4
3rdparty/spirv-cross/spirv_parser.cpp

@@ -31,7 +31,7 @@ namespace SPIRV_CROSS_NAMESPACE
 {
 {
 Parser::Parser(vector<uint32_t> spirv)
 Parser::Parser(vector<uint32_t> spirv)
 {
 {
-	ir.spirv = move(spirv);
+	ir.spirv = std::move(spirv);
 }
 }
 
 
 Parser::Parser(const uint32_t *spirv_data, size_t word_count)
 Parser::Parser(const uint32_t *spirv_data, size_t word_count)
@@ -259,7 +259,7 @@ void Parser::parse(const Instruction &instruction)
 	case OpExtension:
 	case OpExtension:
 	{
 	{
 		auto ext = extract_string(ir.spirv, instruction.offset);
 		auto ext = extract_string(ir.spirv, instruction.offset);
-		ir.declared_extensions.push_back(move(ext));
+		ir.declared_extensions.push_back(std::move(ext));
 		break;
 		break;
 	}
 	}
 
 
@@ -291,7 +291,15 @@ void Parser::parse(const Instruction &instruction)
 	{
 	{
 		// The SPIR-V debug information extended instructions might come at global scope.
 		// The SPIR-V debug information extended instructions might come at global scope.
 		if (current_block)
 		if (current_block)
+		{
 			current_block->ops.push_back(instruction);
 			current_block->ops.push_back(instruction);
+			if (length >= 2)
+			{
+				const auto *type = maybe_get<SPIRType>(ops[0]);
+				if (type)
+					ir.load_type_width.insert({ ops[1], type->width });
+			}
+		}
 		break;
 		break;
 	}
 	}
 
 
@@ -1059,6 +1067,7 @@ void Parser::parse(const Instruction &instruction)
 	}
 	}
 
 
 	case OpKill:
 	case OpKill:
+	case OpTerminateInvocation:
 	{
 	{
 		if (!current_block)
 		if (!current_block)
 			SPIRV_CROSS_THROW("Trying to end a non-existing block.");
 			SPIRV_CROSS_THROW("Trying to end a non-existing block.");
@@ -1211,10 +1220,9 @@ void Parser::parse(const Instruction &instruction)
 		{
 		{
 			const auto *type = maybe_get<SPIRType>(ops[0]);
 			const auto *type = maybe_get<SPIRType>(ops[0]);
 			if (type)
 			if (type)
-			{
 				ir.load_type_width.insert({ ops[1], type->width });
 				ir.load_type_width.insert({ ops[1], type->width });
-			}
 		}
 		}
+
 		if (!current_block)
 		if (!current_block)
 			SPIRV_CROSS_THROW("Currently no block to insert opcode.");
 			SPIRV_CROSS_THROW("Currently no block to insert opcode.");