update SPIRV-Cross to KhronosGroup/SPIRV-Cross@6c09849

2026-08-12 00:20:52 +02:00 · 2026-07-24 16:33:18 -03:00
parent e4649253e8
commit 17258d8179
22 changed files with 12655 additions and 2205 deletions
@@ -0,0 +1,171 @@
+// Copyright (c) 2018-2024 The Khronos Group Inc.
+// 
+// Permission is hereby granted, free of charge, to any person obtaining a copy
+// of this software and/or associated documentation files (the "Materials"),
+// to deal in the Materials without restriction, including without limitation
+// the rights to use, copy, modify, merge, publish, distribute, sublicense,
+// and/or sell copies of the Materials, and to permit persons to whom the
+// Materials are furnished to do so, subject to the following conditions:
+// 
+// The above copyright notice and this permission notice shall be included in
+// all copies or substantial portions of the Materials.
+// 
+// MODIFICATIONS TO THIS FILE MAY MEAN IT NO LONGER ACCURATELY REFLECTS KHRONOS
+// STANDARDS. THE UNMODIFIED, NORMATIVE VERSIONS OF KHRONOS SPECIFICATIONS AND
+// HEADER INFORMATION ARE LOCATED AT https://www.khronos.org/registry/ 
+// 
+// THE MATERIALS ARE PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS
+// OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL
+// THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING
+// FROM,OUT OF OR IN CONNECTION WITH THE MATERIALS OR THE USE OR OTHER DEALINGS
+// IN THE MATERIALS.
+
+#ifndef SPIRV_UNIFIED1_NonSemanticShaderDebugInfo100_H_
+#define SPIRV_UNIFIED1_NonSemanticShaderDebugInfo100_H_
+
+#ifdef __cplusplus
+extern "C" {
+#endif
+
+enum {
+    NonSemanticShaderDebugInfo100Version = 100,
+    NonSemanticShaderDebugInfo100Version_BitWidthPadding = 0x7fffffff
+};
+enum {
+    NonSemanticShaderDebugInfo100Revision = 6,
+    NonSemanticShaderDebugInfo100Revision_BitWidthPadding = 0x7fffffff
+};
+
+enum NonSemanticShaderDebugInfo100Instructions {
+    NonSemanticShaderDebugInfo100DebugInfoNone = 0,
+    NonSemanticShaderDebugInfo100DebugCompilationUnit = 1,
+    NonSemanticShaderDebugInfo100DebugTypeBasic = 2,
+    NonSemanticShaderDebugInfo100DebugTypePointer = 3,
+    NonSemanticShaderDebugInfo100DebugTypeQualifier = 4,
+    NonSemanticShaderDebugInfo100DebugTypeArray = 5,
+    NonSemanticShaderDebugInfo100DebugTypeVector = 6,
+    NonSemanticShaderDebugInfo100DebugTypedef = 7,
+    NonSemanticShaderDebugInfo100DebugTypeFunction = 8,
+    NonSemanticShaderDebugInfo100DebugTypeEnum = 9,
+    NonSemanticShaderDebugInfo100DebugTypeComposite = 10,
+    NonSemanticShaderDebugInfo100DebugTypeMember = 11,
+    NonSemanticShaderDebugInfo100DebugTypeInheritance = 12,
+    NonSemanticShaderDebugInfo100DebugTypePtrToMember = 13,
+    NonSemanticShaderDebugInfo100DebugTypeTemplate = 14,
+    NonSemanticShaderDebugInfo100DebugTypeTemplateParameter = 15,
+    NonSemanticShaderDebugInfo100DebugTypeTemplateTemplateParameter = 16,
+    NonSemanticShaderDebugInfo100DebugTypeTemplateParameterPack = 17,
+    NonSemanticShaderDebugInfo100DebugGlobalVariable = 18,
+    NonSemanticShaderDebugInfo100DebugFunctionDeclaration = 19,
+    NonSemanticShaderDebugInfo100DebugFunction = 20,
+    NonSemanticShaderDebugInfo100DebugLexicalBlock = 21,
+    NonSemanticShaderDebugInfo100DebugLexicalBlockDiscriminator = 22,
+    NonSemanticShaderDebugInfo100DebugScope = 23,
+    NonSemanticShaderDebugInfo100DebugNoScope = 24,
+    NonSemanticShaderDebugInfo100DebugInlinedAt = 25,
+    NonSemanticShaderDebugInfo100DebugLocalVariable = 26,
+    NonSemanticShaderDebugInfo100DebugInlinedVariable = 27,
+    NonSemanticShaderDebugInfo100DebugDeclare = 28,
+    NonSemanticShaderDebugInfo100DebugValue = 29,
+    NonSemanticShaderDebugInfo100DebugOperation = 30,
+    NonSemanticShaderDebugInfo100DebugExpression = 31,
+    NonSemanticShaderDebugInfo100DebugMacroDef = 32,
+    NonSemanticShaderDebugInfo100DebugMacroUndef = 33,
+    NonSemanticShaderDebugInfo100DebugImportedEntity = 34,
+    NonSemanticShaderDebugInfo100DebugSource = 35,
+    NonSemanticShaderDebugInfo100DebugFunctionDefinition = 101,
+    NonSemanticShaderDebugInfo100DebugSourceContinued = 102,
+    NonSemanticShaderDebugInfo100DebugLine = 103,
+    NonSemanticShaderDebugInfo100DebugNoLine = 104,
+    NonSemanticShaderDebugInfo100DebugBuildIdentifier = 105,
+    NonSemanticShaderDebugInfo100DebugStoragePath = 106,
+    NonSemanticShaderDebugInfo100DebugEntryPoint = 107,
+    NonSemanticShaderDebugInfo100DebugTypeMatrix = 108,
+    NonSemanticShaderDebugInfo100InstructionsMax = 0x7fffffff
+};
+
+
+enum NonSemanticShaderDebugInfo100DebugInfoFlags {
+    NonSemanticShaderDebugInfo100None = 0x0000,
+    NonSemanticShaderDebugInfo100FlagIsProtected = 0x01,
+    NonSemanticShaderDebugInfo100FlagIsPrivate = 0x02,
+    NonSemanticShaderDebugInfo100FlagIsPublic = 0x03,
+    NonSemanticShaderDebugInfo100FlagIsLocal = 0x04,
+    NonSemanticShaderDebugInfo100FlagIsDefinition = 0x08,
+    NonSemanticShaderDebugInfo100FlagFwdDecl = 0x10,
+    NonSemanticShaderDebugInfo100FlagArtificial = 0x20,
+    NonSemanticShaderDebugInfo100FlagExplicit = 0x40,
+    NonSemanticShaderDebugInfo100FlagPrototyped = 0x80,
+    NonSemanticShaderDebugInfo100FlagObjectPointer = 0x100,
+    NonSemanticShaderDebugInfo100FlagStaticMember = 0x200,
+    NonSemanticShaderDebugInfo100FlagIndirectVariable = 0x400,
+    NonSemanticShaderDebugInfo100FlagLValueReference = 0x800,
+    NonSemanticShaderDebugInfo100FlagRValueReference = 0x1000,
+    NonSemanticShaderDebugInfo100FlagIsOptimized = 0x2000,
+    NonSemanticShaderDebugInfo100FlagIsEnumClass = 0x4000,
+    NonSemanticShaderDebugInfo100FlagTypePassByValue = 0x8000,
+    NonSemanticShaderDebugInfo100FlagTypePassByReference = 0x10000,
+    NonSemanticShaderDebugInfo100FlagUnknownPhysicalLayout = 0x20000,
+    NonSemanticShaderDebugInfo100DebugInfoFlagsMax = 0x7fffffff
+};
+
+enum NonSemanticShaderDebugInfo100BuildIdentifierFlags {
+    NonSemanticShaderDebugInfo100IdentifierPossibleDuplicates = 0x01,
+    NonSemanticShaderDebugInfo100BuildIdentifierFlagsMax = 0x7fffffff
+};
+
+enum NonSemanticShaderDebugInfo100DebugBaseTypeAttributeEncoding {
+    NonSemanticShaderDebugInfo100Unspecified = 0,
+    NonSemanticShaderDebugInfo100Address = 1,
+    NonSemanticShaderDebugInfo100Boolean = 2,
+    NonSemanticShaderDebugInfo100Float = 3,
+    NonSemanticShaderDebugInfo100Signed = 4,
+    NonSemanticShaderDebugInfo100SignedChar = 5,
+    NonSemanticShaderDebugInfo100Unsigned = 6,
+    NonSemanticShaderDebugInfo100UnsignedChar = 7,
+    NonSemanticShaderDebugInfo100DebugBaseTypeAttributeEncodingMax = 0x7fffffff
+};
+
+enum NonSemanticShaderDebugInfo100DebugCompositeType {
+    NonSemanticShaderDebugInfo100Class = 0,
+    NonSemanticShaderDebugInfo100Structure = 1,
+    NonSemanticShaderDebugInfo100Union = 2,
+    NonSemanticShaderDebugInfo100DebugCompositeTypeMax = 0x7fffffff
+};
+
+enum NonSemanticShaderDebugInfo100DebugTypeQualifier {
+    NonSemanticShaderDebugInfo100ConstType = 0,
+    NonSemanticShaderDebugInfo100VolatileType = 1,
+    NonSemanticShaderDebugInfo100RestrictType = 2,
+    NonSemanticShaderDebugInfo100AtomicType = 3,
+    NonSemanticShaderDebugInfo100DebugTypeQualifierMax = 0x7fffffff
+};
+
+enum NonSemanticShaderDebugInfo100DebugOperation {
+    NonSemanticShaderDebugInfo100Deref = 0,
+    NonSemanticShaderDebugInfo100Plus = 1,
+    NonSemanticShaderDebugInfo100Minus = 2,
+    NonSemanticShaderDebugInfo100PlusUconst = 3,
+    NonSemanticShaderDebugInfo100BitPiece = 4,
+    NonSemanticShaderDebugInfo100Swap = 5,
+    NonSemanticShaderDebugInfo100Xderef = 6,
+    NonSemanticShaderDebugInfo100StackValue = 7,
+    NonSemanticShaderDebugInfo100Constu = 8,
+    NonSemanticShaderDebugInfo100Fragment = 9,
+    NonSemanticShaderDebugInfo100DebugOperationMax = 0x7fffffff
+};
+
+enum NonSemanticShaderDebugInfo100DebugImportedEntity {
+    NonSemanticShaderDebugInfo100ImportedModule = 0,
+    NonSemanticShaderDebugInfo100ImportedDeclaration = 1,
+    NonSemanticShaderDebugInfo100DebugImportedEntityMax = 0x7fffffff
+};
+
+
+#ifdef __cplusplus
+}
+#endif
+
+#endif // SPIRV_UNIFIED1_NonSemanticShaderDebugInfo100_H_
@@ -59,20 +59,26 @@ void CFG::build_immediate_dominators()
 	for (auto i = post_order.size(); i; i--)
 	{
 		uint32_t block = post_order[i - 1];
-		auto &pred = preceding_edges[block];
-		if (pred.empty()) // This is for the entry block, but we've already set up the dominators.
-			continue;

-		for (auto &edge : pred)
+		const auto resolve_preds = [&](const SmallVector<uint32_t> &pred)
 		{
-			if (immediate_dominators[block])
+			if (pred.empty()) // This is for the entry block, but we've already set up the dominators.
+				return;
+
+			for (auto &edge : pred)
 			{
-				assert(immediate_dominators[edge]);
-				immediate_dominators[block] = find_common_dominator(immediate_dominators[block], edge);
+				if (immediate_dominators[block])
+				{
+					assert(immediate_dominators[edge]);
+					immediate_dominators[block] = find_common_dominator(immediate_dominators[block], edge);
+				}
+				else
+					immediate_dominators[block] = edge;
 			}
-			else
-				immediate_dominators[block] = edge;
-		}
+		};
+
+		resolve_preds(preceding_edges[block]);
+		resolve_preds(virtual_dominance_preceding_edges[block]);
 	}
 }

@@ -81,31 +87,105 @@ bool CFG::is_back_edge(uint32_t to) const
 	// We have a back edge if the visit order is set with the temporary magic value 0.
 	// Crossing edges will have already been recorded with a visit order.
 	auto itr = visit_order.find(to);
-	return itr != end(visit_order) && itr->second.get() == 0;
+	return itr != end(visit_order) && itr->second.visited_branches && !itr->second.visited_resolve;
 }

-bool CFG::has_visited_forward_edge(uint32_t to) const
+bool CFG::has_visited_branch(uint32_t to) const
 {
-	// If > 0, we have visited the edge already, and this is not a back edge branch.
 	auto itr = visit_order.find(to);
-	return itr != end(visit_order) && itr->second.get() > 0;
+	return itr != end(visit_order) && itr->second.visited_branches;
 }

-bool CFG::post_order_visit(uint32_t block_id)
+void CFG::post_order_visit_entry(uint32_t block)
 {
-	// If we have already branched to this block (back edge), stop recursion.
-	// If our branches are back-edges, we do not record them.
-	// We have to record crossing edges however.
-	if (has_visited_forward_edge(block_id))
-		return true;
-	else if (is_back_edge(block_id))
-		return false;
+	visit_stack.push_back(block);

-	// Block back-edges from recursively revisiting ourselves.
-	visit_order[block_id].get() = 0;
+	while (!visit_stack.empty())
+	{
+		bool keep_iterating;
+		do
+		{
+			// Reverse the order to allow for stack-like behavior and preserves the visit order from recursive algorithm.
+			// Traverse depth first.
+			uint32_t to_visit = visit_stack.back();
+			last_visited_size = visit_stack.size();
+			post_order_visit_branches(to_visit);
+			keep_iterating = last_visited_size != visit_stack.size();
+			if (keep_iterating)
+				std::reverse(visit_stack.begin() + last_visited_size, visit_stack.end());
+		} while (keep_iterating);

+		// We've reached the end of some tree leaf. Resolve the stack.
+		// Any node which has been visited for real can be popped now.
+		while (!visit_stack.empty() && visit_order[visit_stack.back()].visited_branches)
+		{
+			post_order_visit_resolve(visit_stack.back());
+			visit_stack.pop_back();
+		}
+	}
+}
+
+void CFG::visit_branch(uint32_t block_id)
+{
+	// Prune obvious duplicates.
+	if (std::find(visit_stack.begin() + last_visited_size, visit_stack.end(), block_id) == visit_stack.end() &&
+	    !has_visited_branch(block_id))
+	{
+		visit_stack.push_back(block_id);
+	}
+}
+
+void CFG::post_order_visit_branches(uint32_t block_id)
+{
 	auto &block = compiler.get<SPIRBlock>(block_id);

+	auto &visit = visit_order[block_id];
+	if (visit.visited_branches)
+		return;
+	visit.visited_branches = true;
+
+	if (block.merge == SPIRBlock::MergeLoop)
+		visit_branch(block.merge_block);
+	else if (block.merge == SPIRBlock::MergeSelection)
+		visit_branch(block.next_block);
+
+	// First visit our branch targets.
+	switch (block.terminator)
+	{
+	case SPIRBlock::Direct:
+		visit_branch(block.next_block);
+		break;
+
+	case SPIRBlock::Select:
+		visit_branch(block.true_block);
+		visit_branch(block.false_block);
+		break;
+
+	case SPIRBlock::MultiSelect:
+	{
+		const auto &cases = compiler.get_case_list(block);
+		for (const auto &target : cases)
+			visit_branch(target.block);
+		if (block.default_block)
+			visit_branch(block.default_block);
+		break;
+	}
+
+	default:
+		break;
+	}
+}
+
+void CFG::post_order_visit_resolve(uint32_t block_id)
+{
+	auto &block = compiler.get<SPIRBlock>(block_id);
+
+	auto &visit_block = visit_order[block_id];
+	assert(visit_block.visited_branches);
+	auto &visited = visit_order[block_id].visited_resolve;
+	if (visited)
+		return;
+
 	// If this is a loop header, add an implied branch to the merge target.
 	// This is needed to avoid annoying cases with do { ... } while(false) loops often generated by inliners.
 	// To the CFG, this is linear control flow, but we risk picking the do/while scope as our dominating block.
@@ -116,21 +196,28 @@ bool CFG::post_order_visit(uint32_t block_id)
 	// is lower than inside the loop, which is going to be key for some traversal algorithms like post-dominance analysis.
 	// For selection constructs true/false blocks will end up visiting the merge block directly and it works out fine,
 	// but for loops, only the header might end up actually branching to merge block.
-	if (block.merge == SPIRBlock::MergeLoop && post_order_visit(block.merge_block))
+	if (block.merge == SPIRBlock::MergeLoop && !is_back_edge(block.merge_block))
 		add_branch(block_id, block.merge_block);

+	// Similar case as do/while loops, but expressed in a different form.
+	// if (true) { foo = 1; } else { return/unreachable/kill/blah; } access(foo);
+	// Only consider this branch when computing dominance to avoid breaking other analysis like
+	// parameter preservation.
+	if (block.merge == SPIRBlock::MergeSelection && !is_back_edge(block.next_block))
+		add_virtual_dominance_branch(block_id, block.next_block);
+
 	// First visit our branch targets.
 	switch (block.terminator)
 	{
 	case SPIRBlock::Direct:
-		if (post_order_visit(block.next_block))
+		if (!is_back_edge(block.next_block))
 			add_branch(block_id, block.next_block);
 		break;

 	case SPIRBlock::Select:
-		if (post_order_visit(block.true_block))
+		if (!is_back_edge(block.true_block))
 			add_branch(block_id, block.true_block);
-		if (post_order_visit(block.false_block))
+		if (!is_back_edge(block.false_block))
 			add_branch(block_id, block.false_block);
 		break;

@@ -139,10 +226,10 @@ bool CFG::post_order_visit(uint32_t block_id)
 		const auto &cases = compiler.get_case_list(block);
 		for (const auto &target : cases)
 		{
-			if (post_order_visit(target.block))
+			if (!is_back_edge(target.block))
 				add_branch(block_id, target.block);
 		}
-		if (block.default_block && post_order_visit(block.default_block))
+		if (block.default_block && !is_back_edge(block.default_block))
 			add_branch(block_id, block.default_block);
 		break;
 	}
@@ -157,7 +244,7 @@ bool CFG::post_order_visit(uint32_t block_id)
 	// We can use the variable without a Phi since there is only one possible parent here.
 	// However, in this case, we need to hoist out the inner variable to outside the branch.
 	// Use same strategy as loops.
-	if (block.merge == SPIRBlock::MergeSelection && post_order_visit(block.next_block))
+	if (block.merge == SPIRBlock::MergeSelection && !is_back_edge(block.next_block))
 	{
 		// If there is only one preceding edge to the merge block and it's not ourselves, we need a fixup.
 		// Add a fake branch so any dominator in either the if (), or else () block, or a lone case statement
@@ -201,10 +288,9 @@ bool CFG::post_order_visit(uint32_t block_id)
 		}
 	}

-	// Then visit ourselves. Start counting at one, to let 0 be a magic value for testing back vs. crossing edges.
-	visit_order[block_id].get() = ++visit_count;
+	visited = true;
+	visit_block.order = ++visit_count;
 	post_order.push_back(block_id);
-	return true;
 }

 void CFG::build_post_order_visit_order()
@@ -213,20 +299,29 @@ void CFG::build_post_order_visit_order()
 	visit_count = 0;
 	visit_order.clear();
 	post_order.clear();
-	post_order_visit(block);
+	post_order_visit_entry(block);
+}
+
+static void add_unique(SmallVector <uint32_t> &l, uint32_t value)
+{
+	auto itr = find(begin(l), end(l), value);
+	if (itr == end(l))
+		l.push_back(value);
 }

 void CFG::add_branch(uint32_t from, uint32_t to)
 {
-	const auto add_unique = [](SmallVector<uint32_t> &l, uint32_t value) {
-		auto itr = find(begin(l), end(l), value);
-		if (itr == end(l))
-			l.push_back(value);
-	};
+	assert(from && to);
 	add_unique(preceding_edges[to], from);
 	add_unique(succeeding_edges[from], to);
 }

+void CFG::add_virtual_dominance_branch(uint32_t from, uint32_t to)
+{
+	assert(from && to);
+	add_unique(virtual_dominance_preceding_edges[to], from);
+}
+
 uint32_t CFG::find_loop_dominator(uint32_t block_id) const
 {
 	while (block_id != SPIRBlock::NoDominator)
@@ -68,7 +68,7 @@ public:
 	{
 		auto itr = visit_order.find(block);
 		assert(itr != std::end(visit_order));
-		int v = itr->second.get();
+		int v = itr->second.order;
 		assert(v > 0);
 		return uint32_t(v);
 	}
@@ -114,22 +114,15 @@ public:
 private:
 	struct VisitOrder
 	{
-		int &get()
-		{
-			return v;
-		}
-
-		const int &get() const
-		{
-			return v;
-		}
-
-		int v = -1;
+		int order = -1;
+		bool visited_resolve = false;
+		bool visited_branches = false;
 	};

 	Compiler &compiler;
 	const SPIRFunction &func;
 	std::unordered_map<uint32_t, SmallVector<uint32_t>> preceding_edges;
+	std::unordered_map<uint32_t, SmallVector<uint32_t>> virtual_dominance_preceding_edges;
 	std::unordered_map<uint32_t, SmallVector<uint32_t>> succeeding_edges;
 	std::unordered_map<uint32_t, uint32_t> immediate_dominators;
 	std::unordered_map<uint32_t, VisitOrder> visit_order;
@@ -137,13 +130,20 @@ private:
 	SmallVector<uint32_t> empty_vector;

 	void add_branch(uint32_t from, uint32_t to);
+	void add_virtual_dominance_branch(uint32_t from, uint32_t to);
 	void build_post_order_visit_order();
 	void build_immediate_dominators();
-	bool post_order_visit(uint32_t block);
+	void post_order_visit_branches(uint32_t block);
+	void post_order_visit_resolve(uint32_t block);
+	void post_order_visit_entry(uint32_t block);
 	uint32_t visit_count = 0;

 	bool is_back_edge(uint32_t to) const;
-	bool has_visited_forward_edge(uint32_t to) const;
+	bool has_visited_branch(uint32_t to) const;
+	void visit_branch(uint32_t block_id);
+
+	SmallVector<uint32_t> visit_stack;
+	size_t last_visited_size = 0;
 };

 class DominatorBuilder
@@ -27,8 +27,17 @@
 #ifndef SPV_ENABLE_UTILITY_CODE
 #define SPV_ENABLE_UTILITY_CODE
 #endif
-#include "spirv.hpp"

+// Pragmatic hack to avoid symbol conflicts when including both hpp11 and hpp headers in same translation unit.
+// This is an unfortunate SPIRV-Headers issue that we cannot easily deal with ourselves.
+#ifdef SPIRV_CROSS_SPV_HEADER_NAMESPACE_OVERRIDE
+#define spv SPIRV_CROSS_SPV_HEADER_NAMESPACE_OVERRIDE
+#define SPIRV_CROSS_SPV_HEADER_NAMESPACE SPIRV_CROSS_SPV_HEADER_NAMESPACE_OVERRIDE
+#else
+#define SPIRV_CROSS_SPV_HEADER_NAMESPACE spv
+#endif
+
+#include "spirv.hpp"
 #include "spirv_cross_containers.hpp"
 #include "spirv_cross_error_handling.hpp"
 #include <functional>
@@ -220,7 +229,7 @@ static inline std::string convert_to_string(int32_t value)
 	// INT_MIN is ... special on some backends. If we use a decimal literal, and negate it, we
 	// could accidentally promote the literal to long first, then negate.
 	// To workaround it, emit int(0x80000000) instead.
-	if (value == std::numeric_limits<int32_t>::min())
+	if (value == (std::numeric_limits<int32_t>::min)())
 		return "int(0x80000000)";
 	else
 		return std::to_string(value);
@@ -231,7 +240,7 @@ static inline std::string convert_to_string(int64_t value, const std::string &in
 	// INT64_MIN is ... special on some backends.
 	// If we use a decimal literal, and negate it, we might overflow the representable numbers.
 	// To workaround it, emit int(0x80000000) instead.
-	if (value == std::numeric_limits<int64_t>::min())
+	if (value == (std::numeric_limits<int64_t>::min)())
 		return join(int64_type, "(0x8000000000000000u", (long_long_literal_suffix ? "ll" : "l"), ")");
 	else
 		return std::to_string(value) + (long_long_literal_suffix ? "ll" : "l");
@@ -368,6 +377,7 @@ enum Types
 	TypeAccessChain,
 	TypeUndef,
 	TypeString,
+	TypeDebugLocalVariable,
 	TypeCount
 };

@@ -497,6 +507,18 @@ struct SPIRString : IVariant
 	SPIRV_CROSS_DECLARE_CLONE(SPIRString)
 };

+struct SPIRDebugLocalVariable : IVariant
+{
+	enum
+	{
+		type = TypeDebugLocalVariable
+	};
+
+	uint32_t name_id;
+
+	SPIRV_CROSS_DECLARE_CLONE(SPIRDebugLocalVariable)
+};
+
 // This type is only used by backends which need to access the combined image and sampler IDs separately after
 // the OpSampledImage opcode.
 struct SPIRCombinedImageSampler : IVariant
@@ -574,11 +596,20 @@ struct SPIRType : IVariant
 		Sampler,
 		AccelerationStructure,
 		RayQuery,
+		CoopVecNV,

 		// Keep internal types at the end.
 		ControlPointArray,
 		Interpolant,
-		Char
+		Char,
+		// MSL specific type, that is used by 'object'(analog of 'task' from glsl) shader.
+		MeshGridProperties,
+		BFloat16,
+		FloatE4M3,
+		FloatE5M2,
+
+		Tensor,
+		DescriptorHeapBuffer
 	};

 	// Scalar/vector/matrix support.
@@ -603,6 +634,35 @@ struct SPIRType : IVariant
 	bool pointer = false;
 	bool forward_pointer = false;

+	union
+	{
+		struct
+		{
+			uint32_t use_id;
+			uint32_t rows_id;
+			uint32_t columns_id;
+			uint32_t scope_id;
+		} cooperative;
+
+		struct
+		{
+			uint32_t component_type_id;
+			uint32_t component_count_id;
+		} coopVecNV;
+
+		struct
+		{
+			uint32_t type;
+			uint32_t rank;
+			uint32_t shape;
+		} tensor;
+
+		struct
+		{
+			spv::StorageClass storage;
+		} descriptor_heap_buffer;
+	} ext;
+
 	spv::StorageClass storage = spv::StorageClassGeneric;

 	SmallVector<TypeID> member_types;
@@ -659,6 +719,12 @@ struct SPIRExtension : IVariant
 		NonSemanticGeneric
 	};

+	enum ShaderDebugInfoOps
+	{
+		DebugLine = 103,
+		DebugSource = 35
+	};
+
 	explicit SPIRExtension(Extension ext_)
 	    : ext(ext_)
 	{
@@ -684,6 +750,11 @@ struct SPIREntryPoint
 	FunctionID self = 0;
 	std::string name;
 	std::string orig_name;
+	std::unordered_map<uint32_t, uint32_t> fp_fast_math_defaults;
+	bool signed_zero_inf_nan_preserve_8 = false;
+	bool signed_zero_inf_nan_preserve_16 = false;
+	bool signed_zero_inf_nan_preserve_32 = false;
+	bool signed_zero_inf_nan_preserve_64 = false;
 	SmallVector<VariableID> interface_variables;

 	Bitset flags;
@@ -743,9 +814,20 @@ struct SPIRExpression : IVariant
 	// Whether or not gl_MeshVerticesEXT[].gl_Position (as a whole or .y) is referenced
 	bool access_meshlet_position_y = false;

+	// If this expression represents a OpBufferPointerEXT cast.
+	bool buffer_pointer = false;
+
+	// Temporaries which can remain forwarded as long as this variable is not modified.
+	// Only used for buffer pointers.
+	SmallVector<ID> buffer_pointer_dependees;
+
 	// A list of expressions which this expression depends on.
 	SmallVector<ID> expression_dependencies;

+	// Similar as expression dependencies, but does not stop the tracking for force-temporary variables.
+	// We need to know the full chain from store back to any SSA variable.
+	SmallVector<ID> invariance_dependencies;
+
 	// By reading this expression, we implicitly read these expressions as well.
 	// Used by access chain Store and Load since we read multiple expressions in this case.
 	SmallVector<ID> implied_read_expressions;
@@ -912,6 +994,7 @@ struct SPIRBlock : IVariant
 	// All access to these variables are dominated by this block,
 	// so before branching anywhere we need to make sure that we declare these variables.
 	SmallVector<VariableID> dominated_variables;
+	SmallVector<bool> rearm_dominated_variables;

 	// These are variables which should be declared in a for loop header, if we
 	// fail to use a classic for-loop,
@@ -1020,6 +1103,9 @@ struct SPIRFunction : IVariant
 	// consider arrays value types.
 	SmallVector<ID> constant_arrays_needed_on_stack;

+	// Does this function (or any function called by it), emit geometry?
+	bool emits_geometry = false;
+
 	bool active = false;
 	bool flush_undeclared = true;
 	bool do_combined_parameters = true;
@@ -1103,6 +1189,9 @@ struct SPIRVariable : IVariant
 	// Temporaries which can remain forwarded as long as this variable is not modified.
 	SmallVector<ID> dependees;

+	// ShaderDebugInfo local variables attached to this variable via DebugDeclare
+	SmallVector<ID> debug_local_variables;
+
 	bool deferred_declaration = false;
 	bool phi_variable = false;

@@ -1121,6 +1210,15 @@ struct SPIRVariable : IVariant
 	// Set to true while we're inside the for loop.
 	bool loop_variable_enable = false;

+	// Used to find global LUTs
+	bool is_written_to = false;
+
+	// Untyped pointer. The pointer of the variable is effectively void.
+	// The underlying payload for allocation is in alloca_type, but may be 0 too.
+	// This is mostly here to support descriptor heap proxy.
+	bool untyped = false;
+	ID untyped_alloca_type = 0;
+
 	SPIRFunction::Parameter *parameter = nullptr;

 	SPIRV_CROSS_DECLARE_CLONE(SPIRVariable)
@@ -1217,6 +1315,26 @@ struct SPIRConstant : IVariant
 		return u.f32;
 	}

+	static inline float fe4m3_to_f32(uint8_t v)
+	{
+		if ((v & 0x7f) == 0x7f)
+		{
+			union
+			{
+				float f32;
+				uint32_t u32;
+			} u;
+
+			u.u32 = (v & 0x80) ? 0xffffffffu : 0x7fffffffu;
+			return u.f32;
+		}
+		else
+		{
+			// Reuse the FP16 to FP32 code. Cute bit-hackery.
+			return f16_to_f32((int16_t(int8_t(v)) << 7) & (0xffff ^ 0x4000)) * 256.0f;
+		}
+	}
+
 	inline uint32_t specialization_constant_id(uint32_t col, uint32_t row) const
 	{
 		return m.c[col].id[row];
@@ -1257,6 +1375,24 @@ struct SPIRConstant : IVariant
 		return f16_to_f32(scalar_u16(col, row));
 	}

+	inline float scalar_bf16(uint32_t col = 0, uint32_t row = 0) const
+	{
+		uint32_t v = scalar_u16(col, row) << 16;
+		float fp32;
+		memcpy(&fp32, &v, sizeof(float));
+		return fp32;
+	}
+
+	inline float scalar_floate4m3(uint32_t col = 0, uint32_t row = 0) const
+	{
+		return fe4m3_to_f32(scalar_u8(col, row));
+	}
+
+	inline float scalar_bf8(uint32_t col = 0, uint32_t row = 0) const
+	{
+		return f16_to_f32(uint16_t(scalar_u8(col, row) << 8));
+	}
+
 	inline float scalar_f32(uint32_t col = 0, uint32_t row = 0) const
 	{
 		return m.c[col].r[row].f32;
@@ -1327,9 +1463,10 @@ struct SPIRConstant : IVariant

 	SPIRConstant() = default;

-	SPIRConstant(TypeID constant_type_, const uint32_t *elements, uint32_t num_elements, bool specialized)
+	SPIRConstant(TypeID constant_type_, const uint32_t *elements, uint32_t num_elements, bool specialized, bool replicated_ = false)
 	    : constant_type(constant_type_)
 	    , specialization(specialized)
+	    , replicated(replicated_)
 	{
 		subconstants.reserve(num_elements);
 		for (uint32_t i = 0; i < num_elements; i++)
@@ -1401,15 +1538,25 @@ struct SPIRConstant : IVariant
 	// If true, this is a LUT, and should always be declared in the outer scope.
 	bool is_used_as_lut = false;

+	// If this is a null constant of array type with specialized length.
+	// May require special handling in initializer
+	bool is_null_array_specialized_length = false;
+
 	// For composites which are constant arrays, etc.
 	SmallVector<ConstantID> subconstants;

+	// Whether the subconstants are intended to be replicated (e.g. OpConstantCompositeReplicateEXT)
+	bool replicated = false;
+
 	// Non-Vulkan GLSL, HLSL and sometimes MSL emits defines for each specialization constant,
 	// and uses them to initialize the constant. This allows the user
 	// to still be able to specialize the value by supplying corresponding
 	// preprocessor directives before compiling the shader.
 	std::string specialization_constant_macro_name;

+	// ConstantSizeOfEXT.
+	ID size_of_type = 0;
+
 	SPIRV_CROSS_DECLARE_CLONE(SPIRConstant)
 };

@@ -1595,6 +1742,8 @@ struct AccessChainMeta
 	bool flattened_struct = false;
 	bool relaxed_precision = false;
 	bool access_meshlet_position_y = false;
+	bool chain_is_builtin = false;
+	spv::BuiltIn builtin = {};
 };

 enum ExtendedDecorations
@@ -1668,6 +1817,8 @@ enum ExtendedDecorations
 	// lack of constructors in the 'threadgroup' address space.
 	SPIRVCrossDecorationWorkgroupStruct,

+	SPIRVCrossDecorationOverlappingBinding,
+
 	SPIRVCrossDecorationCount
 };

@@ -1677,7 +1828,7 @@ struct Meta
 	{
 		std::string alias;
 		std::string qualified_alias;
-		std::string hlsl_semantic;
+		std::string user_semantic;
 		std::string user_type;
 		Bitset decoration_flags;
 		spv::BuiltIn builtin_type = spv::BuiltInMax;
@@ -1686,16 +1837,20 @@ struct Meta
 		uint32_t set = 0;
 		uint32_t binding = 0;
 		uint32_t offset = 0;
+		uint32_t offset_id = 0;
 		uint32_t xfb_buffer = 0;
 		uint32_t xfb_stride = 0;
 		uint32_t stream = 0;
 		uint32_t array_stride = 0;
+		uint32_t array_stride_id = 0;
 		uint32_t matrix_stride = 0;
 		uint32_t input_attachment = 0;
 		uint32_t spec_id = 0;
 		uint32_t index = 0;
 		spv::FPRoundingMode fp_rounding_mode = spv::FPRoundingModeMax;
+		spv::FPFastMathModeMask fp_fast_math_mode = spv::FPFastMathModeMaskNone;
 		bool builtin = false;
+		bool qualified_alias_explicit_override = false;

 		struct Extended
 		{
@@ -1749,7 +1904,8 @@ private:

 static inline bool type_is_floating_point(const SPIRType &type)
 {
-	return type.basetype == SPIRType::Half || type.basetype == SPIRType::Float || type.basetype == SPIRType::Double;
+	return type.basetype == SPIRType::Half || type.basetype == SPIRType::Float || type.basetype == SPIRType::Double ||
+	       type.basetype == SPIRType::BFloat16 || type.basetype == SPIRType::FloatE5M2 || type.basetype == SPIRType::FloatE4M3;
 }

 static inline bool type_is_integral(const SPIRType &type)
@@ -1934,4 +2090,7 @@ struct hash<SPIRV_CROSS_NAMESPACE::TypedID<type>>
 };
 } // namespace std

+#ifdef SPIRV_CROSS_SPV_HEADER_NAMESPACE_OVERRIDE
+#undef spv
+#endif
 #endif
@@ -23,7 +23,7 @@

 #include "spirv_cpp.hpp"

-using namespace spv;
+using namespace SPIRV_CROSS_SPV_HEADER_NAMESPACE;
 using namespace SPIRV_CROSS_NAMESPACE;
 using namespace std;

@@ -40,7 +40,7 @@ void CompilerCPP::emit_buffer_block(const SPIRVariable &var)
 	emit_block_struct(type);
 	auto buffer_name = to_name(type.self);

-	statement("internal::Resource<", buffer_name, type_to_array_glsl(type), "> ", instance_name, "__;");
+	statement("internal::Resource<", buffer_name, type_to_array_glsl(type, var.self), "> ", instance_name, "__;");
 	statement_no_indent("#define ", instance_name, " __res->", instance_name, "__.get()");
 	resource_registrations.push_back(
 	    join("s.register_resource(", instance_name, "__", ", ", descriptor_set, ", ", binding, ");"));
@@ -68,7 +68,7 @@ void CompilerCPP::emit_interface_block(const SPIRVariable &var)
 	else
 		buffer_name = type_to_glsl(type);

-	statement("internal::", qual, "<", buffer_name, type_to_array_glsl(type), "> ", instance_name, "__;");
+	statement("internal::", qual, "<", buffer_name, type_to_array_glsl(type, var.self), "> ", instance_name, "__;");
 	statement_no_indent("#define ", instance_name, " __res->", instance_name, "__.get()");
 	resource_registrations.push_back(join("s.register_", lowerqual, "(", instance_name, "__", ", ", location, ");"));
 	statement("");
@@ -100,14 +100,14 @@ void CompilerCPP::emit_uniform(const SPIRVariable &var)
 	if (type.basetype == SPIRType::Image || type.basetype == SPIRType::SampledImage ||
 	    type.basetype == SPIRType::AtomicCounter)
 	{
-		statement("internal::Resource<", type_name, type_to_array_glsl(type), "> ", instance_name, "__;");
+		statement("internal::Resource<", type_name, type_to_array_glsl(type, var.self), "> ", instance_name, "__;");
 		statement_no_indent("#define ", instance_name, " __res->", instance_name, "__.get()");
 		resource_registrations.push_back(
 		    join("s.register_resource(", instance_name, "__", ", ", descriptor_set, ", ", binding, ");"));
 	}
 	else
 	{
-		statement("internal::UniformConstant<", type_name, type_to_array_glsl(type), "> ", instance_name, "__;");
+		statement("internal::UniformConstant<", type_name, type_to_array_glsl(type, var.self), "> ", instance_name, "__;");
 		statement_no_indent("#define ", instance_name, " __res->", instance_name, "__.get()");
 		resource_registrations.push_back(
 		    join("s.register_uniform_constant(", instance_name, "__", ", ", location, ");"));
@@ -130,7 +130,7 @@ void CompilerCPP::emit_push_constant_block(const SPIRVariable &var)
 	auto buffer_name = to_name(type.self);
 	auto instance_name = to_name(var.self);

-	statement("internal::PushConstant<", buffer_name, type_to_array_glsl(type), "> ", instance_name, ";");
+	statement("internal::PushConstant<", buffer_name, type_to_array_glsl(type, var.self), "> ", instance_name, ";");
 	statement_no_indent("#define ", instance_name, " __res->", instance_name, ".get()");
 	resource_registrations.push_back(join("s.register_push_constant(", instance_name, "__", ");"));
 	statement("");
@@ -27,12 +27,20 @@
 #ifndef SPV_ENABLE_UTILITY_CODE
 #define SPV_ENABLE_UTILITY_CODE
 #endif
+
+// Pragmatic hack to avoid symbol conflicts when including both hpp11 and hpp headers in same translation unit.
+// This is an unfortunate SPIRV-Headers issue that we cannot easily deal with ourselves.
+#ifdef SPIRV_CROSS_SPV_HEADER_NAMESPACE_OVERRIDE
+#define spv SPIRV_CROSS_SPV_HEADER_NAMESPACE_OVERRIDE
+#endif
+
 #include "spirv.hpp"
 #include "spirv_cfg.hpp"
 #include "spirv_cross_parsed_ir.hpp"

 namespace SPIRV_CROSS_NAMESPACE
 {
+using namespace SPIRV_CROSS_SPV_HEADER_NAMESPACE;
 struct Resource
 {
 	// Resources are identified with their SPIR-V ID.
@@ -69,7 +77,7 @@ struct BuiltInResource
 	// A builtin present here does not necessarily mean it's considered an active builtin,
 	// since variable ID "activeness" is only tracked on OpVariable level, not Block members.
 	// For that, update_active_builtins() -> has_active_builtin() can be used to further refine the reflection.
-	spv::BuiltIn builtin;
+	BuiltIn builtin;

 	// This is the actual value type of the builtin.
 	// Typically float4, float, array<float, N> for the gl_PerVertex builtins.
@@ -83,6 +91,28 @@ struct BuiltInResource
 	Resource resource;
 };

+// Needs to stay in sync 1:1 with C API.
+enum ResourceType
+{
+	ResourceTypeUnknown = 0,
+	ResourceTypeUniformBuffer = 1,
+	ResourceTypeStorageBuffer = 2,
+	ResourceTypeStageInput = 3,
+	ResourceTypeStageOutput = 4,
+	ResourceTypeSubpassInput = 5,
+	ResourceTypeStorageImage = 6,
+	ResourceTypeSampledImage = 7,
+	ResourceTypeAtomicCounter = 8,
+	ResourceTypePushConstant = 9,
+	ResourceTypeSeparateImage = 10,
+	ResourceTypeSeparateSamplers = 11,
+	ResourceTypeAccelerationStructure = 12,
+	ResourceTypeRayQuery = 13,
+	ResourceTypeShaderRecordBuffer = 14,
+	ResourceTypeGLPlainUniform = 15,
+	ResourceTypeTensor = 16
+};
+
 struct ShaderResources
 {
 	SmallVector<Resource> uniform_buffers;
@@ -95,6 +125,7 @@ struct ShaderResources
 	SmallVector<Resource> atomic_counters;
 	SmallVector<Resource> acceleration_structures;
 	SmallVector<Resource> gl_plain_uniforms;
+	SmallVector<Resource> tensors;

 	// There can only be one push constant block,
 	// but keep the vector in case this restriction is lifted in the future.
@@ -151,7 +182,7 @@ enum BufferPackingStandard
 struct EntryPoint
 {
 	std::string name;
-	spv::ExecutionModel execution_model;
+	ExecutionModel execution_model;
 };

 class Compiler
@@ -182,8 +213,8 @@ public:
 	const std::string &get_name(ID id) const;

 	// Applies a decoration to an ID. Effectively injects OpDecorate.
-	void set_decoration(ID id, spv::Decoration decoration, uint32_t argument = 0);
-	void set_decoration_string(ID id, spv::Decoration decoration, const std::string &argument);
+	void set_decoration(ID id, Decoration decoration, uint32_t argument = 0);
+	void set_decoration_string(ID id, Decoration decoration, const std::string &argument);

 	// Overrides the identifier OpName of an ID.
 	// Identifiers beginning with underscores or identifiers which contain double underscores
@@ -191,22 +222,22 @@ public:
 	void set_name(ID id, const std::string &name);

 	// Gets a bitmask for the decorations which are applied to ID.
-	// I.e. (1ull << spv::DecorationFoo) | (1ull << spv::DecorationBar)
+	// I.e. (1ull << DecorationFoo) | (1ull << DecorationBar)
 	const Bitset &get_decoration_bitset(ID id) const;

 	// Returns whether the decoration has been applied to the ID.
-	bool has_decoration(ID id, spv::Decoration decoration) const;
+	bool has_decoration(ID id, Decoration decoration) const;

 	// Gets the value for decorations which take arguments.
-	// If the decoration is a boolean (i.e. spv::DecorationNonWritable),
+	// If the decoration is a boolean (i.e. DecorationNonWritable),
 	// 1 will be returned.
 	// If decoration doesn't exist or decoration is not recognized,
 	// 0 will be returned.
-	uint32_t get_decoration(ID id, spv::Decoration decoration) const;
-	const std::string &get_decoration_string(ID id, spv::Decoration decoration) const;
+	uint32_t get_decoration(ID id, Decoration decoration) const;
+	const std::string &get_decoration_string(ID id, Decoration decoration) const;

 	// Removes the decoration for an ID.
-	void unset_decoration(ID id, spv::Decoration decoration);
+	void unset_decoration(ID id, Decoration decoration);

 	// Gets the SPIR-V type associated with ID.
 	// Mostly used with Resource::type_id and Resource::base_type_id to parse the underlying type of a resource.
@@ -216,7 +247,7 @@ public:
 	const SPIRType &get_type_from_variable(VariableID id) const;

 	// Gets the underlying storage class for an OpVariable.
-	spv::StorageClass get_storage_class(VariableID id) const;
+	StorageClass get_storage_class(VariableID id) const;

 	// If get_name() is an empty string, get the fallback name which will be used
 	// instead in the disassembled source.
@@ -231,8 +262,8 @@ public:
 	const std::string &get_member_name(TypeID id, uint32_t index) const;

 	// Given an OpTypeStruct in ID, obtain the OpMemberDecoration for member number "index".
-	uint32_t get_member_decoration(TypeID id, uint32_t index, spv::Decoration decoration) const;
-	const std::string &get_member_decoration_string(TypeID id, uint32_t index, spv::Decoration decoration) const;
+	uint32_t get_member_decoration(TypeID id, uint32_t index, Decoration decoration) const;
+	const std::string &get_member_decoration_string(TypeID id, uint32_t index, Decoration decoration) const;

 	// Sets the member identifier for OpTypeStruct ID, member number "index".
 	void set_member_name(TypeID id, uint32_t index, const std::string &name);
@@ -245,15 +276,15 @@ public:
 	const Bitset &get_member_decoration_bitset(TypeID id, uint32_t index) const;

 	// Returns whether the decoration has been applied to a member of a struct.
-	bool has_member_decoration(TypeID id, uint32_t index, spv::Decoration decoration) const;
+	bool has_member_decoration(TypeID id, uint32_t index, Decoration decoration) const;

 	// Similar to set_decoration, but for struct members.
-	void set_member_decoration(TypeID id, uint32_t index, spv::Decoration decoration, uint32_t argument = 0);
-	void set_member_decoration_string(TypeID id, uint32_t index, spv::Decoration decoration,
+	void set_member_decoration(TypeID id, uint32_t index, Decoration decoration, uint32_t argument = 0);
+	void set_member_decoration_string(TypeID id, uint32_t index, Decoration decoration,
 	                                  const std::string &argument);

 	// Unsets a member decoration, similar to unset_decoration.
-	void unset_member_decoration(TypeID id, uint32_t index, spv::Decoration decoration);
+	void unset_member_decoration(TypeID id, uint32_t index, Decoration decoration);

 	// Gets the fallback name for a member, similar to get_fallback_name.
 	virtual const std::string get_fallback_member_name(uint32_t index) const
@@ -339,28 +370,28 @@ public:
 	// Names for entry points in the SPIR-V module may alias if they belong to different execution models.
 	// To disambiguate, we must pass along with the entry point names the execution model.
 	SmallVector<EntryPoint> get_entry_points_and_stages() const;
-	void set_entry_point(const std::string &entry, spv::ExecutionModel execution_model);
+	void set_entry_point(const std::string &entry, ExecutionModel execution_model);

 	// Renames an entry point from old_name to new_name.
 	// If old_name is currently selected as the current entry point, it will continue to be the current entry point,
 	// albeit with a new name.
 	// get_entry_points() is essentially invalidated at this point.
 	void rename_entry_point(const std::string &old_name, const std::string &new_name,
-	                        spv::ExecutionModel execution_model);
-	const SPIREntryPoint &get_entry_point(const std::string &name, spv::ExecutionModel execution_model) const;
-	SPIREntryPoint &get_entry_point(const std::string &name, spv::ExecutionModel execution_model);
+	                        ExecutionModel execution_model);
+	const SPIREntryPoint &get_entry_point(const std::string &name, ExecutionModel execution_model) const;
+	SPIREntryPoint &get_entry_point(const std::string &name, ExecutionModel execution_model);
 	const std::string &get_cleansed_entry_point_name(const std::string &name,
-	                                                 spv::ExecutionModel execution_model) const;
+	                                                 ExecutionModel execution_model) const;

 	// Traverses all reachable opcodes and sets active_builtins to a bitmask of all builtin variables which are accessed in the shader.
 	void update_active_builtins();
-	bool has_active_builtin(spv::BuiltIn builtin, spv::StorageClass storage) const;
+	bool has_active_builtin(BuiltIn builtin, StorageClass storage) const;

 	// Query and modify OpExecutionMode.
 	const Bitset &get_execution_mode_bitset() const;

-	void unset_execution_mode(spv::ExecutionMode mode);
-	void set_execution_mode(spv::ExecutionMode mode, uint32_t arg0 = 0, uint32_t arg1 = 0, uint32_t arg2 = 0);
+	void unset_execution_mode(ExecutionMode mode);
+	void set_execution_mode(ExecutionMode mode, uint32_t arg0 = 0, uint32_t arg1 = 0, uint32_t arg2 = 0);

 	// Gets argument for an execution mode (LocalSize, Invocations, OutputVertices).
 	// For LocalSize or LocalSizeId, the index argument is used to select the dimension (X = 0, Y = 1, Z = 2).
@@ -368,8 +399,8 @@ public:
 	// LocalSizeId query returns an ID. If LocalSizeId execution mode is not used, it returns 0.
 	// LocalSize always returns a literal. If execution mode is LocalSizeId,
 	// the literal (spec constant or not) is still returned.
-	uint32_t get_execution_mode_argument(spv::ExecutionMode mode, uint32_t index = 0) const;
-	spv::ExecutionModel get_execution_model() const;
+	uint32_t get_execution_mode_argument(ExecutionMode mode, uint32_t index = 0) const;
+	ExecutionModel get_execution_model() const;

 	bool is_tessellation_shader() const;
 	bool is_tessellating_triangles() const;
@@ -482,7 +513,7 @@ public:
 	// If the decoration was declared, sets the word_offset to an offset into the provided SPIR-V binary buffer and returns true,
 	// otherwise, returns false.
 	// If the decoration does not have any value attached to it (e.g. DecorationRelaxedPrecision), this function will also return false.
-	bool get_binary_offset_for_decoration(VariableID id, spv::Decoration decoration, uint32_t &word_offset) const;
+	bool get_binary_offset_for_decoration(VariableID id, Decoration decoration, uint32_t &word_offset) const;

 	// HLSL counter buffer reflection interface.
 	// Append/Consume/Increment/Decrement in HLSL is implemented as two "neighbor" buffer objects where
@@ -508,7 +539,7 @@ public:
 	bool buffer_get_hlsl_counter_buffer(VariableID id, uint32_t &counter_id) const;

 	// Gets the list of all SPIR-V Capabilities which were declared in the SPIR-V module.
-	const SmallVector<spv::Capability> &get_declared_capabilities() const;
+	const SmallVector<Capability> &get_declared_capabilities() const;

 	// Gets the list of all SPIR-V extensions which were declared in the SPIR-V module.
 	const SmallVector<std::string> &get_declared_extensions() const;
@@ -540,6 +571,9 @@ public:
 		return position_invariant;
 	}

+	const ParsedIR &get_ir() const { return ir; }
+	uint32_t evaluate_constant_u32(uint32_t id) const;
+
 protected:
 	const uint32_t *stream(const Instruction &instr) const
 	{
@@ -573,6 +607,7 @@ protected:
 	// (SSBO, image load store, etc)
 	SmallVector<uint32_t> global_variables;
 	SmallVector<uint32_t> aliased_variables;
+	SmallVector<uint32_t> buffer_pointer_variables;

 	SPIRFunction *current_function = nullptr;
 	SPIRBlock *current_block = nullptr;
@@ -671,26 +706,29 @@ protected:

 	const SPIREntryPoint &get_entry_point() const;
 	SPIREntryPoint &get_entry_point();
-	static bool is_tessellation_shader(spv::ExecutionModel model);
+	static bool is_tessellation_shader(ExecutionModel model);

 	virtual std::string to_name(uint32_t id, bool allow_alias = true) const;
 	bool is_builtin_variable(const SPIRVariable &var) const;
 	bool is_builtin_type(const SPIRType &type) const;
 	bool is_hidden_variable(const SPIRVariable &var, bool include_builtins = false) const;
 	bool is_immutable(uint32_t id) const;
-	bool is_member_builtin(const SPIRType &type, uint32_t index, spv::BuiltIn *builtin) const;
+	bool is_member_builtin(const SPIRType &type, uint32_t index, BuiltIn *builtin) const;
 	bool is_scalar(const SPIRType &type) const;
 	bool is_vector(const SPIRType &type) const;
 	bool is_matrix(const SPIRType &type) const;
 	bool is_array(const SPIRType &type) const;
 	bool is_pointer(const SPIRType &type) const;
 	bool is_physical_pointer(const SPIRType &type) const;
+	bool is_physical_or_buffer_pointer(const SPIRType &type) const;
+	bool is_physical_pointer_to_buffer_block(const SPIRType &type) const;
 	static bool is_runtime_size_array(const SPIRType &type);
 	uint32_t expression_type_id(uint32_t id) const;
 	const SPIRType &expression_type(uint32_t id) const;
 	bool expression_is_lvalue(uint32_t id) const;
 	bool variable_storage_is_aliased(const SPIRVariable &var);
 	SPIRVariable *maybe_get_backing_variable(uint32_t chain);
+	SPIRExpression *maybe_get_backing_buffer_pointer(uint32_t chain);

 	void register_read(uint32_t expr, uint32_t chain, bool forwarded);
 	void register_write(uint32_t chain);
@@ -725,6 +763,7 @@ protected:

 	// Dependency tracking for temporaries read from variables.
 	void flush_dependees(SPIRVariable &var);
+	void flush_dependees(SPIRExpression &expr);
 	void flush_all_active_variables();
 	void flush_control_dependent_expressions(uint32_t block);
 	void flush_all_atomic_capable_variables();
@@ -743,6 +782,8 @@ protected:

 	bool function_is_pure(const SPIRFunction &func);
 	bool block_is_pure(const SPIRBlock &block);
+	bool function_is_control_dependent(const SPIRFunction &func);
+	bool block_is_control_dependent(const SPIRBlock &block);

 	bool execution_is_branchless(const SPIRBlock &from, const SPIRBlock &to) const;
 	bool execution_is_direct_branch(const SPIRBlock &from, const SPIRBlock &to) const;
@@ -784,11 +825,12 @@ protected:
 	// Used internally to implement various traversals for queries.
 	struct OpcodeHandler
 	{
+		explicit OpcodeHandler(Compiler &compiler_) : compiler(compiler_) {}
 		virtual ~OpcodeHandler() = default;

 		// Return true if traversal should continue.
 		// If false, traversal will end immediately.
-		virtual bool handle(spv::Op opcode, const uint32_t *args, uint32_t length) = 0;
+		virtual bool handle(Op opcode, const uint32_t *args, uint32_t length) = 0;
 		virtual bool handle_terminator(const SPIRBlock &)
 		{
 			return true;
@@ -819,20 +861,40 @@ protected:
 		{
 			return true;
 		}
+
+		Compiler &compiler;
+		std::unordered_map<uint32_t, uint32_t> result_types;
+		const SPIRType *get_expression_result_type(uint32_t id) const;
+		bool enable_result_types = false;
+
+		template <typename T> T &get(uint32_t id)
+		{
+			return compiler.get<T>(id);
+		}
+
+		template <typename T> const T &get(uint32_t id) const
+		{
+			return compiler.get<T>(id);
+		}
+
+		template <typename T, typename... P>
+		T &set(uint32_t id, P &&... args)
+		{
+			return compiler.set<T>(id, std::forward<P>(args)...);
+		}
 	};

 	struct BufferAccessHandler : OpcodeHandler
 	{
 		BufferAccessHandler(const Compiler &compiler_, SmallVector<BufferRange> &ranges_, uint32_t id_)
-		    : compiler(compiler_)
+		    : OpcodeHandler(const_cast<Compiler &>(compiler_))
 		    , ranges(ranges_)
 		    , id(id_)
 		{
 		}

-		bool handle(spv::Op opcode, const uint32_t *args, uint32_t length) override;
+		bool handle(Op opcode, const uint32_t *args, uint32_t length) override;

-		const Compiler &compiler;
 		SmallVector<BufferRange> &ranges;
 		uint32_t id;

@@ -842,29 +904,26 @@ protected:
 	struct InterfaceVariableAccessHandler : OpcodeHandler
 	{
 		InterfaceVariableAccessHandler(const Compiler &compiler_, std::unordered_set<VariableID> &variables_)
-		    : compiler(compiler_)
+		    : OpcodeHandler(const_cast<Compiler &>(compiler_))
 		    , variables(variables_)
 		{
 		}

-		bool handle(spv::Op opcode, const uint32_t *args, uint32_t length) override;
+		bool handle(Op opcode, const uint32_t *args, uint32_t length) override;

-		const Compiler &compiler;
 		std::unordered_set<VariableID> &variables;
 	};

 	struct CombinedImageSamplerHandler : OpcodeHandler
 	{
-		CombinedImageSamplerHandler(Compiler &compiler_)
-		    : compiler(compiler_)
+		explicit CombinedImageSamplerHandler(Compiler &compiler_)
+		    : OpcodeHandler(compiler_)
 		{
 		}
-		bool handle(spv::Op opcode, const uint32_t *args, uint32_t length) override;
+		bool handle(Op opcode, const uint32_t *args, uint32_t length) override;
 		bool begin_function_scope(const uint32_t *args, uint32_t length) override;
 		bool end_function_scope(const uint32_t *args, uint32_t length) override;

-		Compiler &compiler;
-
 		// Each function in the call stack needs its own remapping for parameters so we can deduce which global variable each texture/sampler the parameter is statically bound to.
 		std::stack<std::unordered_map<uint32_t, uint32_t>> parameter_remapping;
 		std::stack<SPIRFunction *> functions;
@@ -878,27 +937,24 @@ protected:

 	struct DummySamplerForCombinedImageHandler : OpcodeHandler
 	{
-		DummySamplerForCombinedImageHandler(Compiler &compiler_)
-		    : compiler(compiler_)
+		explicit DummySamplerForCombinedImageHandler(Compiler &compiler_)
+		    : OpcodeHandler(compiler_)
 		{
 		}
-		bool handle(spv::Op opcode, const uint32_t *args, uint32_t length) override;
-
-		Compiler &compiler;
+		bool handle(Op opcode, const uint32_t *args, uint32_t length) override;
 		bool need_dummy_sampler = false;
 	};

 	struct ActiveBuiltinHandler : OpcodeHandler
 	{
-		ActiveBuiltinHandler(Compiler &compiler_)
-		    : compiler(compiler_)
+		explicit ActiveBuiltinHandler(Compiler &compiler_)
+		    : OpcodeHandler(compiler_)
 		{
 		}

-		bool handle(spv::Op opcode, const uint32_t *args, uint32_t length) override;
-		Compiler &compiler;
+		bool handle(Op opcode, const uint32_t *args, uint32_t length) override;

-		void handle_builtin(const SPIRType &type, spv::BuiltIn builtin, const Bitset &decoration_flags);
+		void handle_builtin(const SPIRType &type, BuiltIn builtin, const Bitset &decoration_flags);
 		void add_if_builtin(uint32_t id);
 		void add_if_builtin_or_block(uint32_t id);
 		void add_if_builtin(uint32_t id, bool allow_blocks);
@@ -950,13 +1006,12 @@ protected:

 	struct CombinedImageSamplerDrefHandler : OpcodeHandler
 	{
-		CombinedImageSamplerDrefHandler(Compiler &compiler_)
-		    : compiler(compiler_)
+		explicit CombinedImageSamplerDrefHandler(Compiler &compiler_)
+		    : OpcodeHandler(compiler_)
 		{
 		}
-		bool handle(spv::Op opcode, const uint32_t *args, uint32_t length) override;
+		bool handle(Op opcode, const uint32_t *args, uint32_t length) override;

-		Compiler &compiler;
 		std::unordered_set<uint32_t> dref_combined_samplers;
 	};

@@ -964,14 +1019,13 @@ protected:
 	{
 		CombinedImageSamplerUsageHandler(Compiler &compiler_,
 		                                 const std::unordered_set<uint32_t> &dref_combined_samplers_)
-		    : compiler(compiler_)
+		    : OpcodeHandler(compiler_)
 		    , dref_combined_samplers(dref_combined_samplers_)
 		{
 		}

 		bool begin_function_scope(const uint32_t *args, uint32_t length) override;
-		bool handle(spv::Op opcode, const uint32_t *args, uint32_t length) override;
-		Compiler &compiler;
+		bool handle(Op opcode, const uint32_t *args, uint32_t length) override;
 		const std::unordered_set<uint32_t> &dref_combined_samplers;

 		std::unordered_map<uint32_t, std::unordered_set<uint32_t>> dependency_hierarchy;
@@ -993,8 +1047,7 @@ protected:
 		explicit CFGBuilder(Compiler &compiler_);

 		bool follow_function_call(const SPIRFunction &func) override;
-		bool handle(spv::Op op, const uint32_t *args, uint32_t length) override;
-		Compiler &compiler;
+		bool handle(Op op, const uint32_t *args, uint32_t length) override;
 		std::unordered_map<uint32_t, std::unique_ptr<CFG>> function_cfgs;
 	};

@@ -1008,10 +1061,9 @@ protected:
 		void notify_variable_access(uint32_t id, uint32_t block);
 		bool id_is_phi_variable(uint32_t id) const;
 		bool id_is_potential_temporary(uint32_t id) const;
-		bool handle(spv::Op op, const uint32_t *args, uint32_t length) override;
+		bool handle(Op op, const uint32_t *args, uint32_t length) override;
 		bool handle_terminator(const SPIRBlock &block) override;

-		Compiler &compiler;
 		SPIRFunction &entry;
 		std::unordered_map<uint32_t, std::unordered_set<uint32_t>> accessed_variables_to_block;
 		std::unordered_map<uint32_t, std::unordered_set<uint32_t>> accessed_temporaries_to_block;
@@ -1029,9 +1081,8 @@ protected:
 	{
 		StaticExpressionAccessHandler(Compiler &compiler_, uint32_t variable_id_);
 		bool follow_function_call(const SPIRFunction &) override;
-		bool handle(spv::Op op, const uint32_t *args, uint32_t length) override;
+		bool handle(Op op, const uint32_t *args, uint32_t length) override;

-		Compiler &compiler;
 		uint32_t variable_id;
 		uint32_t static_expression = 0;
 		uint32_t write_count = 0;
@@ -1045,12 +1096,12 @@ protected:
 	struct PhysicalStorageBufferPointerHandler : OpcodeHandler
 	{
 		explicit PhysicalStorageBufferPointerHandler(Compiler &compiler_);
-		bool handle(spv::Op op, const uint32_t *args, uint32_t length) override;
-		Compiler &compiler;
+		bool handle(Op op, const uint32_t *args, uint32_t length) override;

 		std::unordered_set<uint32_t> non_block_types;
 		std::unordered_map<uint32_t, PhysicalBlockMeta> physical_block_type_meta;
 		std::unordered_map<uint32_t, PhysicalBlockMeta *> access_chain_to_physical_block;
+		std::unordered_set<uint32_t> analyzed_type_ids;

 		void mark_aligned_access(uint32_t id, const uint32_t *args, uint32_t length);
 		PhysicalBlockMeta *find_block_meta(uint32_t id) const;
@@ -1064,11 +1115,43 @@ protected:
 	SmallVector<uint32_t> physical_storage_non_block_pointer_types;
 	std::unordered_map<uint32_t, PhysicalBlockMeta> physical_storage_type_to_alignment;

+	struct DescriptorHeapMeta
+	{
+		TypeID type;
+		bool hlsl_style_stride;
+
+		// For buffers
+		ID buffer_pointer_id;
+		StorageClass storage;
+		bool nonwritable;
+		bool nonreadable;
+		bool coherent;
+		bool is_volatile;
+		bool is_restrict;
+	};
+	std::vector<DescriptorHeapMeta> descriptor_heap_types;
+	void analyze_descriptor_heap_types();
+
 	void analyze_variable_scope(SPIRFunction &function, AnalyzeVariableScopeAccessHandler &handler);
 	void find_function_local_luts(SPIRFunction &function, const AnalyzeVariableScopeAccessHandler &handler,
 	                              bool single_function);
 	bool may_read_undefined_variable_in_block(const SPIRBlock &block, uint32_t var);

+	struct GeometryEmitDisocveryHandler : OpcodeHandler
+	{
+		explicit GeometryEmitDisocveryHandler(Compiler &compiler_)
+		    : OpcodeHandler(compiler_)
+		{
+		}
+
+		bool handle(Op opcode, const uint32_t *args, uint32_t length) override;
+		bool begin_function_scope(const uint32_t *, uint32_t) override;
+		bool end_function_scope(const uint32_t *, uint32_t) override;
+		SmallVector<SPIRFunction *> function_stack;
+	};
+
+	void discover_geometry_emitters();
+
 	// Finds all resources that are written to from inside the critical section, if present.
 	// The critical section is delimited by OpBeginInvocationInterlockEXT and
 	// OpEndInvocationInterlockEXT instructions. In MSL and HLSL, any resources written
@@ -1076,16 +1159,15 @@ protected:
 	struct InterlockedResourceAccessHandler : OpcodeHandler
 	{
 		InterlockedResourceAccessHandler(Compiler &compiler_, uint32_t entry_point_id)
-		    : compiler(compiler_)
+		    : OpcodeHandler(compiler_)
 		{
 			call_stack.push_back(entry_point_id);
 		}

-		bool handle(spv::Op op, const uint32_t *args, uint32_t length) override;
+		bool handle(Op op, const uint32_t *args, uint32_t length) override;
 		bool begin_function_scope(const uint32_t *args, uint32_t length) override;
 		bool end_function_scope(const uint32_t *args, uint32_t length) override;

-		Compiler &compiler;
 		bool in_crit_sec = false;

 		uint32_t interlock_function_id = 0;
@@ -1101,17 +1183,16 @@ protected:
 	struct InterlockedResourceAccessPrepassHandler : OpcodeHandler
 	{
 		InterlockedResourceAccessPrepassHandler(Compiler &compiler_, uint32_t entry_point_id)
-		    : compiler(compiler_)
+		    : OpcodeHandler(compiler_)
 		{
 			call_stack.push_back(entry_point_id);
 		}

 		void rearm_current_block(const SPIRBlock &block) override;
-		bool handle(spv::Op op, const uint32_t *args, uint32_t length) override;
+		bool handle(Op op, const uint32_t *args, uint32_t length) override;
 		bool begin_function_scope(const uint32_t *args, uint32_t length) override;
 		bool end_function_scope(const uint32_t *args, uint32_t length) override;

-		Compiler &compiler;
 		uint32_t interlock_function_id = 0;
 		uint32_t current_block_id = 0;
 		bool split_function_case = false;
@@ -1128,11 +1209,11 @@ protected:

 	std::unordered_map<uint32_t, std::string> declared_block_names;

-	bool instruction_to_result_type(uint32_t &result_type, uint32_t &result_id, spv::Op op, const uint32_t *args,
-	                                uint32_t length);
+	static bool instruction_to_result_type(
+			uint32_t &result_type, uint32_t &result_id, Op op, const uint32_t *args, uint32_t length);

 	Bitset combined_decoration_for_member(const SPIRType &type, uint32_t index) const;
-	static bool is_desktop_only_format(spv::ImageFormat format);
+	static bool is_desktop_only_format(ImageFormat format);

 	bool is_depth_image(const SPIRType &type, uint32_t id) const;

@@ -1151,6 +1232,7 @@ protected:
 	bool type_contains_recursion(const SPIRType &type);
 	bool type_is_array_of_pointers(const SPIRType &type) const;
 	bool type_is_block_like(const SPIRType &type) const;
+	bool type_is_explicit_layout(const SPIRType &type) const;
 	bool type_is_top_level_block(const SPIRType &type) const;
 	bool type_is_opaque_value(const SPIRType &type) const;

@@ -1160,7 +1242,6 @@ protected:
 	bool flush_phi_required(BlockID from, BlockID to) const;

 	uint32_t evaluate_spec_constant_u32(const SPIRConstantOp &spec) const;
-	uint32_t evaluate_constant_u32(uint32_t id) const;

 	bool is_vertex_like_shader() const;

@@ -1176,4 +1257,8 @@ private:
 };
 } // namespace SPIRV_CROSS_NAMESPACE

+#ifdef SPIRV_CROSS_SPV_HEADER_NAMESPACE_OVERRIDE
+#undef spv
+#endif
+
 #endif
@@ -576,6 +576,7 @@ public:
 			if (!ptr)
 				return nullptr;

+			vacants.reserve(num_objects);
 			for (unsigned i = 0; i < num_objects; i++)
 				vacants.push_back(&ptr[i]);

@@ -66,6 +66,11 @@ public:
 	    : std::runtime_error(str)
 	{
 	}
+
+	explicit CompilerError(const char *str)
+	    : std::runtime_error(str)
+	{
+	}
 };

 #define SPIRV_CROSS_THROW(x) throw CompilerError(x)
@@ -26,7 +26,7 @@
 #include <assert.h>

 using namespace std;
-using namespace spv;
+using namespace SPIRV_CROSS_SPV_HEADER_NAMESPACE;

 namespace SPIRV_CROSS_NAMESPACE
 {
@@ -49,6 +49,7 @@ ParsedIR::ParsedIR()
 	pool_group->pools[TypeAccessChain].reset(new ObjectPool<SPIRAccessChain>);
 	pool_group->pools[TypeUndef].reset(new ObjectPool<SPIRUndef>);
 	pool_group->pools[TypeString].reset(new ObjectPool<SPIRString>);
+	pool_group->pools[TypeDebugLocalVariable].reset(new ObjectPool<SPIRDebugLocalVariable>);
 }

 // Should have been default-implemented, but need this on MSVC 2013.
@@ -78,7 +79,10 @@ ParsedIR &ParsedIR::operator=(ParsedIR &&other) SPIRV_CROSS_NOEXCEPT
 		memory_model = other.memory_model;

 		default_entry_point = other.default_entry_point;
-		source = other.source;
+		is_library_module = other.is_library_module;
+		library_exports = std::move(other.library_exports);
+		library_exported_functions = std::move(other.library_exported_functions);
+		sources = std::move(other.sources);
 		loop_iteration_depth_hard = other.loop_iteration_depth_hard;
 		loop_iteration_depth_soft = other.loop_iteration_depth_soft;

@@ -110,7 +114,10 @@ ParsedIR &ParsedIR::operator=(const ParsedIR &other)
 		continue_block_to_loop_header = other.continue_block_to_loop_header;
 		entry_points = other.entry_points;
 		default_entry_point = other.default_entry_point;
-		source = other.source;
+		is_library_module = other.is_library_module;
+		library_exports = other.library_exports;
+		library_exported_functions = other.library_exported_functions;
+		sources = other.sources;
 		loop_iteration_depth_hard = other.loop_iteration_depth_hard;
 		loop_iteration_depth_soft = other.loop_iteration_depth_soft;
 		addressing_model = other.addressing_model;
@@ -366,8 +373,8 @@ void ParsedIR::set_decoration_string(ID id, Decoration decoration, const string

 	switch (decoration)
 	{
-	case DecorationHlslSemanticGOOGLE:
-		dec.hlsl_semantic = argument;
+	case DecorationUserSemantic:
+		dec.user_semantic = argument;
 		break;

 	case DecorationUserTypeGOOGLE:
@@ -403,6 +410,10 @@ void ParsedIR::set_decoration(ID id, Decoration decoration, uint32_t argument)
 		dec.offset = argument;
 		break;

+	case DecorationOffsetIdEXT:
+		dec.offset_id = argument;
+		break;
+
 	case DecorationXfbBuffer:
 		dec.xfb_buffer = argument;
 		break;
@@ -419,6 +430,10 @@ void ParsedIR::set_decoration(ID id, Decoration decoration, uint32_t argument)
 		dec.array_stride = argument;
 		break;

+	case DecorationArrayStrideIdEXT:
+		dec.array_stride_id = argument;
+		break;
+
 	case DecorationMatrixStride:
 		dec.matrix_stride = argument;
 		break;
@@ -452,6 +467,10 @@ void ParsedIR::set_decoration(ID id, Decoration decoration, uint32_t argument)
 		dec.fp_rounding_mode = static_cast<FPRoundingMode>(argument);
 		break;

+	case DecorationFPFastMathMode:
+		dec.fp_fast_math_mode = static_cast<FPFastMathModeMask>(argument);
+		break;
+
 	default:
 		break;
 	}
@@ -487,6 +506,10 @@ void ParsedIR::set_member_decoration(TypeID id, uint32_t index, Decoration decor
 		dec.offset = argument;
 		break;

+	case DecorationOffsetIdEXT:
+		dec.offset_id = argument;
+		break;
+
 	case DecorationXfbBuffer:
 		dec.xfb_buffer = argument;
 		break;
@@ -523,8 +546,27 @@ void ParsedIR::mark_used_as_array_length(ID id)
 	switch (ids[id].get_type())
 	{
 	case TypeConstant:
-		get<SPIRConstant>(id).is_used_as_array_length = true;
+	{
+		auto &c = get<SPIRConstant>(id);
+		c.is_used_as_array_length = true;
+
+		// Mark composite dependencies as well.
+		for (auto &sub_id: c.m.id)
+			if (sub_id)
+				mark_used_as_array_length(sub_id);
+
+		for (uint32_t col = 0; col < c.m.columns; col++)
+		{
+			for (auto &sub_id : c.m.c[col].id)
+				if (sub_id)
+					mark_used_as_array_length(sub_id);
+		}
+
+		for (auto &sub_id : c.subconstants)
+			if (sub_id)
+				mark_used_as_array_length(sub_id);
 		break;
+	}

 	case TypeConstantOp:
 	{
@@ -564,7 +606,8 @@ Bitset ParsedIR::get_buffer_block_type_flags(const SPIRType &type) const
 Bitset ParsedIR::get_buffer_block_flags(const SPIRVariable &var) const
 {
 	auto &type = get<SPIRType>(var.basetype);
-	assert(type.basetype == SPIRType::Struct);
+	if (type.basetype != SPIRType::Struct)
+		SPIRV_CROSS_THROW("Cannot get buffer block flags for non-buffer variable.");

 	// Some flags like non-writable, non-readable are actually found
 	// as member decorations. If all members have a decoration set, propagate
@@ -620,6 +663,8 @@ uint32_t ParsedIR::get_decoration(ID id, Decoration decoration) const
 		return dec.component;
 	case DecorationOffset:
 		return dec.offset;
+	case DecorationOffsetIdEXT:
+		return dec.offset_id;
 	case DecorationXfbBuffer:
 		return dec.xfb_buffer;
 	case DecorationXfbStride:
@@ -636,12 +681,16 @@ uint32_t ParsedIR::get_decoration(ID id, Decoration decoration) const
 		return dec.spec_id;
 	case DecorationArrayStride:
 		return dec.array_stride;
+	case DecorationArrayStrideIdEXT:
+		return dec.array_stride_id;
 	case DecorationMatrixStride:
 		return dec.matrix_stride;
 	case DecorationIndex:
 		return dec.index;
 	case DecorationFPRoundingMode:
 		return dec.fp_rounding_mode;
+	case DecorationFPFastMathMode:
+		return dec.fp_fast_math_mode;
 	default:
 		return 1;
 	}
@@ -660,8 +709,8 @@ const string &ParsedIR::get_decoration_string(ID id, Decoration decoration) cons

 	switch (decoration)
 	{
-	case DecorationHlslSemanticGOOGLE:
-		return dec.hlsl_semantic;
+	case DecorationUserSemantic:
+		return dec.user_semantic;

 	case DecorationUserTypeGOOGLE:
 		return dec.user_type;
@@ -693,6 +742,10 @@ void ParsedIR::unset_decoration(ID id, Decoration decoration)
 		dec.offset = 0;
 		break;

+	case DecorationOffsetIdEXT:
+		dec.offset_id = 0;
+		break;
+
 	case DecorationXfbBuffer:
 		dec.xfb_buffer = 0;
 		break;
@@ -721,14 +774,18 @@ void ParsedIR::unset_decoration(ID id, Decoration decoration)
 		dec.spec_id = 0;
 		break;

-	case DecorationHlslSemanticGOOGLE:
-		dec.hlsl_semantic.clear();
+	case DecorationUserSemantic:
+		dec.user_semantic.clear();
 		break;

 	case DecorationFPRoundingMode:
 		dec.fp_rounding_mode = FPRoundingModeMax;
 		break;

+	case DecorationFPFastMathMode:
+		dec.fp_fast_math_mode = FPFastMathModeMaskNone;
+		break;
+
 	case DecorationHlslCounterBufferGOOGLE:
 	{
 		auto &counter = meta[id].hlsl_magic_counter_buffer;
@@ -775,6 +832,8 @@ uint32_t ParsedIR::get_member_decoration(TypeID id, uint32_t index, Decoration d
 		return dec.binding;
 	case DecorationOffset:
 		return dec.offset;
+	case DecorationOffsetIdEXT:
+		return dec.offset_id;
 	case DecorationXfbBuffer:
 		return dec.xfb_buffer;
 	case DecorationXfbStride:
@@ -783,6 +842,8 @@ uint32_t ParsedIR::get_member_decoration(TypeID id, uint32_t index, Decoration d
 		return dec.stream;
 	case DecorationSpecId:
 		return dec.spec_id;
+	case DecorationMatrixStride:
+		return dec.matrix_stride;
 	case DecorationIndex:
 		return dec.index;
 	default:
@@ -811,8 +872,8 @@ void ParsedIR::set_member_decoration_string(TypeID id, uint32_t index, Decoratio

 	switch (decoration)
 	{
-	case DecorationHlslSemanticGOOGLE:
-		dec.hlsl_semantic = argument;
+	case DecorationUserSemantic:
+		dec.user_semantic = argument;
 		break;

 	default:
@@ -832,8 +893,8 @@ const string &ParsedIR::get_member_decoration_string(TypeID id, uint32_t index,

 		switch (decoration)
 		{
-		case DecorationHlslSemanticGOOGLE:
-			return dec.hlsl_semantic;
+		case DecorationUserSemantic:
+			return dec.user_semantic;

 		default:
 			return empty_string;
@@ -870,6 +931,10 @@ void ParsedIR::unset_member_decoration(TypeID id, uint32_t index, Decoration dec
 		dec.offset = 0;
 		break;

+	case DecorationOffsetIdEXT:
+		dec.offset_id = 0;
+		break;
+
 	case DecorationXfbBuffer:
 		dec.xfb_buffer = 0;
 		break;
@@ -886,8 +951,8 @@ void ParsedIR::unset_member_decoration(TypeID id, uint32_t index, Decoration dec
 		dec.spec_id = 0;
 		break;

-	case DecorationHlslSemanticGOOGLE:
-		dec.hlsl_semantic.clear();
+	case DecorationUserSemantic:
+		dec.user_semantic.clear();
 		break;

 	default:
@@ -925,6 +990,8 @@ void ParsedIR::reset_all_of_type(Types type)

 void ParsedIR::add_typed_id(Types type, ID id)
 {
+	assert(id < ids.size());
+
 	if (loop_iteration_depth_hard != 0)
 		SPIRV_CROSS_THROW("Cannot add typed ID while looping over it.");

@@ -1027,6 +1094,8 @@ ParsedIR::LoopLock &ParsedIR::LoopLock::operator=(LoopLock &&other) SPIRV_CROSS_

 void ParsedIR::make_constant_null(uint32_t id, uint32_t type, bool add_to_typed_id_set)
 {
+	assert(id < ids.size());
+
 	auto &constant_type = get<SPIRType>(type);

 	if (constant_type.pointer)
@@ -1043,16 +1112,21 @@ void ParsedIR::make_constant_null(uint32_t id, uint32_t type, bool add_to_typed_
 		uint32_t parent_id = increase_bound_by(1);
 		make_constant_null(parent_id, constant_type.parent_type, add_to_typed_id_set);

-		if (!constant_type.array_size_literal.back())
-			SPIRV_CROSS_THROW("Array size of OpConstantNull must be a literal.");
+		// The array size of OpConstantNull can be either literal or specialization constant.
+		// In the latter case, we cannot take the value as-is, as it can be changed to anything.
+		// Rather, we assume it to be *one* for the sake of initializer.
+		bool is_literal_array_size = constant_type.array_size_literal.back();
+		uint32_t count = is_literal_array_size ? constant_type.array.back() : 1;

-		SmallVector<uint32_t> elements(constant_type.array.back());
-		for (uint32_t i = 0; i < constant_type.array.back(); i++)
+		SmallVector<uint32_t> elements(count);
+		for (uint32_t i = 0; i < count; i++)
 			elements[i] = parent_id;

 		if (add_to_typed_id_set)
 			add_typed_id(TypeConstant, id);
-		variant_set<SPIRConstant>(ids[id], type, elements.data(), uint32_t(elements.size()), false).self = id;
+		auto& constant = variant_set<SPIRConstant>(ids[id], type, elements.data(), uint32_t(elements.size()), false);
+		constant.self = id;
+		constant.is_null_array_specialized_length = !is_literal_array_size;
 	}
 	else if (!constant_type.member_types.empty())
 	{
@@ -30,6 +30,7 @@

 namespace SPIRV_CROSS_NAMESPACE
 {
+using namespace SPIRV_CROSS_SPV_HEADER_NAMESPACE;

 // This data structure holds all information needed to perform cross-compilation and reflection.
 // It is the output of the Parser, but any implementation could create this structure.
@@ -87,7 +88,7 @@ public:

 	// Declared capabilities and extensions in the SPIR-V module.
 	// Not really used except for reflection at the moment.
-	SmallVector<spv::Capability> declared_capabilities;
+	SmallVector<Capability> declared_capabilities;
 	SmallVector<std::string> declared_extensions;

 	// Meta data about blocks. The cross-compiler needs to query if a block is either of these types.
@@ -109,20 +110,43 @@ public:
 	std::unordered_map<FunctionID, SPIREntryPoint> entry_points;
 	FunctionID default_entry_point = 0;

+	// A "library" module has no OpEntryPoint and instead exports symbols via
+	// OpDecorate ... LinkageAttributes ... Export. These vectors keep track
+	// of all these exports and specifically the function exports.
+	bool is_library_module = false;
+	SmallVector<uint32_t> library_exports;
+	SmallVector<FunctionID> library_exported_functions;
+
 	struct Source
 	{
+		SourceLanguage lang = SourceLanguageUnknown;
 		uint32_t version = 0;
 		bool es = false;
 		bool known = false;
 		bool hlsl = false;

+		ID file_id = 0; // string
+		ID define_id = 0; // only non-zero for DebugSource
+		std::string source;
+
+		struct Marker
+		{
+			ID line; // in source
+			ID col; // in source
+			ID offset; // in spirv stream
+			ID function_id;
+			ID block_id;
+		};
+
+		SmallVector<Marker> line_markers; // sorted by line
+
 		Source() = default;
 	};

-	Source source;
+	std::vector<Source> sources;

-	spv::AddressingModel addressing_model = spv::AddressingModelMax;
-	spv::MemoryModel memory_model = spv::MemoryModelMax;
+	AddressingModel addressing_model = AddressingModelMax;
+	MemoryModel memory_model = MemoryModelMax;

 	// Decoration handling methods.
 	// Can be useful for simple "raw" reflection.
@@ -130,25 +154,25 @@ public:
 	// and might as well just have the whole suite of decoration/name handling in one place.
 	void set_name(ID id, const std::string &name);
 	const std::string &get_name(ID id) const;
-	void set_decoration(ID id, spv::Decoration decoration, uint32_t argument = 0);
-	void set_decoration_string(ID id, spv::Decoration decoration, const std::string &argument);
-	bool has_decoration(ID id, spv::Decoration decoration) const;
-	uint32_t get_decoration(ID id, spv::Decoration decoration) const;
-	const std::string &get_decoration_string(ID id, spv::Decoration decoration) const;
+	void set_decoration(ID id, Decoration decoration, uint32_t argument = 0);
+	void set_decoration_string(ID id, Decoration decoration, const std::string &argument);
+	bool has_decoration(ID id, Decoration decoration) const;
+	uint32_t get_decoration(ID id, Decoration decoration) const;
+	const std::string &get_decoration_string(ID id, Decoration decoration) const;
 	const Bitset &get_decoration_bitset(ID id) const;
-	void unset_decoration(ID id, spv::Decoration decoration);
+	void unset_decoration(ID id, Decoration decoration);

 	// Decoration handling methods (for members of a struct).
 	void set_member_name(TypeID id, uint32_t index, const std::string &name);
 	const std::string &get_member_name(TypeID id, uint32_t index) const;
-	void set_member_decoration(TypeID id, uint32_t index, spv::Decoration decoration, uint32_t argument = 0);
-	void set_member_decoration_string(TypeID id, uint32_t index, spv::Decoration decoration,
+	void set_member_decoration(TypeID id, uint32_t index, Decoration decoration, uint32_t argument = 0);
+	void set_member_decoration_string(TypeID id, uint32_t index, Decoration decoration,
 	                                  const std::string &argument);
-	uint32_t get_member_decoration(TypeID id, uint32_t index, spv::Decoration decoration) const;
-	const std::string &get_member_decoration_string(TypeID id, uint32_t index, spv::Decoration decoration) const;
-	bool has_member_decoration(TypeID id, uint32_t index, spv::Decoration decoration) const;
+	uint32_t get_member_decoration(TypeID id, uint32_t index, Decoration decoration) const;
+	const std::string &get_member_decoration_string(TypeID id, uint32_t index, Decoration decoration) const;
+	bool has_member_decoration(TypeID id, uint32_t index, Decoration decoration) const;
 	const Bitset &get_member_decoration_bitset(TypeID id, uint32_t index) const;
-	void unset_member_decoration(TypeID id, uint32_t index, spv::Decoration decoration);
+	void unset_member_decoration(TypeID id, uint32_t index, Decoration decoration);

 	void mark_used_as_array_length(ID id);
 	uint32_t increase_bound_by(uint32_t count);
@@ -169,7 +193,7 @@ public:
 		~LoopLock();

 	private:
-		uint32_t *lock;
+		uint32_t *lock = nullptr;
 	};

 	// This must be held while iterating over a type ID array.
@@ -231,7 +255,6 @@ public:

 	uint32_t get_spirv_version() const;

-private:
 	template <typename T>
 	T &get(uint32_t id)
 	{
@@ -244,6 +267,7 @@ private:
 		return variant_get<T>(ids[id]);
 	}

+private:
 	mutable uint32_t loop_iteration_depth_hard = 0;
 	mutable uint32_t loop_iteration_depth_soft = 0;
 	std::string empty_string;
@@ -24,7 +24,7 @@
 #include "spirv_cross_util.hpp"
 #include "spirv_common.hpp"

-using namespace spv;
+using namespace SPIRV_CROSS_SPV_HEADER_NAMESPACE;
 using namespace SPIRV_CROSS_NAMESPACE;

 namespace spirv_cross_util
@@ -34,10 +34,10 @@ void rename_interface_variable(Compiler &compiler, const SmallVector<Resource> &
 {
 	for (auto &v : resources)
 	{
-		if (!compiler.has_decoration(v.id, spv::DecorationLocation))
+		if (!compiler.has_decoration(v.id, DecorationLocation))
 			continue;

-		auto loc = compiler.get_decoration(v.id, spv::DecorationLocation);
+		auto loc = compiler.get_decoration(v.id, DecorationLocation);
 		if (loc != location)
 			continue;

@@ -61,16 +61,16 @@ void inherit_combined_sampler_bindings(Compiler &compiler)
 	auto &samplers = compiler.get_combined_image_samplers();
 	for (auto &s : samplers)
 	{
-		if (compiler.has_decoration(s.image_id, spv::DecorationDescriptorSet))
+		if (compiler.has_decoration(s.image_id, DecorationDescriptorSet))
 		{
-			uint32_t set = compiler.get_decoration(s.image_id, spv::DecorationDescriptorSet);
-			compiler.set_decoration(s.combined_id, spv::DecorationDescriptorSet, set);
+			uint32_t set = compiler.get_decoration(s.image_id, DecorationDescriptorSet);
+			compiler.set_decoration(s.combined_id, DecorationDescriptorSet, set);
 		}

-		if (compiler.has_decoration(s.image_id, spv::DecorationBinding))
+		if (compiler.has_decoration(s.image_id, DecorationBinding))
 		{
-			uint32_t binding = compiler.get_decoration(s.image_id, spv::DecorationBinding);
-			compiler.set_decoration(s.combined_id, spv::DecorationBinding, binding);
+			uint32_t binding = compiler.get_decoration(s.image_id, DecorationBinding);
+			compiler.set_decoration(s.combined_id, DecorationBinding, binding);
 		}
 	}
 }
@@ -32,6 +32,9 @@

 namespace SPIRV_CROSS_NAMESPACE
 {
+using namespace SPIRV_CROSS_SPV_HEADER_NAMESPACE;
+struct GlslConstantNameMapping;
+
 enum PlsFormat
 {
 	PlsNone = 0,
@@ -66,7 +69,9 @@ enum AccessChainFlagBits
 	ACCESS_CHAIN_SKIP_REGISTER_EXPRESSION_READ_BIT = 1 << 3,
 	ACCESS_CHAIN_LITERAL_MSB_FORCE_ID = 1 << 4,
 	ACCESS_CHAIN_FLATTEN_ALL_MEMBERS_BIT = 1 << 5,
-	ACCESS_CHAIN_FORCE_COMPOSITE_BIT = 1 << 6
+	ACCESS_CHAIN_FORCE_COMPOSITE_BIT = 1 << 6,
+	ACCESS_CHAIN_PTR_CHAIN_POINTER_ARITH_BIT = 1 << 7,
+	ACCESS_CHAIN_PTR_CHAIN_CAST_TO_SCALAR_BIT = 1 << 8
 };
 typedef uint32_t AccessChainFlags;

@@ -154,6 +159,9 @@ public:
 		// If non-zero, controls layout(num_views = N) in; in GL_OVR_multiview2.
 		uint32_t ovr_multiview_view_count = 0;

+		// Emit the entry point name in SPIR-V rather than "main".
+		bool use_entry_point_name = false;
+
 		enum Precision
 		{
 			DontCare,
@@ -285,7 +293,7 @@ public:
 	// This option is only meaningful for MSL and HLSL, since GLSL matches by location directly.
 	// Masking builtins only takes effect if the builtin in question is part of the stage output interface.
 	void mask_stage_output_by_location(uint32_t location, uint32_t component);
-	void mask_stage_output_by_builtin(spv::BuiltIn builtin);
+	void mask_stage_output_by_builtin(BuiltIn builtin);

 	// Allow to control how to format float literals in the output.
 	// Set to "nullptr" to use the default "convert_to_string" function.
@@ -295,6 +303,22 @@ public:
 		float_formatter = formatter;
 	}

+	// Returns the macro name corresponding to constant id
+	std::string constant_value_macro_name(uint32_t id) const;
+
+	// Rather than using layout(descriptor_heap), emit layout(set, binding).
+	// This intended to be compatible with descriptor buffers, legacy descriptor indexing,
+	// or when the heap descriptors require unusual kinds of mapping in the Vulkan API
+	// which is not expressible by GLSL directly.
+	//
+	// ResourceTypeUnknown can be used as a default catch-all mapping.
+	// dim can be used to disambiguate between texel buffers and images since they are both image types,
+	// but use different descriptor types in the Vulkan API.
+	// No distinction is made between 1D/2D/3D/Cube textures.
+	// The default argument of DimMax maps to both texel buffers and images.
+	// dim is ignored for ResourceTypeUnknown.
+	void remap_descriptor_heap(ResourceType type, uint32_t desc_set, uint32_t binding, Dim dim = DimMax);
+
 protected:
 	struct ShaderSubgroupSupportHelper
 	{
@@ -382,7 +406,7 @@ protected:
 	};

 	// TODO remove this function when all subgroup ops are supported (or make it always return true)
-	static bool is_supported_subgroup_op_in_opengl(spv::Op op, const uint32_t *ops);
+	static bool is_supported_subgroup_op_in_opengl(Op op, const uint32_t *ops);

 	void reset(uint32_t iteration_count);
 	void emit_function(SPIRFunction &func, const Bitset &return_flags);
@@ -409,7 +433,7 @@ protected:

 	// For relax_nan_checks.
 	GLSLstd450 get_remapped_glsl_op(GLSLstd450 std450_op) const;
-	spv::Op get_remapped_spirv_op(spv::Op op) const;
+	Op get_remapped_spirv_op(Op op) const;

 	virtual void emit_glsl_op(uint32_t result_type, uint32_t result_id, uint32_t op, const uint32_t *args,
 	                          uint32_t count);
@@ -421,6 +445,8 @@ protected:
 	                                                   const uint32_t *args, uint32_t count);
 	virtual void emit_spv_amd_gcn_shader_op(uint32_t result_type, uint32_t result_id, uint32_t op, const uint32_t *args,
 	                                        uint32_t count);
+	void emit_non_semantic_shader_debug_info(uint32_t result_type, uint32_t result_id, uint32_t op,
+	                                         const uint32_t *args, uint32_t count);
 	virtual void emit_header();
 	void emit_line_directive(uint32_t file_id, uint32_t line_literal);
 	void build_workgroup_size(SmallVector<std::string> &arguments, const SpecializationConstant &x,
@@ -434,11 +460,10 @@ protected:
 	                                  SmallVector<uint32_t> &inherited_expressions);
 	virtual void emit_subgroup_op(const Instruction &i);
 	virtual std::string type_to_glsl(const SPIRType &type, uint32_t id = 0);
-	virtual std::string builtin_to_glsl(spv::BuiltIn builtin, spv::StorageClass storage);
+	virtual std::string builtin_to_glsl(BuiltIn builtin, StorageClass storage);
 	virtual void emit_struct_member(const SPIRType &type, uint32_t member_type_id, uint32_t index,
 	                                const std::string &qualifier = "", uint32_t base_offset = 0);
-	virtual void emit_struct_padding_target(const SPIRType &type);
-	virtual std::string image_type_glsl(const SPIRType &type, uint32_t id = 0);
+	virtual std::string image_type_glsl(const SPIRType &type, uint32_t id = 0, bool member = false);
 	std::string constant_expression(const SPIRConstant &c,
 	                                bool inside_block_like_struct_scope = false,
 	                                bool inside_struct_scope = false);
@@ -446,8 +471,9 @@ protected:
 	virtual std::string constant_expression_vector(const SPIRConstant &c, uint32_t vector);
 	virtual void emit_fixup();
 	virtual std::string variable_decl(const SPIRType &type, const std::string &name, uint32_t id = 0);
-	virtual bool variable_decl_is_remapped_storage(const SPIRVariable &var, spv::StorageClass storage) const;
+	virtual bool variable_decl_is_remapped_storage(const SPIRVariable &var, StorageClass storage) const;
 	virtual std::string to_func_call_arg(const SPIRFunction::Parameter &arg, uint32_t id);
+	virtual void emit_workgroup_initialization(const SPIRVariable &var);

 	struct TextureFunctionBaseArguments
 	{
@@ -477,7 +503,7 @@ protected:
 		uint32_t coord = 0, coord_components = 0, dref = 0;
 		uint32_t grad_x = 0, grad_y = 0, lod = 0, offset = 0;
 		uint32_t bias = 0, component = 0, sample = 0, sparse_texel = 0, min_lod = 0;
-		bool nonuniform_expression = false;
+		bool nonuniform_expression = false, has_array_offsets = false;
 	};
 	virtual std::string to_function_args(const TextureFunctionArguments &args, bool *p_forward);

@@ -490,7 +516,7 @@ protected:
 	virtual std::string unpack_expression_type(std::string expr_str, const SPIRType &type, uint32_t physical_type_id,
 	                                           bool packed_type, bool row_major);

-	virtual bool builtin_translates_to_nonarray(spv::BuiltIn builtin) const;
+	virtual bool builtin_translates_to_nonarray(BuiltIn builtin) const;

 	virtual bool is_user_type_structured(uint32_t id) const;

@@ -564,8 +590,8 @@ protected:

 	Options options;

-	virtual std::string type_to_array_glsl(
-	    const SPIRType &type); // Allow Metal to use the array<T> template to make arrays a value type
+	// Allow Metal to use the array<T> template to make arrays a value type
+	virtual std::string type_to_array_glsl(const SPIRType &type, uint32_t variable_id);
 	std::string to_array_size(const SPIRType &type, uint32_t index);
 	uint32_t to_array_size_literal(const SPIRType &type, uint32_t index) const;
 	uint32_t to_array_size_literal(const SPIRType &type) const;
@@ -620,6 +646,8 @@ protected:
 		const char *uint16_t_literal_suffix = "us";
 		const char *nonuniform_qualifier = "nonuniformEXT";
 		const char *boolean_mix_function = "mix";
+		const char *printf_function = "debugPrintfEXT";
+		std::string constant_null_initializer = "";
 		SPIRType::BaseType boolean_in_struct_remapped_type = SPIRType::Boolean;
 		bool swizzle_is_function = false;
 		bool shared_is_implied = false;
@@ -627,6 +655,7 @@ protected:
 		bool explicit_struct_type = false;
 		bool use_initializer_list = false;
 		bool use_typed_initializer_list = false;
+		bool requires_matching_array_initializer = false;
 		bool can_declare_struct_inline = true;
 		bool can_declare_arrays_inline = true;
 		bool native_row_major_matrix = true;
@@ -653,20 +682,24 @@ protected:
 		bool workgroup_size_is_hidden = false;
 		bool requires_relaxed_precision_analysis = false;
 		bool implicit_c_integer_promotion_rules = false;
+		bool supports_spec_constant_array_size = true;
+		bool requires_phi_undef_zero_init = false;
 	} backend;

 	void emit_struct(SPIRType &type);
 	void emit_resources();
-	void emit_extension_workarounds(spv::ExecutionModel model);
-	void emit_subgroup_arithmetic_workaround(const std::string &func, spv::Op op, spv::GroupOperation group_op);
+	void emit_extension_workarounds(ExecutionModel model);
+	void emit_subgroup_arithmetic_workaround(const std::string &func, Op op, GroupOperation group_op);
 	void emit_polyfills(uint32_t polyfills, bool relaxed);
-	void emit_buffer_block_native(const SPIRVariable &var);
+	void emit_buffer_block_native(const SPIRVariable *var, const DescriptorHeapMeta *heap_meta = nullptr);
+	std::string to_buffer_pointer_name_prefix(uint32_t ptr_id) const;
+	static std::string heap_meta_to_prefix(const DescriptorHeapMeta &meta);
 	void emit_buffer_reference_block(uint32_t type_id, bool forward_declaration);
 	void emit_buffer_block_legacy(const SPIRVariable &var);
 	void emit_buffer_block_flattened(const SPIRVariable &type);
-	void fixup_implicit_builtin_block_names(spv::ExecutionModel model);
-	void emit_declared_builtin_block(spv::StorageClass storage, spv::ExecutionModel model);
-	bool should_force_emit_builtin_block(spv::StorageClass storage);
+	void fixup_implicit_builtin_block_names(ExecutionModel model);
+	void emit_declared_builtin_block(StorageClass storage, ExecutionModel model);
+	bool should_force_emit_builtin_block(StorageClass storage);
 	void emit_push_constant_block_vulkan(const SPIRVariable &var);
 	void emit_push_constant_block_glsl(const SPIRVariable &var);
 	void emit_interface_block(const SPIRVariable &type);
@@ -676,8 +709,9 @@ protected:
 	void emit_flattened_io_block_member(const std::string &basename, const SPIRType &type, const char *qual,
 	                                    const SmallVector<uint32_t> &indices);
 	void emit_block_chain(SPIRBlock &block);
+	BlockID emit_block_chain_inner(SPIRBlock &block);
+	void emit_block_chain_cleanup(SPIRBlock &block);
 	void emit_hoisted_temporaries(SmallVector<std::pair<TypeID, ID>> &temporaries);
-	std::string constant_value_macro_name(uint32_t id);
 	int get_constant_mapping_to_workgroup_component(const SPIRConstant &constant) const;
 	void emit_constant(const SPIRConstant &constant);
 	void emit_specialization_constant_op(const SPIRConstantOp &constant);
@@ -693,6 +727,7 @@ protected:
 	void emit_variable_temporary_copies(const SPIRVariable &var);

 	bool should_dereference(uint32_t id);
+	bool should_dereference_caller_param(uint32_t id);
 	bool should_forward(uint32_t id) const;
 	bool should_suppress_usage_tracking(uint32_t id) const;
 	void emit_mix_op(uint32_t result_type, uint32_t id, uint32_t left, uint32_t right, uint32_t lerp);
@@ -746,22 +781,28 @@ protected:
 	bool expression_read_implies_multiple_reads(uint32_t id) const;
 	SPIRExpression &emit_op(uint32_t result_type, uint32_t result_id, const std::string &rhs, bool forward_rhs,
 	                        bool suppress_usage_tracking = false);
+	void emit_transposed_op(uint32_t result_type, uint32_t result_id, const std::string &rhs, bool forward_rhs);

 	void access_chain_internal_append_index(std::string &expr, uint32_t base, const SPIRType *type,
 	                                        AccessChainFlags flags, bool &access_chain_is_arrayed, uint32_t index);

 	std::string access_chain_internal(uint32_t base, const uint32_t *indices, uint32_t count, AccessChainFlags flags,
-	                                  AccessChainMeta *meta);
+	                                  AccessChainMeta *meta, const SPIRType *untyped_data_type);

-	spv::StorageClass get_expression_effective_storage_class(uint32_t ptr);
+	// Only meaningful on backends with physical pointer support ala MSL.
+	// Relevant for PtrAccessChain / BDA.
+	virtual uint32_t get_physical_type_id_stride(TypeID type_id) const;
+
+	StorageClass get_expression_effective_storage_class(uint32_t ptr);
 	virtual bool access_chain_needs_stage_io_builtin_translation(uint32_t base);

-	virtual void check_physical_type_cast(std::string &expr, const SPIRType *type, uint32_t physical_type);
+	virtual bool check_physical_type_cast(std::string &expr, const SPIRType *type, uint32_t physical_type);
 	virtual bool prepare_access_chain_for_scalar_access(std::string &expr, const SPIRType &type,
-	                                                    spv::StorageClass storage, bool &is_packed);
+	                                                    StorageClass storage, bool &is_packed);

 	std::string access_chain(uint32_t base, const uint32_t *indices, uint32_t count, const SPIRType &target_type,
-	                         AccessChainMeta *meta = nullptr, bool ptr_chain = false);
+	                         AccessChainMeta *meta = nullptr, bool ptr_chain = false,
+	                         const SPIRType *untyped_data_type = nullptr);

 	std::string flattened_access_chain(uint32_t base, const uint32_t *indices, uint32_t count,
 	                                   const SPIRType &target_type, uint32_t offset, uint32_t matrix_stride,
@@ -784,10 +825,16 @@ protected:
 	const char *index_to_swizzle(uint32_t index);
 	std::string remap_swizzle(const SPIRType &result_type, uint32_t input_components, const std::string &expr);
 	std::string declare_temporary(uint32_t type, uint32_t id);
+	bool can_declare_inline_temporary(uint32_t id) const;
 	void emit_uninitialized_temporary(uint32_t type, uint32_t id);
 	SPIRExpression &emit_uninitialized_temporary_expression(uint32_t type, uint32_t id);
-	void append_global_func_args(const SPIRFunction &func, uint32_t index, SmallVector<std::string> &arglist);
+	virtual void append_global_func_args(const SPIRFunction &func, uint32_t index, SmallVector<std::string> &arglist);
 	std::string to_non_uniform_aware_expression(uint32_t id);
+	std::string to_atomic_ptr_expression(uint32_t id);
+	std::string to_pretty_expression_if_int_constant(
+			uint32_t id,
+			const GlslConstantNameMapping *mapping_start, const GlslConstantNameMapping *mapping_end,
+			bool register_expression_read = true);
 	std::string to_expression(uint32_t id, bool register_expression_read = true);
 	std::string to_composite_constructor_expression(const SPIRType &parent_type, uint32_t id, bool block_like_type);
 	std::string to_rerolled_array_expression(const SPIRType &parent_type, const std::string &expr, const SPIRType &type);
@@ -816,15 +863,15 @@ protected:
 	void emit_output_variable_initializer(const SPIRVariable &var);
 	std::string to_precision_qualifiers_glsl(uint32_t id);
 	virtual const char *to_storage_qualifiers_glsl(const SPIRVariable &var);
-	std::string flags_to_qualifiers_glsl(const SPIRType &type, const Bitset &flags);
-	const char *format_to_glsl(spv::ImageFormat format);
+	std::string flags_to_qualifiers_glsl(const SPIRType &type, uint32_t id, const Bitset &flags);
+	const char *format_to_glsl(ImageFormat format);
 	virtual std::string layout_for_member(const SPIRType &type, uint32_t index);
 	virtual std::string to_interpolation_qualifiers(const Bitset &flags);
 	std::string layout_for_variable(const SPIRVariable &variable);
 	std::string to_combined_image_sampler(VariableID image_id, VariableID samp_id);
 	virtual bool skip_argument(uint32_t id) const;
 	virtual bool emit_array_copy(const char *expr, uint32_t lhs_id, uint32_t rhs_id,
-	                             spv::StorageClass lhs_storage, spv::StorageClass rhs_storage);
+	                             StorageClass lhs_storage, StorageClass rhs_storage);
 	virtual void emit_block_hints(const SPIRBlock &block);
 	virtual std::string to_initializer_expression(const SPIRVariable &var);
 	virtual std::string to_zero_initialized_expression(uint32_t type_id);
@@ -833,7 +880,9 @@ protected:
 	bool buffer_is_packing_standard(const SPIRType &type, BufferPackingStandard packing,
 	                                uint32_t *failed_index = nullptr, uint32_t start_offset = 0,
 	                                uint32_t end_offset = ~(0u));
-	std::string buffer_to_packing_standard(const SPIRType &type, bool support_std430_without_scalar_layout);
+	std::string buffer_to_packing_standard(const SPIRType &type,
+	                                       bool support_std430_without_scalar_layout,
+	                                       bool support_enhanced_layouts);

 	uint32_t type_to_packed_base_size(const SPIRType &type, BufferPackingStandard packing);
 	uint32_t type_to_packed_alignment(const SPIRType &type, const Bitset &flags, BufferPackingStandard packing);
@@ -931,12 +980,32 @@ protected:
 		PolyfillMatrixInverse2x2 = 1 << 6,
 		PolyfillMatrixInverse3x3 = 1 << 7,
 		PolyfillMatrixInverse4x4 = 1 << 8,
+		PolyfillNMin16 = 1 << 9,
+		PolyfillNMin32 = 1 << 10,
+		PolyfillNMin64 = 1 << 11,
+		PolyfillNMax16 = 1 << 12,
+		PolyfillNMax32 = 1 << 13,
+		PolyfillNMax64 = 1 << 14,
+		PolyfillNClamp16 = 1 << 15,
+		PolyfillNClamp32 = 1 << 16,
+		PolyfillNClamp64 = 1 << 17,
 	};

 	uint32_t required_polyfills = 0;
 	uint32_t required_polyfills_relaxed = 0;
 	void require_polyfill(Polyfill polyfill, bool relaxed);

+	struct IntegerDotProduct
+	{
+		Id result_type;
+		Id argument_type[2];
+		Op op;
+	};
+	SmallVector<IntegerDotProduct> integer_dot_products_polyfills;
+	void add_integer_dot_product_polyfill(const IntegerDotProduct &idot);
+	std::string integer_dot_product_entry_point(const IntegerDotProduct &idot);
+	void emit_polyfills_integer_dot_product();
+
 	bool ray_tracing_is_khr = false;
 	bool barycentric_is_nv = false;
 	void ray_tracing_khr_fixup_locations();
@@ -982,16 +1051,19 @@ protected:

 	std::string emit_for_loop_initializers(const SPIRBlock &block);
 	void emit_while_loop_initializers(const SPIRBlock &block);
+	std::string undef_loop_variable_initializer_suffix(const SPIRVariable &var);
 	bool for_loop_initializers_are_same_type(const SPIRBlock &block);
 	bool optimize_read_modify_write(const SPIRType &type, const std::string &lhs, const std::string &rhs);
 	void fixup_image_load_store_access();

 	bool type_is_empty(const SPIRType &type);

-	bool can_use_io_location(spv::StorageClass storage, bool block);
+	bool can_use_io_location(StorageClass storage, bool block);
 	const Instruction *get_next_instruction_in_block(const Instruction &instr);
 	static uint32_t mask_relevant_memory_semantics(uint32_t semantics);

+	std::string convert_floate4m3_to_string(const SPIRConstant &value, uint32_t col, uint32_t row);
+	std::string convert_floate5m2_to_string(const SPIRConstant &value, uint32_t col, uint32_t row);
 	std::string convert_half_to_string(const SPIRConstant &value, uint32_t col, uint32_t row);
 	std::string convert_float_to_string(const SPIRConstant &value, uint32_t col, uint32_t row);
 	std::string convert_double_to_string(const SPIRConstant &value, uint32_t col, uint32_t row);
@@ -1001,7 +1073,7 @@ protected:
 	// Builtins in GLSL are always specific signedness, but the SPIR-V can declare them
 	// as either unsigned or signed.
 	// Sometimes we will need to automatically perform casts on load and store to make this work.
-	virtual SPIRType::BaseType get_builtin_basetype(spv::BuiltIn builtin, SPIRType::BaseType default_type);
+	virtual SPIRType::BaseType get_builtin_basetype(BuiltIn builtin, SPIRType::BaseType default_type);
 	virtual void cast_to_variable_store(uint32_t target_id, std::string &expr, const SPIRType &expr_type);
 	virtual void cast_from_variable_load(uint32_t source_id, std::string &expr, const SPIRType &expr_type);
 	void unroll_array_from_complex_load(uint32_t target_id, uint32_t source_id, std::string &expr);
@@ -1012,7 +1084,7 @@ protected:
 	void disallow_forwarding_in_expression_chain(const SPIRExpression &expr);

 	bool expression_is_constant_null(uint32_t id) const;
-	bool expression_is_non_value_type_array(uint32_t ptr);
+	bool expression_is_non_value_type_array(uint32_t value_type_id, uint32_t ptr);
 	virtual void emit_store_statement(uint32_t lhs_expression, uint32_t rhs_expression);

 	uint32_t get_integer_width_for_instruction(const Instruction &instr) const;
@@ -1030,7 +1102,7 @@ protected:
 	static const char *vector_swizzle(int vecsize, int index);

 	bool is_stage_output_location_masked(uint32_t location, uint32_t component) const;
-	bool is_stage_output_builtin_masked(spv::BuiltIn builtin) const;
+	bool is_stage_output_builtin_masked(BuiltIn builtin) const;
 	bool is_stage_output_variable_masked(const SPIRVariable &var) const;
 	bool is_stage_output_block_member_masked(const SPIRVariable &var, uint32_t index, bool strip_array) const;
 	bool is_per_primitive_variable(const SPIRVariable &var) const;
@@ -1043,6 +1115,20 @@ protected:
 	std::string format_float(float value) const;
 	std::string format_double(double value) const;

+	uint32_t get_fp_fast_math_flags_for_op(uint32_t result_type, uint32_t id) const;
+	bool has_legacy_nocontract(uint32_t result_type, uint32_t id) const;
+
+	struct DescriptorHeapMapping
+	{
+		ResourceType type;
+		uint32_t desc_set;
+		uint32_t binding;
+		Dim dim;
+	};
+	SmallVector<DescriptorHeapMapping> descriptor_heap_mappings;
+	bool is_descriptor_non_uniform(uint32_t id) const;
+	std::string to_descriptor_heap_layout(const SPIRType &type, StorageClass storage = StorageClassUniformConstant) const;
+
 private:
 	void init();

@@ -29,6 +29,7 @@

 namespace SPIRV_CROSS_NAMESPACE
 {
+using namespace SPIRV_CROSS_SPV_HEADER_NAMESPACE;
 // Interface which remaps vertex inputs to a fixed semantic name to make linking easier.
 struct HLSLVertexAttributeRemap
 {
@@ -87,7 +88,7 @@ using HLSLBindingFlags = uint32_t;
 // For deeper control of push constants, set_root_constant_layouts() can be used instead.
 struct HLSLResourceBinding
 {
-	spv::ExecutionModel stage = spv::ExecutionModelMax;
+	ExecutionModel stage = ExecutionModelMax;
 	uint32_t desc_set = 0;
 	uint32_t binding = 0;

@@ -150,6 +151,9 @@ public:
 		// This relies on UserTypeGOOGLE to encode the buffer type either as "structuredbuffer" or "rwstructuredbuffer"
 		// whereas the type can be extended with an optional subtype, e.g. "structuredbuffer:int".
 		bool preserve_structured_buffers = false;
+
+		// Use UserSemantic decoration info (if specified), otherwise use default mechanism (such as add_vertex_attribute_remap or TEXCOORD#).
+		bool user_semantic = false;
 	};

 	explicit CompilerHLSL(std::vector<uint32_t> spirv_)
@@ -216,7 +220,7 @@ public:
 	// is_hlsl_resource_binding_used() will return true after calling ::compile() if
 	// the set/binding combination was used by the HLSL code.
 	void add_hlsl_resource_binding(const HLSLResourceBinding &resource);
-	bool is_hlsl_resource_binding_used(spv::ExecutionModel model, uint32_t set, uint32_t binding) const;
+	bool is_hlsl_resource_binding_used(ExecutionModel model, uint32_t set, uint32_t binding) const;

 	// Controls which storage buffer bindings will be forced to be declared as UAVs.
 	void set_hlsl_force_storage_buffer_as_uav(uint32_t desc_set, uint32_t binding);
@@ -231,6 +235,7 @@ private:
 	std::string image_type_hlsl(const SPIRType &type, uint32_t id);
 	std::string image_type_hlsl_modern(const SPIRType &type, uint32_t id);
 	std::string image_type_hlsl_legacy(const SPIRType &type, uint32_t id);
+	uint32_t input_vertices_from_execution_mode(SPIREntryPoint &execution) const;
 	void emit_function_prototype(SPIRFunction &func, const Bitset &return_flags) override;
 	void emit_hlsl_entry_point();
 	void emit_header() override;
@@ -254,11 +259,13 @@ private:
 	void emit_specialization_constants_and_structs();
 	void emit_composite_constants();
 	void emit_fixup() override;
-	std::string builtin_to_glsl(spv::BuiltIn builtin, spv::StorageClass storage) override;
+	std::string builtin_to_glsl(BuiltIn builtin, StorageClass storage) override;
 	std::string layout_for_member(const SPIRType &type, uint32_t index) override;
 	std::string to_interpolation_qualifiers(const Bitset &flags) override;
 	std::string bitcast_glsl_op(const SPIRType &result_type, const SPIRType &argument_type) override;
 	bool emit_complex_bitcast(uint32_t result_type, uint32_t id, uint32_t op0) override;
+	void append_global_func_args(const SPIRFunction &func, uint32_t index, SmallVector<std::string> &arglist) override;
+
 	std::string to_func_call_arg(const SPIRFunction::Parameter &arg, uint32_t id) override;
 	std::string to_sampler_expression(uint32_t id);
 	std::string to_resource_binding(const SPIRVariable &var);
@@ -278,7 +285,7 @@ private:
 	                              const SmallVector<uint32_t> &composite_chain);
 	std::string write_access_chain_value(uint32_t value, const SmallVector<uint32_t> &composite_chain, bool enclose);
 	void emit_store(const Instruction &instruction);
-	void emit_atomic(const uint32_t *ops, uint32_t length, spv::Op op);
+	void emit_atomic(const uint32_t *ops, uint32_t length, Op op);
 	void emit_subgroup_op(const Instruction &i) override;
 	void emit_block_hints(const SPIRBlock &block) override;

@@ -286,13 +293,15 @@ private:
 	                        uint32_t base_offset = 0) override;
 	void emit_rayquery_function(const char *commited, const char *candidate, const uint32_t *ops);
 	void emit_mesh_tasks(SPIRBlock &block) override;
+	void emit_geometry_stream_append();

 	const char *to_storage_qualifiers_glsl(const SPIRVariable &var) override;
 	void replace_illegal_names() override;

-	SPIRType::BaseType get_builtin_basetype(spv::BuiltIn builtin, SPIRType::BaseType default_type) override;
+	SPIRType::BaseType get_builtin_basetype(BuiltIn builtin, SPIRType::BaseType default_type) override;

 	bool is_hlsl_force_storage_buffer_as_uav(ID id) const;
+	bool is_hidden_io_variable(const SPIRVariable &var) const;

 	Options hlsl_options;

@@ -375,7 +384,7 @@ private:

 	uint32_t type_to_consumed_locations(const SPIRType &type) const;

-	std::string to_semantic(uint32_t location, spv::ExecutionModel em, spv::StorageClass sc);
+	std::string to_semantic(uint32_t location, ExecutionModel em, StorageClass sc);

 	uint32_t num_workgroups_builtin = 0;
 	HLSLBindingFlags resource_binding_flags = 0;
@@ -408,6 +417,8 @@ private:
 	std::vector<TypeID> composite_selection_workaround_types;

 	std::string get_inner_entry_point_name() const;
+
+	void cast_to_variable_store(uint32_t target_id, std::string &expr, const SPIRType &expr_type) override;
 };
 } // namespace SPIRV_CROSS_NAMESPACE

@@ -33,6 +33,7 @@

 namespace SPIRV_CROSS_NAMESPACE
 {
+using namespace SPIRV_CROSS_SPV_HEADER_NAMESPACE;

 // Indicates the format of a shader interface variable. Currently limited to specifying
 // if the input is an 8-bit unsigned integer, 16-bit unsigned integer, or
@@ -78,7 +79,7 @@ struct MSLShaderInterfaceVariable
 	uint32_t location = 0;
 	uint32_t component = 0;
 	MSLShaderVariableFormat format = MSL_SHADER_VARIABLE_FORMAT_OTHER;
-	spv::BuiltIn builtin = spv::BuiltInMax;
+	BuiltIn builtin = BuiltInMax;
 	uint32_t vecsize = 0;
 	MSLShaderVariableRate rate = MSL_SHADER_VARIABLE_RATE_PER_VERTEX;
 };
@@ -104,7 +105,7 @@ struct MSLShaderInterfaceVariable
 // become a [[buffer(N)]], [[texture(N)]] or [[sampler(N)]] depending on the resource types used.
 struct MSLResourceBinding
 {
-	spv::ExecutionModel stage = spv::ExecutionModelMax;
+	ExecutionModel stage = ExecutionModelMax;
 	SPIRType::BaseType basetype = SPIRType::Unknown;
 	uint32_t desc_set = 0;
 	uint32_t binding = 0;
@@ -287,9 +288,6 @@ static const uint32_t kArgumentBufferBinding = ~(3u);

 static const uint32_t kMaxArgumentBuffers = 8;

-// The arbitrary maximum for the nesting of array of array copies.
-static const uint32_t kArrayCopyMultidimMax = 6;
-
 // Decompiles SPIR-V to Metal Shading Language
 class CompilerMSL : public CompilerGLSL
 {
@@ -319,6 +317,8 @@ public:
 		uint32_t shader_input_buffer_index = 22;
 		uint32_t shader_index_buffer_index = 21;
 		uint32_t shader_patch_input_buffer_index = 20;
+        uint32_t draw_id_buffer_index = 19;
+		uint32_t reversed_depth_viewport_buffer_index = 18;
 		uint32_t shader_input_wg_index = 0;
 		uint32_t device_index = 0;
 		uint32_t enable_frag_output_mask = 0xffffffff;
@@ -327,6 +327,8 @@ public:
 		// of the shader with the additional fixed sample mask.
 		uint32_t additional_fixed_sample_mask = 0xffffffff;
 		bool enable_point_size_builtin = true;
+		bool enable_point_size_default = false;
+		float default_point_size = 1.0f;
 		bool enable_frag_depth_builtin = true;
 		bool enable_frag_stencil_ref_builtin = true;
 		bool disable_rasterization = false;
@@ -338,6 +340,7 @@ public:
 		bool view_index_from_device_index = false;
 		bool dispatch_base = false;
 		bool texture_1D_as_2D = false;
+		bool emulate_reversed_depth_viewport = false;

 		// Enable use of Metal argument buffers.
 		// MSL 2.0 must also be enabled.
@@ -519,6 +522,34 @@ public:
 		// transformed.
 		bool agx_manual_cube_grad_fixup = false;

+		// Metal will discard fragments with side effects under certain circumstances prematurely.
+		// Example: CTS test dEQP-VK.fragment_operations.early_fragment.discard_no_early_fragment_tests_depth
+		// Test will render a full screen quad with varying depth [0,1] for each fragment.
+		// Each fragment will do an operation with side effects, modify the depth value and
+		// discard the fragment. The test expects the fragment to be run due to:
+		// https://registry.khronos.org/vulkan/specs/1.0-extensions/html/vkspec.html#fragops-shader-depthreplacement
+		// which states that the fragment shader must be run due to replacing the depth in shader.
+		// However, Metal may prematurely discards fragments without executing them
+		// (I believe this to be due to a greedy optimization on their end) making the test fail.
+		// This option enforces fragment execution for such cases where the fragment has operations
+		// with side effects. Provided as an option hoping Metal will fix this issue in the future.
+		bool force_fragment_with_side_effects_execution = false;
+
+		// If set, adds a depth pass through statement to circumvent the following issue:
+		// When the same depth/stencil is used as input and depth/stencil attachment, we need to
+		// force Metal to perform the depth/stencil write after fragment execution. Otherwise,
+		// Metal will write to the depth attachment before fragment execution. This happens
+		// if the fragment does not modify the depth value.
+		bool input_attachment_is_ds_attachment = false;
+
+		// If BuiltInPosition is not written, automatically disable rasterization.
+		// The result can be queried with get_is_rasterization_disabled.
+		bool auto_disable_rasterization = false;
+
+		// Use Fast Math pragmas in MSL code, based on SPIR-V float controls and FP ExecutionModes.
+		// Requires MSL 3.2 or above, and has no effect with earlier MSL versions.
+		bool use_fast_math_pragmas = false;
+
 		bool is_ios() const
 		{
 			return platform == iOS;
@@ -564,9 +595,9 @@ public:
 	// rasterization if vertex shader requires rasterization to be disabled.
 	bool get_is_rasterization_disabled() const
 	{
-		return is_rasterization_disabled && (get_entry_point().model == spv::ExecutionModelVertex ||
-		                                     get_entry_point().model == spv::ExecutionModelTessellationControl ||
-		                                     get_entry_point().model == spv::ExecutionModelTessellationEvaluation);
+		return is_rasterization_disabled && (get_entry_point().model == ExecutionModelVertex ||
+		                                     get_entry_point().model == ExecutionModelTessellationControl ||
+		                                     get_entry_point().model == ExecutionModelTessellationEvaluation);
 	}

 	// Provide feedback to calling API to allow it to pass an auxiliary
@@ -583,6 +614,14 @@ public:
 		return !buffers_requiring_array_length.empty();
 	}

+	// Provide feedback to calling API to determine if the vertex shader writes
+	// to PointSize. This allows the API to avoid declaring a point size output
+	// when it is not needed.
+	bool get_writes_to_point_size() const
+	{
+		return writes_to_point_size;
+	}
+
 	bool buffer_requires_array_length(VariableID id) const
 	{
 		return buffers_requiring_array_length.count(id) != 0;
@@ -679,20 +718,20 @@ public:
 	// This is typical for tessellation builtin inputs such as tess levels, gl_Position, etc.
 	// This returns k_unknown_location if the location was explicitly assigned with
 	// add_msl_shader_input or the builtin is not used, otherwise returns N in [[attribute(N)]].
-	uint32_t get_automatic_builtin_input_location(spv::BuiltIn builtin) const;
+	uint32_t get_automatic_builtin_input_location(BuiltIn builtin) const;

 	// If not using add_msl_shader_output, it's possible
 	// that certain builtin attributes need to be automatically assigned locations.
 	// This is typical for tessellation builtin outputs such as tess levels, gl_Position, etc.
 	// This returns k_unknown_location if the location were explicitly assigned with
 	// add_msl_shader_output or the builtin were not used, otherwise returns N in [[attribute(N)]].
-	uint32_t get_automatic_builtin_output_location(spv::BuiltIn builtin) const;
+	uint32_t get_automatic_builtin_output_location(BuiltIn builtin) const;

 	// NOTE: Only resources which are remapped using add_msl_resource_binding will be reported here.
 	// Constexpr samplers are always assumed to be emitted.
 	// No specific MSLResourceBinding remapping is required for constexpr samplers as long as they are remapped
 	// by remap_constexpr_sampler(_by_binding).
-	bool is_msl_resource_binding_used(spv::ExecutionModel model, uint32_t set, uint32_t binding) const;
+	bool is_msl_resource_binding_used(ExecutionModel model, uint32_t set, uint32_t binding) const;

 	// This must only be called after a successful call to CompilerMSL::compile().
 	// For a variable resource ID obtained through reflection API, report the automatically assigned resource index.
@@ -739,6 +778,19 @@ public:
 	void set_combined_sampler_suffix(const char *suffix);
 	const char *get_combined_sampler_suffix() const;

+	// Information about specialization constants that are translated into MSL macros
+	// instead of using function constant
+	// These must only be called after a successful call to CompilerMSL::compile().
+	bool specialization_constant_is_macro(uint32_t constant_id) const;
+
+	// Returns a mask of SPIR-V FP Fast Math Mode flags, that represents the set of flags that can be applied
+	// across all floating-point types. Each FPFastMathDefault execution mode operation identifies the flags
+	// for one floating-point type, and the value returned here is a bitwise-AND combination across all types.
+	// If incl_ops is enabled, the FPFastMathMode of any SPIR-V operations are also included in the bitwise-AND
+	// to determine the minimal fast-math that applies to all default execution modes and all operations.
+	// The returned value is also affected by execution modes SignedZeroInfNanPreserve and ContractionOff.
+	uint32_t get_fp_fast_math_flags(bool incl_ops) const;
+
 protected:
 	// An enum of SPIR-V functions that are implemented in additional
 	// source code that is added to the shader if necessary.
@@ -746,21 +798,17 @@ protected:
 	{
 		SPVFuncImplNone,
 		SPVFuncImplMod,
+		SPVFuncImplSMod,
 		SPVFuncImplRadians,
 		SPVFuncImplDegrees,
 		SPVFuncImplFindILsb,
 		SPVFuncImplFindSMsb,
 		SPVFuncImplFindUMsb,
 		SPVFuncImplSSign,
-		SPVFuncImplArrayCopyMultidimBase,
-		// Unfortunately, we cannot use recursive templates in the MSL compiler properly,
-		// so stamp out variants up to some arbitrary maximum.
-		SPVFuncImplArrayCopy = SPVFuncImplArrayCopyMultidimBase + 1,
-		SPVFuncImplArrayOfArrayCopy2Dim = SPVFuncImplArrayCopyMultidimBase + 2,
-		SPVFuncImplArrayOfArrayCopy3Dim = SPVFuncImplArrayCopyMultidimBase + 3,
-		SPVFuncImplArrayOfArrayCopy4Dim = SPVFuncImplArrayCopyMultidimBase + 4,
-		SPVFuncImplArrayOfArrayCopy5Dim = SPVFuncImplArrayCopyMultidimBase + 5,
-		SPVFuncImplArrayOfArrayCopy6Dim = SPVFuncImplArrayCopyMultidimBase + 6,
+		SPVFuncImplArrayCopy,
+		SPVFuncImplArrayCopyMultidim,
+		SPVFuncImplArrayCopyExtendedSrc,
+		SPVFuncImplArrayCopyExtendedDst,
 		SPVFuncImplTexelBufferCoords,
 		SPVFuncImplImage2DAtomicCoords, // Emulate texture2D atomic operations
 		SPVFuncImplGradientCube,
@@ -774,14 +822,15 @@ protected:
 		SPVFuncImplInverse4x4,
 		SPVFuncImplInverse3x3,
 		SPVFuncImplInverse2x2,
-		// It is very important that this come before *Swizzle and ChromaReconstruct*, to ensure it's
-		// emitted before them.
-		SPVFuncImplForwardArgs,
-		// Likewise, this must come before *Swizzle.
+		// It is very important that this come before *Swizzle, to ensure it's emitted before them.
 		SPVFuncImplGetSwizzle,
 		SPVFuncImplTextureSwizzle,
+		SPVFuncImplGatherReturn,
+		SPVFuncImplGatherCompareReturn,
 		SPVFuncImplGatherSwizzle,
 		SPVFuncImplGatherCompareSwizzle,
+		SPVFuncImplGatherConstOffsets,
+		SPVFuncImplGatherCompareConstOffsets,
 		SPVFuncImplSubgroupBroadcast,
 		SPVFuncImplSubgroupBroadcastFirst,
 		SPVFuncImplSubgroupBallot,
@@ -794,6 +843,30 @@ protected:
 		SPVFuncImplSubgroupShuffleXor,
 		SPVFuncImplSubgroupShuffleUp,
 		SPVFuncImplSubgroupShuffleDown,
+		SPVFuncImplSubgroupRotate,
+		SPVFuncImplSubgroupClusteredAdd,
+		SPVFuncImplSubgroupClusteredFAdd = SPVFuncImplSubgroupClusteredAdd,
+		SPVFuncImplSubgroupClusteredIAdd = SPVFuncImplSubgroupClusteredAdd,
+		SPVFuncImplSubgroupClusteredMul,
+		SPVFuncImplSubgroupClusteredFMul = SPVFuncImplSubgroupClusteredMul,
+		SPVFuncImplSubgroupClusteredIMul = SPVFuncImplSubgroupClusteredMul,
+		SPVFuncImplSubgroupClusteredMin,
+		SPVFuncImplSubgroupClusteredFMin = SPVFuncImplSubgroupClusteredMin,
+		SPVFuncImplSubgroupClusteredSMin = SPVFuncImplSubgroupClusteredMin,
+		SPVFuncImplSubgroupClusteredUMin = SPVFuncImplSubgroupClusteredMin,
+		SPVFuncImplSubgroupClusteredMax,
+		SPVFuncImplSubgroupClusteredFMax = SPVFuncImplSubgroupClusteredMax,
+		SPVFuncImplSubgroupClusteredSMax = SPVFuncImplSubgroupClusteredMax,
+		SPVFuncImplSubgroupClusteredUMax = SPVFuncImplSubgroupClusteredMax,
+		SPVFuncImplSubgroupClusteredAnd,
+		SPVFuncImplSubgroupClusteredBitwiseAnd = SPVFuncImplSubgroupClusteredAnd,
+		SPVFuncImplSubgroupClusteredLogicalAnd = SPVFuncImplSubgroupClusteredAnd,
+		SPVFuncImplSubgroupClusteredOr,
+		SPVFuncImplSubgroupClusteredBitwiseOr = SPVFuncImplSubgroupClusteredOr,
+		SPVFuncImplSubgroupClusteredLogicalOr = SPVFuncImplSubgroupClusteredOr,
+		SPVFuncImplSubgroupClusteredXor,
+		SPVFuncImplSubgroupClusteredBitwiseXor = SPVFuncImplSubgroupClusteredXor,
+		SPVFuncImplSubgroupClusteredLogicalXor = SPVFuncImplSubgroupClusteredXor,
 		SPVFuncImplQuadBroadcast,
 		SPVFuncImplQuadSwap,
 		SPVFuncImplReflectScalar,
@@ -823,7 +896,14 @@ protected:
 		SPVFuncImplVariableDescriptor,
 		SPVFuncImplVariableSizedDescriptor,
 		SPVFuncImplVariableDescriptorArray,
-		SPVFuncImplPaddedStd140
+		SPVFuncImplPaddedStd140,
+		SPVFuncImplPaddedArrayElement,
+		SPVFuncImplReduceAdd,
+		SPVFuncImplImageFence,
+		SPVFuncImplTextureCast,
+		SPVFuncImplMulExtended,
+		SPVFuncImplSetMeshOutputsEXT,
+		SPVFuncImplAssume,
 	};

 	// If the underlying resource has been used for comparison then duplicate loads of that resource must be too
@@ -841,6 +921,11 @@ protected:
 	void emit_function_prototype(SPIRFunction &func, const Bitset &return_flags) override;
 	void emit_sampled_image_op(uint32_t result_type, uint32_t result_id, uint32_t image_id, uint32_t samp_id) override;
 	void emit_subgroup_op(const Instruction &i) override;
+	void emit_subgroup_cluster_op(uint32_t result_type, uint32_t result_id, uint32_t cluster_size, uint32_t op0,
+	                              const char *op);
+	void emit_subgroup_cluster_op_cast(uint32_t result_type, uint32_t result_id, uint32_t cluster_size, uint32_t op0,
+	                                   const char *op, SPIRType::BaseType input_type,
+	                                   SPIRType::BaseType expected_result_type);
 	std::string to_texture_op(const Instruction &i, bool sparse, bool *forward,
 	                          SmallVector<uint32_t> &inherited_expressions) override;
 	void emit_fixup() override;
@@ -848,23 +933,26 @@ protected:
 	                             const std::string &qualifier = "");
 	void emit_struct_member(const SPIRType &type, uint32_t member_type_id, uint32_t index,
 	                        const std::string &qualifier = "", uint32_t base_offset = 0) override;
-	void emit_struct_padding_target(const SPIRType &type) override;
 	std::string type_to_glsl(const SPIRType &type, uint32_t id, bool member);
 	std::string type_to_glsl(const SPIRType &type, uint32_t id = 0) override;
 	void emit_block_hints(const SPIRBlock &block) override;
+	void emit_mesh_entry_point();
+	void emit_mesh_outputs();
+	void emit_mesh_tasks(SPIRBlock &block) override;
+	void emit_workgroup_initialization(const SPIRVariable &var) override;

 	// Allow Metal to use the array<T> template to make arrays a value type
-	std::string type_to_array_glsl(const SPIRType &type) override;
+	std::string type_to_array_glsl(const SPIRType &type, uint32_t variable_id) override;
 	std::string constant_op_expression(const SPIRConstantOp &cop) override;

-	bool variable_decl_is_remapped_storage(const SPIRVariable &variable, spv::StorageClass storage) const override;
+	bool variable_decl_is_remapped_storage(const SPIRVariable &variable, StorageClass storage) const override;

 	// GCC workaround of lambdas calling protected functions (for older GCC versions)
 	std::string variable_decl(const SPIRType &type, const std::string &name, uint32_t id = 0) override;

-	std::string image_type_glsl(const SPIRType &type, uint32_t id = 0) override;
-	std::string sampler_type(const SPIRType &type, uint32_t id);
-	std::string builtin_to_glsl(spv::BuiltIn builtin, spv::StorageClass storage) override;
+	std::string image_type_glsl(const SPIRType &type, uint32_t id, bool member) override;
+	std::string sampler_type(const SPIRType &type, uint32_t id, bool member);
+	std::string builtin_to_glsl(BuiltIn builtin, StorageClass storage) override;
 	std::string to_func_call_arg(const SPIRFunction::Parameter &arg, uint32_t id) override;
 	std::string to_name(uint32_t id, bool allow_alias = true) const override;
 	std::string to_function_name(const TextureFunctionNameArguments &args) override;
@@ -876,7 +964,7 @@ protected:
 	                                   bool is_packed, bool row_major) override;

 	// Returns true for BuiltInSampleMask because gl_SampleMask[] is an array in SPIR-V, but [[sample_mask]] is a scalar in Metal.
-	bool builtin_translates_to_nonarray(spv::BuiltIn builtin) const override;
+	bool builtin_translates_to_nonarray(BuiltIn builtin) const override;

 	std::string bitcast_glsl_op(const SPIRType &result_type, const SPIRType &argument_type) override;
 	bool emit_complex_bitcast(uint32_t result_id, uint32_t id, uint32_t op0) override;
@@ -903,6 +991,7 @@ protected:

 	bool is_tesc_shader() const;
 	bool is_tese_shader() const;
+	bool is_mesh_shader() const;

 	void preprocess_op_codes();
 	void localize_global_variables();
@@ -915,8 +1004,9 @@ protected:
 	void extract_global_variables_from_function(uint32_t func_id, std::set<uint32_t> &added_arg_ids,
 	                                            std::unordered_set<uint32_t> &global_var_ids,
 	                                            std::unordered_set<uint32_t> &processed_func_ids);
-	uint32_t add_interface_block(spv::StorageClass storage, bool patch = false);
-	uint32_t add_interface_block_pointer(uint32_t ib_var_id, spv::StorageClass storage);
+	uint32_t add_interface_block(StorageClass storage, bool patch = false);
+	uint32_t add_interface_block_pointer(uint32_t ib_var_id, StorageClass storage);
+	uint32_t add_meshlet_block(bool per_primitive);

 	struct InterfaceBlockMeta
 	{
@@ -936,37 +1026,39 @@ protected:

 	std::string to_tesc_invocation_id();
 	void emit_local_masked_variable(const SPIRVariable &masked_var, bool strip_array);
-	void add_variable_to_interface_block(spv::StorageClass storage, const std::string &ib_var_ref, SPIRType &ib_type,
+	void add_variable_to_interface_block(StorageClass storage, const std::string &ib_var_ref, SPIRType &ib_type,
 	                                     SPIRVariable &var, InterfaceBlockMeta &meta);
-	void add_composite_variable_to_interface_block(spv::StorageClass storage, const std::string &ib_var_ref,
+	void add_composite_variable_to_interface_block(StorageClass storage, const std::string &ib_var_ref,
 	                                               SPIRType &ib_type, SPIRVariable &var, InterfaceBlockMeta &meta);
-	void add_plain_variable_to_interface_block(spv::StorageClass storage, const std::string &ib_var_ref,
+	void add_plain_variable_to_interface_block(StorageClass storage, const std::string &ib_var_ref,
 	                                           SPIRType &ib_type, SPIRVariable &var, InterfaceBlockMeta &meta);
-	bool add_component_variable_to_interface_block(spv::StorageClass storage, const std::string &ib_var_ref,
+	bool add_component_variable_to_interface_block(StorageClass storage, const std::string &ib_var_ref,
 	                                               SPIRVariable &var, const SPIRType &type,
 	                                               InterfaceBlockMeta &meta);
-	void add_plain_member_variable_to_interface_block(spv::StorageClass storage,
+	void add_plain_member_variable_to_interface_block(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,
+	void add_composite_member_variable_to_interface_block(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);
+	                                                      uint32_t &location, uint32_t &var_mbr_idx,
+	                                                      const Bitset &interpolation_qual);
 	void add_tess_level_input_to_interface_block(const std::string &ib_var_ref, SPIRType &ib_type, SPIRVariable &var);
 	void add_tess_level_input(const std::string &base_ref, const std::string &mbr_name, SPIRVariable &var);

-	void fix_up_interface_member_indices(spv::StorageClass storage, uint32_t ib_type_id);
+	void ensure_struct_members_valid_vecsizes(SPIRType &struct_type, uint32_t &location);
+	void fix_up_interface_member_indices(StorageClass storage, uint32_t ib_type_id);

 	void mark_location_as_used_by_shader(uint32_t location, const SPIRType &type,
-	                                     spv::StorageClass storage, bool fallback = false);
-	uint32_t ensure_correct_builtin_type(uint32_t type_id, spv::BuiltIn builtin);
+	                                     StorageClass storage, bool fallback = false);
+	uint32_t ensure_correct_builtin_type(uint32_t type_id, BuiltIn builtin);
 	uint32_t ensure_correct_input_type(uint32_t type_id, uint32_t location, uint32_t component,
 	                                   uint32_t num_components, bool strip_array);

@@ -981,6 +1073,9 @@ protected:

 	void fix_up_shader_inputs_outputs();

+	bool entry_point_is_vertex() const;
+	bool entry_point_returns_stage_output() const;
+	bool entry_point_requires_const_device_buffers() const;
 	std::string func_type_decl(SPIRType &type);
 	std::string entry_point_args_classic(bool append_comma);
 	std::string entry_point_args_argument_buffer(bool append_comma);
@@ -994,31 +1089,33 @@ protected:
 	std::string to_buffer_size_expression(uint32_t id);
 	bool is_sample_rate() const;
 	bool is_intersection_query() const;
-	bool is_direct_input_builtin(spv::BuiltIn builtin);
-	std::string builtin_qualifier(spv::BuiltIn builtin);
-	std::string builtin_type_decl(spv::BuiltIn builtin, uint32_t id = 0);
-	std::string built_in_func_arg(spv::BuiltIn builtin, bool prefix_comma);
+	bool is_direct_input_builtin(BuiltIn builtin);
+	std::string builtin_qualifier(BuiltIn builtin);
+	std::string builtin_type_decl(BuiltIn builtin, uint32_t id = 0);
+	std::string built_in_func_arg(BuiltIn builtin, bool prefix_comma);
 	std::string member_attribute_qualifier(const SPIRType &type, uint32_t index);
 	std::string member_location_attribute_qualifier(const SPIRType &type, uint32_t index);
 	std::string argument_decl(const SPIRFunction::Parameter &arg);
-	const char *descriptor_address_space(uint32_t id, spv::StorageClass storage, const char *plain_address_space) const;
+	const char *descriptor_address_space(uint32_t id, StorageClass storage, const char *plain_address_space) const;
 	std::string round_fp_tex_coords(std::string tex_coords, bool coord_is_fp);
 	uint32_t get_metal_resource_index(SPIRVariable &var, SPIRType::BaseType basetype, uint32_t plane = 0);
 	uint32_t get_member_location(uint32_t type_id, uint32_t index, uint32_t *comp = nullptr) const;
-	uint32_t get_or_allocate_builtin_input_member_location(spv::BuiltIn builtin,
+	uint32_t get_or_allocate_builtin_input_member_location(BuiltIn builtin,
 	                                                       uint32_t type_id, uint32_t index, uint32_t *comp = nullptr);
-	uint32_t get_or_allocate_builtin_output_member_location(spv::BuiltIn builtin,
+	uint32_t get_or_allocate_builtin_output_member_location(BuiltIn builtin,
 	                                                        uint32_t type_id, uint32_t index, uint32_t *comp = nullptr);

-	uint32_t get_physical_tess_level_array_size(spv::BuiltIn builtin) const;
+	uint32_t get_physical_tess_level_array_size(BuiltIn builtin) const;
+
+	uint32_t get_physical_type_id_stride(TypeID type_id) const override;

 	// MSL packing rules. These compute the effective packing rules as observed by the MSL compiler in the MSL output.
 	// These values can change depending on various extended decorations which control packing rules.
 	// We need to make these rules match up with SPIR-V declared rules.
-	uint32_t get_declared_type_size_msl(const SPIRType &type, bool packed, bool row_major) const;
-	uint32_t get_declared_type_array_stride_msl(const SPIRType &type, bool packed, bool row_major) const;
-	uint32_t get_declared_type_matrix_stride_msl(const SPIRType &type, bool packed, bool row_major) const;
-	uint32_t get_declared_type_alignment_msl(const SPIRType &type, bool packed, bool row_major) const;
+	uint32_t get_declared_type_size_msl(TypeID type_id, const SPIRType *special_type, bool packed, bool row_major) const;
+	uint32_t get_declared_type_array_stride_msl(TypeID type_id, const SPIRType *special_type, bool packed, bool row_major) const;
+	uint32_t get_declared_type_matrix_stride_msl(TypeID type_id, const SPIRType *special_type, bool packed, bool row_major) const;
+	uint32_t get_declared_type_alignment_msl(TypeID type_id, const SPIRType *special_type, bool packed, bool row_major) const;

 	uint32_t get_declared_struct_member_size_msl(const SPIRType &struct_type, uint32_t index) const;
 	uint32_t get_declared_struct_member_array_stride_msl(const SPIRType &struct_type, uint32_t index) const;
@@ -1030,11 +1127,10 @@ protected:
 	uint32_t get_declared_input_matrix_stride_msl(const SPIRType &struct_type, uint32_t index) const;
 	uint32_t get_declared_input_alignment_msl(const SPIRType &struct_type, uint32_t index) const;

-	const SPIRType &get_physical_member_type(const SPIRType &struct_type, uint32_t index) const;
+	TypeID get_physical_member_type_id(const SPIRType &struct_type, uint32_t index) const;
 	SPIRType get_presumed_input_type(const SPIRType &struct_type, uint32_t index) const;

-	uint32_t get_declared_struct_size_msl(const SPIRType &struct_type, bool ignore_alignment = false,
-	                                      bool ignore_padding = false) const;
+	uint32_t get_declared_struct_size_msl(const SPIRType &struct_type) const;

 	std::string to_component_argument(uint32_t id);
 	void align_struct(SPIRType &ib_type, std::unordered_set<uint32_t> &aligned_structs);
@@ -1042,8 +1138,12 @@ protected:
 	void mark_struct_members_packed(const SPIRType &type);
 	void ensure_member_packing_rules_msl(SPIRType &ib_type, uint32_t index);
 	bool validate_member_packing_rules_msl(const SPIRType &type, uint32_t index) const;
-	std::string get_argument_address_space(const SPIRVariable &argument);
+	std::string get_variable_address_space(const SPIRVariable &argument);
+	// Special case of get_variable_address_space which is only used for leaf functions.
+	std::string get_leaf_argument_address_space(const SPIRVariable &argument);
 	std::string get_type_address_space(const SPIRType &type, uint32_t id, bool argument = false);
+	bool decoration_flags_signal_volatile(const Bitset &flags) const;
+	bool decoration_flags_signal_coherent(const Bitset &flags) const;
 	const char *to_restrict(uint32_t id, bool space);
 	SPIRType &get_stage_in_struct_type();
 	SPIRType &get_stage_out_struct_type();
@@ -1052,15 +1152,15 @@ protected:
 	std::string get_tess_factor_struct_name();
 	SPIRType &get_uint_type();
 	uint32_t get_uint_type_id();
-	void emit_atomic_func_op(uint32_t result_type, uint32_t result_id, const char *op, spv::Op opcode,
+	void emit_atomic_func_op(uint32_t result_type, uint32_t result_id, const char *op, Op opcode,
 	                         uint32_t mem_order_1, uint32_t mem_order_2, bool has_mem_order_2, uint32_t op0, uint32_t op1 = 0,
 	                         bool op1_is_pointer = false, bool op1_is_literal = false, uint32_t op2 = 0);
 	const char *get_memory_order(uint32_t spv_mem_sem);
-	void add_pragma_line(const std::string &line);
+	void add_pragma_line(const std::string &line, bool recompile_on_unique);
 	void add_typedef_line(const std::string &line);
 	void emit_barrier(uint32_t id_exe_scope, uint32_t id_mem_scope, uint32_t id_mem_sem);
 	bool emit_array_copy(const char *expr, uint32_t lhs_id, uint32_t rhs_id,
-	                     spv::StorageClass lhs_storage, spv::StorageClass rhs_storage) override;
+	                     StorageClass lhs_storage, StorageClass rhs_storage) override;
 	void build_implicit_builtins();
 	uint32_t build_constant_uint_array_pointer();
 	void emit_entry_point_declarations() override;
@@ -1084,11 +1184,18 @@ protected:
 	uint32_t builtin_stage_input_size_id = 0;
 	uint32_t builtin_local_invocation_index_id = 0;
 	uint32_t builtin_workgroup_size_id = 0;
+	uint32_t builtin_mesh_primitive_indices_id = 0;
+	uint32_t builtin_mesh_sizes_id = 0;
+	uint32_t builtin_task_grid_id = 0;
+	uint32_t builtin_frag_depth_id = 0;
 	uint32_t swizzle_buffer_id = 0;
 	uint32_t buffer_size_buffer_id = 0;
 	uint32_t view_mask_buffer_id = 0;
+	uint32_t draw_index_buffer_id = 0;
 	uint32_t dynamic_offsets_buffer_id = 0;
 	uint32_t uint_type_id = 0;
+	uint32_t shared_uint_type_id = 0;
+	uint32_t meshlet_type_id = 0;
 	uint32_t argument_buffer_padding_buffer_type_id = 0;
 	uint32_t argument_buffer_padding_image_type_id = 0;
 	uint32_t argument_buffer_padding_sampler_type_id = 0;
@@ -1101,20 +1208,21 @@ protected:
 	void emit_store_statement(uint32_t lhs_expression, uint32_t rhs_expression) override;

 	void analyze_sampled_image_usage();
+	void analyze_workgroup_variables();

 	bool access_chain_needs_stage_io_builtin_translation(uint32_t base) override;
-	bool prepare_access_chain_for_scalar_access(std::string &expr, const SPIRType &type, spv::StorageClass storage,
+	bool prepare_access_chain_for_scalar_access(std::string &expr, const SPIRType &type, StorageClass storage,
 	                                            bool &is_packed) override;
 	void fix_up_interpolant_access_chain(const uint32_t *ops, uint32_t length);
-	void check_physical_type_cast(std::string &expr, const SPIRType *type, uint32_t physical_type) override;
+	bool check_physical_type_cast(std::string &expr, const SPIRType *type, uint32_t physical_type) override;

 	bool emit_tessellation_access_chain(const uint32_t *ops, uint32_t length);
 	bool emit_tessellation_io_load(uint32_t result_type, uint32_t id, uint32_t ptr);
 	bool is_out_of_bounds_tessellation_level(uint32_t id_lhs);

-	void ensure_builtin(spv::StorageClass storage, spv::BuiltIn builtin);
+	void ensure_builtin(StorageClass storage, BuiltIn builtin);

-	void mark_implicit_builtin(spv::StorageClass storage, spv::BuiltIn builtin, uint32_t id);
+	void mark_implicit_builtin(StorageClass storage, BuiltIn builtin, uint32_t id);

 	std::string convert_to_f32(const std::string &expr, uint32_t components);

@@ -1132,9 +1240,10 @@ protected:
 	std::unordered_map<uint32_t, uint32_t> fragment_output_components;
 	std::unordered_map<uint32_t, uint32_t> builtin_to_automatic_input_location;
 	std::unordered_map<uint32_t, uint32_t> builtin_to_automatic_output_location;
-	std::set<std::string> pragma_lines;
-	std::set<std::string> typedef_lines;
+	std::vector<std::string> pragma_lines;
+	std::vector<std::string> typedef_lines;
 	SmallVector<uint32_t> vars_needing_early_declaration;
+	std::unordered_set<uint32_t> constant_macro_ids;

 	std::unordered_map<StageSetBinding, std::pair<MSLResourceBinding, bool>, InternalHasher> resource_bindings;
 	std::unordered_map<StageSetBinding, uint32_t, InternalHasher> resource_arg_buff_idx_to_binding_number;
@@ -1153,6 +1262,8 @@ protected:
 	VariableID stage_out_ptr_var_id = 0;
 	VariableID tess_level_inner_var_id = 0;
 	VariableID tess_level_outer_var_id = 0;
+	VariableID mesh_out_per_vertex = 0;
+	VariableID mesh_out_per_primitive = 0;
 	VariableID stage_out_masked_builtin_type_id = 0;

 	// Handle HLSL-style 0-based vertex/instance index.
@@ -1172,15 +1283,22 @@ protected:
 	bool using_builtin_array() const;

 	bool is_rasterization_disabled = false;
+	bool has_descriptor_side_effects_buffer = false;
 	bool capture_output_to_buffer = false;
 	bool needs_swizzle_buffer_def = false;
 	bool used_swizzle_buffer = false;
 	bool added_builtin_tess_level = false;
+	bool needs_local_invocation_index = false;
 	bool needs_subgroup_invocation_id = false;
 	bool needs_subgroup_size = false;
 	bool needs_sample_id = false;
 	bool needs_helper_invocation = false;
+	bool needs_workgroup_zero_init = false;
+	bool needs_point_size_output = false;
+	bool writes_to_depth = false;
+	bool writes_to_point_size = false;
 	std::string qual_pos_var_name;
+	std::string qual_viewport_idx_var_name;
 	std::string stage_in_var_name = "in";
 	std::string stage_out_var_name = "out";
 	std::string patch_stage_in_var_name = "patchIn";
@@ -1196,7 +1314,7 @@ protected:
 	std::string patch_output_buffer_var_name = "spvPatchOut";
 	std::string tess_factor_buffer_var_name = "spvTessLevel";
 	std::string index_buffer_var_name = "spvIndices";
-	spv::Op previous_instruction_opcode = spv::OpNop;
+	Op previous_instruction_opcode = OpNop;

 	// Must be ordered since declaration is in a specific order.
 	std::map<uint32_t, MSLConstexprSampler> constexpr_samplers_by_id;
@@ -1204,7 +1322,6 @@ protected:
 	const MSLConstexprSampler *find_constexpr_sampler(uint32_t id) const;

 	std::unordered_set<uint32_t> buffers_requiring_array_length;
-	SmallVector<std::pair<uint32_t, uint32_t>> buffer_aliases_argument;
 	SmallVector<uint32_t> buffer_aliases_discrete;
 	std::unordered_set<uint32_t> atomic_image_vars_emulated; // Emulate texture2D atomic operations
 	std::unordered_set<uint32_t> pull_model_inputs;
@@ -1213,7 +1330,13 @@ protected:
 	SmallVector<SPIRVariable *> entry_point_bindings;

 	// Must be ordered since array is in a specific order.
-	std::map<SetBindingPair, std::pair<uint32_t, uint32_t>> buffers_requiring_dynamic_offset;
+	struct DynamicBuffer
+	{
+		uint32_t base_index;
+		uint32_t var_id;
+		std::string mbr_name;
+	};
+	std::map<SetBindingPair, DynamicBuffer> buffers_requiring_dynamic_offset;

 	SmallVector<uint32_t> disabled_frag_outputs;

@@ -1223,9 +1346,12 @@ protected:
 	uint32_t argument_buffer_discrete_mask = 0;
 	uint32_t argument_buffer_device_storage_mask = 0;

+	void emit_argument_buffer_aliased_descriptor(const SPIRVariable &aliased_var,
+	                                             const SPIRVariable &base_var);
+
 	void analyze_argument_buffers();
 	bool descriptor_set_is_argument_buffer(uint32_t desc_set) const;
-	MSLResourceBinding &get_argument_buffer_resource(uint32_t desc_set, uint32_t arg_idx);
+	const MSLResourceBinding &get_argument_buffer_resource(uint32_t desc_set, uint32_t arg_idx) const;
 	void add_argument_buffer_padding_buffer_type(SPIRType &struct_type, uint32_t &mbr_idx, uint32_t &arg_buff_index, MSLResourceBinding &rez_bind);
 	void add_argument_buffer_padding_image_type(SPIRType &struct_type, uint32_t &mbr_idx, uint32_t &arg_buff_index, MSLResourceBinding &rez_bind);
 	void add_argument_buffer_padding_sampler_type(SPIRType &struct_type, uint32_t &mbr_idx, uint32_t &arg_buff_index, MSLResourceBinding &rez_bind);
@@ -1237,17 +1363,17 @@ protected:
 	uint32_t build_msl_interpolant_type(uint32_t type_id, bool is_noperspective);

 	bool suppress_missing_prototypes = false;
+	bool suppress_incompatible_pointer_types_discard_qualifiers = false;
+	bool suppress_sometimes_unitialized = false;

 	void add_spv_func_and_recompile(SPVFuncImpl spv_func);

 	void activate_argument_buffer_resources();

 	bool type_is_msl_framebuffer_fetch(const SPIRType &type) const;
-	bool type_is_pointer(const SPIRType &type) const;
-	bool type_is_pointer_to_pointer(const SPIRType &type) const;
 	bool is_supported_argument_buffer_type(const SPIRType &type) const;

-	bool variable_storage_requires_stage_io(spv::StorageClass storage) const;
+	bool variable_storage_requires_stage_io(StorageClass storage) const;

 	bool needs_manual_helper_invocation_updates() const
 	{
@@ -1255,7 +1381,7 @@ protected:
 	}
 	bool needs_frag_discard_checks() const
 	{
-		return get_execution_model() == spv::ExecutionModelFragment && msl_options.supports_msl_version(2, 3) &&
+		return get_execution_model() == ExecutionModelFragment && msl_options.supports_msl_version(2, 3) &&
 		       msl_options.check_discarded_frag_stores && frag_shader_needs_discard_checks;
 	}

@@ -1265,40 +1391,41 @@ protected:
 	// OpcodeHandler that handles several MSL preprocessing operations.
 	struct OpCodePreprocessor : OpcodeHandler
 	{
-		OpCodePreprocessor(CompilerMSL &compiler_)
-		    : compiler(compiler_)
+		explicit OpCodePreprocessor(CompilerMSL &compiler_)
+		    : OpcodeHandler(compiler_), self(compiler_)
 		{
+			enable_result_types = true;
 		}

-		bool handle(spv::Op opcode, const uint32_t *args, uint32_t length) override;
-		CompilerMSL::SPVFuncImpl get_spv_func_impl(spv::Op opcode, const uint32_t *args);
+		bool handle(Op opcode, const uint32_t *args, uint32_t length) override;
+		CompilerMSL::SPVFuncImpl get_spv_func_impl(Op opcode, const uint32_t *args, uint32_t length);
 		void check_resource_write(uint32_t var_id);

-		CompilerMSL &compiler;
-		std::unordered_map<uint32_t, uint32_t> result_types;
+		CompilerMSL &self;
 		std::unordered_map<uint32_t, uint32_t> image_pointers_emulated; // Emulate texture2D atomic operations
 		bool suppress_missing_prototypes = false;
 		bool uses_atomics = false;
 		bool uses_image_write = false;
 		bool uses_buffer_write = false;
 		bool uses_discard = false;
+		bool needs_local_invocation_index = false;
 		bool needs_subgroup_invocation_id = false;
 		bool needs_subgroup_size = false;
 		bool needs_sample_id = false;
 		bool needs_helper_invocation = false;
+		bool uses_cooperative_matrix = false;
 	};

 	// OpcodeHandler that scans for uses of sampled images
 	struct SampledImageScanner : OpcodeHandler
 	{
-		SampledImageScanner(CompilerMSL &compiler_)
-		    : compiler(compiler_)
+		explicit SampledImageScanner(CompilerMSL &compiler_)
+		    : OpcodeHandler(compiler_), self(compiler_)
 		{
 		}

-		bool handle(spv::Op opcode, const uint32_t *args, uint32_t) override;
-
-		CompilerMSL &compiler;
+		CompilerMSL &self;
+		bool handle(Op opcode, const uint32_t *args, uint32_t) override;
 	};

 	// Sorts the members of a SPIRType and associated Meta info based on a settable sorting
@@ -22,10 +22,11 @@
 */

 #include "spirv_parser.hpp"
+#include "NonSemanticShaderDebugInfo100.h"
 #include <assert.h>

 using namespace std;
-using namespace spv;
+using namespace SPIRV_CROSS_SPV_HEADER_NAMESPACE;

 namespace SPIRV_CROSS_NAMESPACE
 {
@@ -43,7 +44,7 @@ static bool decoration_is_string(Decoration decoration)
 {
 	switch (decoration)
 	{
-	case DecorationHlslSemanticGOOGLE:
+	case DecorationUserSemantic:
 		return true;

 	default:
@@ -137,12 +138,42 @@ void Parser::parse()
 	}
 	forward_pointer_fixups.clear();

+	for (auto &source : ir.sources)
+	{
+		auto cmp = [](const ParsedIR::Source::Marker &a, const ParsedIR::Source::Marker &b) {
+			return a.line < b.line;
+		};
+
+		std::sort(source.line_markers.begin(), source.line_markers.end(), cmp);
+	}
+
 	if (current_function)
 		SPIRV_CROSS_THROW("Function was not terminated.");
 	if (current_block)
 		SPIRV_CROSS_THROW("Block was not terminated.");
+
+	// Now that all definitions are bound to a kind, we can filter the library
+	// exports and populate the exported functions.
+	for (uint32_t id : ir.library_exports)
+	{
+		if (ir.ids[id].get_type() == TypeFunction)
+			ir.library_exported_functions.push_back(id);
+	}
+
 	if (ir.default_entry_point == 0)
-		SPIRV_CROSS_THROW("There is no entry point in the SPIR-V module.");
+	{
+		if (ir.library_exported_functions.empty())
+			SPIRV_CROSS_THROW("There is no entry point in the SPIR-V module.");
+
+		// No OpEntryPoint, but the module exports functions. Treat as a library
+		// module: designate the first exported function as the default entry
+		// point so analyses keyed on default_entry_point can run.
+		ir.is_library_module = true;
+		ir.default_entry_point = ir.library_exported_functions.front();
+		auto &name = ir.get_name(ir.default_entry_point);
+		ir.entry_points.insert(std::make_pair(ir.default_entry_point,
+		                                      SPIREntryPoint(ir.default_entry_point, ExecutionModelGLCompute, name)));
+	}
 }

 const uint32_t *Parser::stream(const Instruction &instr) const
@@ -194,7 +225,6 @@ void Parser::parse(const Instruction &instruction)

 	switch (op)
 	{
-	case OpSourceContinued:
 	case OpSourceExtension:
 	case OpNop:
 	case OpModuleProcessed:
@@ -213,38 +243,53 @@ void Parser::parse(const Instruction &instruction)

 	case OpSource:
 	{
-		auto lang = static_cast<SourceLanguage>(ops[0]);
-		switch (lang)
+		ir.sources.emplace_back();
+		auto &source = ir.sources.back();
+		source.lang = static_cast<SourceLanguage>(ops[0]);
+
+		switch (source.lang)
 		{
 		case SourceLanguageESSL:
-			ir.source.es = true;
-			ir.source.version = ops[1];
-			ir.source.known = true;
-			ir.source.hlsl = false;
+			source.es = true;
+			source.version = ops[1];
+			source.known = true;
+			source.hlsl = false;
 			break;

 		case SourceLanguageGLSL:
-			ir.source.es = false;
-			ir.source.version = ops[1];
-			ir.source.known = true;
-			ir.source.hlsl = false;
+			source.es = false;
+			source.version = ops[1];
+			source.known = true;
+			source.hlsl = false;
 			break;

 		case SourceLanguageHLSL:
 			// For purposes of cross-compiling, this is GLSL 450.
-			ir.source.es = false;
-			ir.source.version = 450;
-			ir.source.known = true;
-			ir.source.hlsl = true;
+			source.es = false;
+			source.version = 450;
+			source.known = true;
+			source.hlsl = true;
 			break;

 		default:
-			ir.source.known = false;
+			source.known = false;
 			break;
 		}
+
+		if (length >= 3)
+			source.file_id = ops[2];
+
+		if (length >= 4)
+			source.source = extract_string(ir.spirv, instruction.offset + 3);
+
 		break;
 	}

+	case OpSourceContinued:
+		if (!ir.sources.empty())
+			ir.sources.back().source += extract_string(ir.spirv, instruction.offset);
+		break;
+
 	case OpUndef:
 	{
 		uint32_t result_type = ops[0];
@@ -305,6 +350,7 @@ void Parser::parse(const Instruction &instruction)
 	}

 	case OpExtInst:
+	case OpExtInstWithForwardRefsKHR:
 	{
 		// The SPIR-V debug information extended instructions might come at global scope.
 		if (current_block)
@@ -317,6 +363,72 @@ void Parser::parse(const Instruction &instruction)
 					ir.load_type_width.insert({ ops[1], type->width });
 			}
 		}
+
+		if (op == OpExtInst && length > 4)
+		{
+			// Don't want to deal with ForwardRefs here.
+			auto &ext = get<SPIRExtension>(ops[2]);
+			if (ext.ext == SPIRExtension::NonSemanticShaderDebugInfo)
+			{
+				const auto instr = ops[3];
+				if (instr == NonSemanticShaderDebugInfo100DebugSource)
+				{
+					set<SPIRString>(ops[1], get<SPIRString>(ops[4]).str);
+
+					ir.sources.emplace_back();
+					auto &source = ir.sources.back();
+					source.file_id = ops[4];
+					source.define_id = ops[1];
+					if (length >= 6)
+						source.source = ir.get<SPIRString>(ops[5]).str;
+				}
+				else if (instr == NonSemanticShaderDebugInfo100DebugSourceContinued)
+				{
+					if (length < 5)
+						SPIRV_CROSS_THROW("Invalid arguments for ShaderDebugInfo100DebugSourceContinued");
+					if (!ir.sources.empty())
+						ir.sources.back().source += ir.get<SPIRString>(ops[4]).str;
+				}
+				else if (instr == NonSemanticShaderDebugInfo100DebugLine)
+				{
+					if (length < 9)
+						SPIRV_CROSS_THROW("Invalid arguments for ShaderDebugInfo100DebugLine");
+					auto source_id = ops[4];
+					auto line_start = ir.get<SPIRConstant>(ops[5]).scalar_i32();
+					auto col_start = ir.get<SPIRConstant>(ops[7]).scalar_i32();
+
+					for (auto &source : ir.sources)
+					{
+						if (source.define_id != source_id)
+							continue;
+
+						source.line_markers.emplace_back();
+						auto &marker = source.line_markers.back();
+						marker.line = line_start;
+						marker.col = col_start;
+						marker.offset = instruction.offset - 1;
+						marker.function_id = current_function ? current_function->self : ID(0);
+						marker.block_id = current_block ? current_block->self : ID(0);
+						break;
+					}
+				}
+				else if (instr == NonSemanticShaderDebugInfo100DebugLocalVariable)
+				{
+					if (length < 11)
+						SPIRV_CROSS_THROW("Invalid arguments for ShaderDebugInfo100DebugLocalVariable");
+					auto &lvar = set<SPIRDebugLocalVariable>(ops[1]);
+					lvar.name_id = ops[4];
+				}
+				else if (instr == NonSemanticShaderDebugInfo100DebugDeclare)
+				{
+					if (length < 7)
+						SPIRV_CROSS_THROW("Invalid arguments for ShaderDebugInfo100DebugDeclare");
+					auto &lvar = get<SPIRDebugLocalVariable>(ops[4]);
+					auto &var = get<SPIRVariable>(ops[5]);
+					var.debug_local_variables.push_back(lvar.self);
+				}
+			}
+		}
 		break;
 	}

@@ -368,6 +480,30 @@ void Parser::parse(const Instruction &instruction)
 			execution.output_primitives = ops[2];
 			break;

+		case ExecutionModeSignedZeroInfNanPreserve:
+			switch (ops[2])
+			{
+			case 8:
+				execution.signed_zero_inf_nan_preserve_8 = true;
+				break;
+
+			case 16:
+				execution.signed_zero_inf_nan_preserve_16 = true;
+				break;
+
+			case 32:
+				execution.signed_zero_inf_nan_preserve_32 = true;
+				break;
+
+			case 64:
+				execution.signed_zero_inf_nan_preserve_64 = true;
+				break;
+
+			default:
+				SPIRV_CROSS_THROW("Invalid bit-width for SignedZeroInfNanPreserve.");
+			}
+			break;
+
 		default:
 			break;
 		}
@@ -380,13 +516,21 @@ void Parser::parse(const Instruction &instruction)
 		auto mode = static_cast<ExecutionMode>(ops[1]);
 		execution.flags.set(mode);

-		if (mode == ExecutionModeLocalSizeId)
+		switch (mode)
 		{
+		case ExecutionModeLocalSizeId:
 			execution.workgroup_size.id_x = ops[2];
 			execution.workgroup_size.id_y = ops[3];
 			execution.workgroup_size.id_z = ops[4];
-		}
+			break;

+		case ExecutionModeFPFastMathDefault:
+			execution.fp_fast_math_defaults[ops[2]] = ops[3];
+			break;
+
+		default:
+			break;
+		}
 		break;
 	}

@@ -481,6 +625,19 @@ void Parser::parse(const Instruction &instruction)
 		else
 			ir.set_decoration(id, decoration);

+		// Track exported functions so we can compile library modules that have no OpEntryPoint.
+		// LinkageAttributes layout: literal-string (variable words) followed by LinkageType.
+		if (decoration == DecorationLinkageAttributes && length >= 4 &&
+		    static_cast<LinkageType>(ops[length - 1]) == LinkageTypeExport)
+		{
+			ir.library_exports.push_back(id);
+
+			// If OpName was stripped (e.g. by spirv-opt --strip-debug), fall back
+			// to the linkage name so the emitted function keeps its export name.
+			if (ir.get_name(id).empty())
+				ir.set_name(id, extract_string(ir.spirv, instruction.offset + 2));
+		}
+
 		break;
 	}

@@ -504,6 +661,12 @@ void Parser::parse(const Instruction &instruction)
 		break;
 	}

+	// MemberDecorateIdEXT only applies to OffsetIdEXT when descriptors are packed in structs.
+	// This is currently unsupported and will fail in compilation.
+	// Pass it through in case someone just needs reflection.
+	case OpMemberDecorateIdEXT:
+		break;
+
 	case OpMemberDecorateStringGOOGLE:
 	{
 		uint32_t id = ops[0];
@@ -536,12 +699,37 @@ void Parser::parse(const Instruction &instruction)
 		uint32_t id = ops[0];
 		uint32_t width = ops[1];
 		auto &type = set<SPIRType>(id, op);
+
+		if (width != 16 && width != 8 && length > 2)
+			SPIRV_CROSS_THROW("Unrecognized FP encoding mode for OpTypeFloat.");
+
 		if (width == 64)
 			type.basetype = SPIRType::Double;
 		else if (width == 32)
 			type.basetype = SPIRType::Float;
 		else if (width == 16)
-			type.basetype = SPIRType::Half;
+		{
+			if (length > 2)
+			{
+				if (ops[2] == FPEncodingBFloat16KHR)
+					type.basetype = SPIRType::BFloat16;
+				else
+					SPIRV_CROSS_THROW("Unrecognized encoding for OpTypeFloat 16.");
+			}
+			else
+				type.basetype = SPIRType::Half;
+		}
+		else if (width == 8)
+		{
+			if (length < 2)
+				SPIRV_CROSS_THROW("Missing encoding for OpTypeFloat 8.");
+			else if (ops[2] == FPEncodingFloat8E4M3EXT)
+				type.basetype = SPIRType::FloatE4M3;
+			else if (ops[2] == FPEncodingFloat8E5M2EXT)
+				type.basetype = SPIRType::FloatE5M2;
+			else
+				SPIRV_CROSS_THROW("Invalid encoding for OpTypeFloat 8.");
+		}
 		else
 			SPIRV_CROSS_THROW("Unrecognized bit-width of floating point type.");
 		type.width = width;
@@ -592,6 +780,40 @@ void Parser::parse(const Instruction &instruction)
 		break;
 	}

+	case OpTypeCooperativeMatrixKHR:
+	{
+		uint32_t id = ops[0];
+		auto &base = get<SPIRType>(ops[1]);
+		auto &matrixbase = set<SPIRType>(id, base);
+
+		matrixbase.op = op;
+		matrixbase.ext.cooperative.scope_id = ops[2];
+		matrixbase.ext.cooperative.rows_id = ops[3];
+		matrixbase.ext.cooperative.columns_id = ops[4];
+		matrixbase.ext.cooperative.use_id = ops[5];
+		matrixbase.self = id;
+		matrixbase.parent_type = ops[1];
+		break;
+	}
+
+	case OpTypeCooperativeVectorNV:
+	{
+		uint32_t id = ops[0];
+		auto &type = set<SPIRType>(id, op);
+
+		type.basetype = SPIRType::CoopVecNV;
+		type.op = op;
+		type.ext.coopVecNV.component_type_id = ops[1];
+		type.ext.coopVecNV.component_count_id = ops[2];
+		type.parent_type = ops[1];
+
+		// CoopVec-Nv can be used with integer operations like SMax where
+		// where spirv-opt does explicit checks on integer bitwidth
+		auto component_type = get<SPIRType>(type.ext.coopVecNV.component_type_id);
+		type.width = component_type.width;
+		break;
+	}
+
 	case OpTypeArray:
 	{
 		uint32_t id = ops[0];
@@ -677,6 +899,7 @@ void Parser::parse(const Instruction &instruction)
 		break;
 	}

+	case OpTypeUntypedPointerKHR:
 	case OpTypePointer:
 	{
 		uint32_t id = ops[0];
@@ -684,7 +907,7 @@ void Parser::parse(const Instruction &instruction)
 		// Very rarely, we might receive a FunctionPrototype here.
 		// We won't be able to compile it, but we shouldn't crash when parsing.
 		// We should be able to reflect.
-		auto *base = maybe_get<SPIRType>(ops[2]);
+		auto *base = op == OpTypePointer ? maybe_get<SPIRType>(ops[2]) : nullptr;
 		auto &ptrbase = set<SPIRType>(id, op);

 		if (base)
@@ -703,7 +926,10 @@ void Parser::parse(const Instruction &instruction)
 		if (base && base->forward_pointer)
 			forward_pointer_fixups.push_back({ id, ops[2] });

-		ptrbase.parent_type = ops[2];
+		if (op == OpTypePointer)
+			ptrbase.parent_type = ops[2];
+		else
+			ptrbase.basetype = SPIRType::Void;

 		// Do NOT set ptrbase.self!
 		break;
@@ -789,6 +1015,20 @@ void Parser::parse(const Instruction &instruction)
 		break;
 	}

+	case OpTypeTensorARM:
+	{
+		uint32_t id = ops[0];
+		auto &type = set<SPIRType>(id, op);
+		type.basetype = SPIRType::Tensor;
+		type.ext.tensor = {};
+		type.ext.tensor.type = ops[1];
+		if (length >= 3)
+			type.ext.tensor.rank = ops[2];
+		if (length >= 4)
+			type.ext.tensor.shape = ops[3];
+		break;
+	}
+
 	// Variable declaration
 	// All variables are essentially pointers with a storage qualifier.
 	case OpVariable:
@@ -809,6 +1049,27 @@ void Parser::parse(const Instruction &instruction)
 		break;
 	}

+	case OpUntypedVariableKHR:
+	{
+		uint32_t type = ops[0];
+		uint32_t id = ops[1];
+		auto storage = static_cast<StorageClass>(ops[2]);
+		uint32_t data_type = length >= 4 ? ops[3] : 0;
+		uint32_t initializer = length >= 5 ? ops[4] : 0;
+
+		if (storage == StorageClassFunction)
+		{
+			if (!current_function)
+				SPIRV_CROSS_THROW("No function currently in scope");
+			current_function->add_local_variable(id);
+		}
+
+		auto &v = set<SPIRVariable>(id, type, storage, initializer);
+		v.untyped = true;
+		v.untyped_alloca_type = data_type;
+		break;
+	}
+
 	// OpPhi
 	// OpPhi is a fairly magical opcode.
 	// It selects temporary variables based on which parent block we *came from*.
@@ -825,7 +1086,7 @@ void Parser::parse(const Instruction &instruction)
 		uint32_t id = ops[1];

 		// Instead of a temporary, create a new function-wide temporary with this ID instead.
-		auto &var = set<SPIRVariable>(id, result_type, spv::StorageClassFunction);
+		auto &var = set<SPIRVariable>(id, result_type, StorageClassFunction);
 		var.phi_variable = true;

 		current_function->add_local_variable(id);
@@ -835,17 +1096,27 @@ void Parser::parse(const Instruction &instruction)
 		break;
 	}

-		// Constants
+	// Constants
 	case OpSpecConstant:
 	case OpConstant:
+	case OpConstantCompositeReplicateEXT:
+	case OpSpecConstantCompositeReplicateEXT:
 	{
 		uint32_t id = ops[1];
 		auto &type = get<SPIRType>(ops[0]);
-
-		if (type.width > 32)
-			set<SPIRConstant>(id, ops[0], ops[2] | (uint64_t(ops[3]) << 32), op == OpSpecConstant);
+		if (op == OpConstantCompositeReplicateEXT || op == OpSpecConstantCompositeReplicateEXT)
+		{
+			auto subconstant = uint32_t(ops[2]);
+			set<SPIRConstant>(id, ops[0], &subconstant, 1, op == OpSpecConstantCompositeReplicateEXT, true);
+		}
 		else
-			set<SPIRConstant>(id, ops[0], ops[2], op == OpSpecConstant);
+		{
+
+			if (type.width > 32)
+				set<SPIRConstant>(id, ops[0], ops[2] | (uint64_t(ops[3]) << 32), op == OpSpecConstant);
+			else
+				set<SPIRConstant>(id, ops[0], ops[2], op == OpSpecConstant);
+		}
 		break;
 	}

@@ -929,6 +1200,24 @@ void Parser::parse(const Instruction &instruction)
 		break;
 	}

+	case OpConstantSizeOfEXT:
+	{
+		uint32_t id = ops[1];
+		uint32_t type = ops[0];
+		auto &c = set<SPIRConstant>(id, type);
+		c.size_of_type = ops[2];
+		break;
+	}
+
+	case OpTypeBufferEXT:
+	{
+		uint32_t type = ops[0];
+		auto &t = set<SPIRType>(type, OpTypeBufferEXT);
+		t.basetype = SPIRType::DescriptorHeapBuffer;
+		t.ext.descriptor_heap_buffer.storage = static_cast<StorageClass>(ops[1]);
+		break;
+	}
+
 	// Functions
 	case OpFunction:
 	{
@@ -1249,6 +1538,24 @@ void Parser::parse(const Instruction &instruction)
 				current_function->entry_line.line_literal = ops[1];
 			}
 		}
+
+		uint32_t file = ops[0];
+		uint32_t line = ops[1];
+
+		for (auto &source : ir.sources)
+		{
+			if (source.file_id == file)
+			{
+				source.line_markers.emplace_back();
+				auto &marker = source.line_markers.back();
+				marker.line = line;
+				marker.offset = instruction.offset - 1;
+				marker.function_id = current_function ? current_function->self : ID(0);
+				marker.block_id = current_block ? current_block->self : ID(0);
+				break;
+			}
+		}
+
 		break;
 	}

@@ -25,7 +25,7 @@
 #include "spirv_glsl.hpp"
 #include <iomanip>

-using namespace spv;
+using namespace SPIRV_CROSS_SPV_HEADER_NAMESPACE;
 using namespace SPIRV_CROSS_NAMESPACE;
 using namespace std;

@@ -449,7 +449,7 @@ void CompilerReflection::emit_type_member_qualifiers(const SPIRType &type, uint3
 	}
 }

-string CompilerReflection::execution_model_to_str(spv::ExecutionModel model)
+string CompilerReflection::execution_model_to_str(ExecutionModel model)
 {
 	switch (model)
 	{
@@ -477,6 +477,12 @@ string CompilerReflection::execution_model_to_str(spv::ExecutionModel model)
 		return "rmiss";
 	case ExecutionModelCallableNV:
 		return "rcall";
+	case ExecutionModelMeshNV:
+	case ExecutionModelMeshEXT:
+		return "mesh";
+	case ExecutionModelTaskNV:
+	case ExecutionModelTaskEXT:
+		return "task";
 	default:
 		return "???";
 	}
@@ -504,7 +510,9 @@ void CompilerReflection::emit_entry_points()
 			json_stream->begin_json_object();
 			json_stream->emit_json_key_value("name", e.name);
 			json_stream->emit_json_key_value("mode", execution_model_to_str(e.execution_model));
-			if (e.execution_model == ExecutionModelGLCompute)
+			if (e.execution_model == ExecutionModelGLCompute || e.execution_model == ExecutionModelMeshEXT ||
+			    e.execution_model == ExecutionModelMeshNV || e.execution_model == ExecutionModelTaskEXT ||
+			    e.execution_model == ExecutionModelTaskNV)
 			{
 				const auto &spv_entry = get_entry_point(e.name, e.execution_model);

@@ -547,6 +555,7 @@ void CompilerReflection::emit_resources()
 	emit_resources("push_constants", res.push_constant_buffers);
 	emit_resources("counters", res.atomic_counters);
 	emit_resources("acceleration_structures", res.acceleration_structures);
+	emit_resources("tensors", res.tensors);
 }

 void CompilerReflection::emit_resources(const char *tag, const SmallVector<Resource> &resources)
@@ -633,6 +642,12 @@ void CompilerReflection::emit_resources(const char *tag, const SmallVector<Resou
 			                                 get_decoration(res.id, DecorationInputAttachmentIndex));
 		if (mask.get(DecorationOffset))
 			json_stream->emit_json_key_value("offset", get_decoration(res.id, DecorationOffset));
+		if (mask.get(DecorationWeightTextureQCOM))
+			json_stream->emit_json_key_value("WeightTextureQCOM", get_decoration(res.id, DecorationWeightTextureQCOM));
+		if (mask.get(DecorationBlockMatchTextureQCOM))
+			json_stream->emit_json_key_value("BlockMatchTextureQCOM", get_decoration(res.id, DecorationBlockMatchTextureQCOM));
+		if (mask.get(DecorationBlockMatchSamplerQCOM))
+			json_stream->emit_json_key_value("BlockMatchSamplerQCOM", get_decoration(res.id, DecorationBlockMatchSamplerQCOM));

 		// For images, the type itself adds a layout qualifer.
 		// Only emit the format for storage images.
@@ -34,6 +34,7 @@ class Stream;

 namespace SPIRV_CROSS_NAMESPACE
 {
+using namespace SPIRV_CROSS_SPV_HEADER_NAMESPACE;
 class CompilerReflection : public CompilerGLSL
 {
 	using Parent = CompilerGLSL;
@@ -67,7 +68,7 @@ public:
 	std::string compile() override;

 private:
-	static std::string execution_model_to_str(spv::ExecutionModel model);
+	static std::string execution_model_to_str(ExecutionModel model);

 	void emit_entry_points();
 	void emit_types();