1	/*
2	* Copyright 2015-2021 Arm Limited
3	* SPDX-License-Identifier: Apache-2.0 OR MIT
4	*
5	* Licensed under the Apache License, Version 2.0 (the "License");
6	* you may not use this file except in compliance with the License.
7	* You may obtain a copy of the License at
8	*
9	* http://www.apache.org/licenses/LICENSE-2.0
10	*
11	* Unless required by applicable law or agreed to in writing, software
12	* distributed under the License is distributed on an "AS IS" BASIS,
13	* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
14	* See the License for the specific language governing permissions and
15	* limitations under the License.
16	*/
17
18	/*
19	* At your option, you may choose to accept this material under either:
20	* 1. The Apache License, Version 2.0, found at <http://www.apache.org/licenses/LICENSE-2.0>, or
21	* 2. The MIT License, found at <http://opensource.org/licenses/MIT>.
22	*/
23
24	#ifndef SPIRV_CROSS_COMMON_HPP
25	#define SPIRV_CROSS_COMMON_HPP
26
27	#ifndef SPV_ENABLE_UTILITY_CODE
28	#define SPV_ENABLE_UTILITY_CODE
29	#endif
30	#include "spirv.hpp"
31
32	#include "spirv_cross_containers.hpp"
33	#include "spirv_cross_error_handling.hpp"
34	#include <functional>
35
36	// A bit crude, but allows projects which embed SPIRV-Cross statically to
37	// effectively hide all the symbols from other projects.
38	// There is a case where we have:
39	// - Project A links against SPIRV-Cross statically.
40	// - Project A links against Project B statically.
41	// - Project B links against SPIRV-Cross statically (might be a different version).
42	// This leads to a conflict with extremely bizarre results.
43	// By overriding the namespace in one of the project builds, we can work around this.
44	// If SPIRV-Cross is embedded in dynamic libraries,
45	// prefer using -fvisibility=hidden on GCC/Clang instead.
46	#ifdef SPIRV_CROSS_NAMESPACE_OVERRIDE
47	#define SPIRV_CROSS_NAMESPACE SPIRV_CROSS_NAMESPACE_OVERRIDE
48	#else
49	#define SPIRV_CROSS_NAMESPACE spirv_cross
50	#endif
51
52	namespace SPIRV_CROSS_NAMESPACE
53	{
54	namespace inner
55	{
56	template <typename T>
57	void join_helper(StringStream<> &stream, T &&t)
58	{
59	stream << std::forward<T>(t);
60	}
61
62	template <typename T, typename... Ts>
63	void join_helper(StringStream<> &stream, T &&t, Ts &&... ts)
64	{
65	stream << std::forward<T>(t);
66	join_helper(stream, std::forward<Ts>(ts)...);
67	}
68	} // namespace inner
69
70	class Bitset
71	{
72	public:
73	Bitset() = default;
74	explicit inline Bitset(uint64_t lower_)
75	: lower(lower_)
76	{
77	}
78
79	inline bool get(uint32_t bit) const
80	{
81	if (bit < 64)
82	return (lower & (1ull << bit)) != 0;
83	else
84	return higher.count(x: bit) != 0;
85	}
86
87	inline void set(uint32_t bit)
88	{
89	if (bit < 64)
90	lower |= 1ull << bit;
91	else
92	higher.insert(x: bit);
93	}
94
95	inline void clear(uint32_t bit)
96	{
97	if (bit < 64)
98	lower &= ~(1ull << bit);
99	else
100	higher.erase(x: bit);
101	}
102
103	inline uint64_t get_lower() const
104	{
105	return lower;
106	}
107
108	inline void reset()
109	{
110	lower = 0;
111	higher.clear();
112	}
113
114	inline void merge_and(const Bitset &other)
115	{
116	lower &= other.lower;
117	std::unordered_set<uint32_t> tmp_set;
118	for (auto &v : higher)
119	if (other.higher.count(x: v) != 0)
120	tmp_set.insert(x: v);
121	higher = std::move(tmp_set);
122	}
123
124	inline void merge_or(const Bitset &other)
125	{
126	lower |= other.lower;
127	for (auto &v : other.higher)
128	higher.insert(x: v);
129	}
130
131	inline bool operator==(const Bitset &other) const
132	{
133	if (lower != other.lower)
134	return false;
135
136	if (higher.size() != other.higher.size())
137	return false;
138
139	for (auto &v : higher)
140	if (other.higher.count(x: v) == 0)
141	return false;
142
143	return true;
144	}
145
146	inline bool operator!=(const Bitset &other) const
147	{
148	return !(*this == other);
149	}
150
151	template <typename Op>
152	void for_each_bit(const Op &op) const
153	{
154	// TODO: Add ctz-based iteration.
155	for (uint32_t i = 0; i < 64; i++)
156	{
157	if (lower & (1ull << i))
158	op(i);
159	}
160
161	if (higher.empty())
162	return;
163
164	// Need to enforce an order here for reproducible results,
165	// but hitting this path should happen extremely rarely, so having this slow path is fine.
166	SmallVector<uint32_t> bits;
167	bits.reserve(count: higher.size());
168	for (auto &v : higher)
169	bits.push_back(t: v);
170	std::sort(first: std::begin(cont&: bits), last: std::end(cont&: bits));
171
172	for (auto &v : bits)
173	op(v);
174	}
175
176	inline bool empty() const
177	{
178	return lower == 0 && higher.empty();
179	}
180
181	private:
182	// The most common bits to set are all lower than 64,
183	// so optimize for this case. Bits spilling outside 64 go into a slower data structure.
184	// In almost all cases, higher data structure will not be used.
185	uint64_t lower = 0;
186	std::unordered_set<uint32_t> higher;
187	};
188
189	// Helper template to avoid lots of nasty string temporary munging.
190	template <typename... Ts>
191	std::string join(Ts &&... ts)
192	{
193	StringStream<> stream;
194	inner::join_helper(stream, std::forward<Ts>(ts)...);
195	return stream.str();
196	}
197
198	inline std::string merge(const SmallVector<std::string> &list, const char *between = ", ")
199	{
200	StringStream<> stream;
201	for (auto &elem : list)
202	{
203	stream << elem;
204	if (&elem != &list.back())
205	stream << between;
206	}
207	return stream.str();
208	}
209
210	// Make sure we don't accidentally call this with float or doubles with SFINAE.
211	// Have to use the radix-aware overload.
212	template <typename T, typename std::enable_if<!std::is_floating_point<T>::value, int>::type = 0>
213	inline std::string convert_to_string(const T &t)
214	{
215	return std::to_string(t);
216	}
217
218	static inline std::string convert_to_string(int32_t value)
219	{
220	// INT_MIN is ... special on some backends. If we use a decimal literal, and negate it, we
221	// could accidentally promote the literal to long first, then negate.
222	// To workaround it, emit int(0x80000000) instead.
223	if (value == (std::numeric_limits<int32_t>::min)())
224	return "int(0x80000000)";
225	else
226	return std::to_string(val: value);
227	}
228
229	static inline std::string convert_to_string(int64_t value, const std::string &int64_type, bool long_long_literal_suffix)
230	{
231	// INT64_MIN is ... special on some backends.
232	// If we use a decimal literal, and negate it, we might overflow the representable numbers.
233	// To workaround it, emit int(0x80000000) instead.
234	if (value == (std::numeric_limits<int64_t>::min)())
235	return join(ts: int64_type, ts: "(0x8000000000000000u", ts: (long_long_literal_suffix ? "ll" : "l"), ts: ")");
236	else
237	return std::to_string(val: value) + (long_long_literal_suffix ? "ll" : "l");
238	}
239
240	// Allow implementations to set a convenient standard precision
241	#ifndef SPIRV_CROSS_FLT_FMT
242	#define SPIRV_CROSS_FLT_FMT "%.32g"
243	#endif
244
245	// Disable sprintf and strcat warnings.
246	// We cannot rely on snprintf and family existing because, ..., MSVC.
247	#if defined(__clang__) || defined(__GNUC__)
248	#pragma GCC diagnostic push
249	#pragma GCC diagnostic ignored "-Wdeprecated-declarations"
250	#elif defined(_MSC_VER)
251	#pragma warning(push)
252	#pragma warning(disable : 4996)
253	#endif
254
255	static inline void fixup_radix_point(char *str, char radix_point)
256	{
257	// Setting locales is a very risky business in multi-threaded program,
258	// so just fixup locales instead. We only need to care about the radix point.
259	if (radix_point != '.')
260	{
261	while (*str != '\0')
262	{
263	if (*str == radix_point)
264	*str = '.';
265	str++;
266	}
267	}
268	}
269
270	inline std::string convert_to_string(float t, char locale_radix_point)
271	{
272	// std::to_string for floating point values is broken.
273	// Fallback to something more sane.
274	char buf[64];
275	sprintf(s: buf, SPIRV_CROSS_FLT_FMT, t);
276	fixup_radix_point(str: buf, radix_point: locale_radix_point);
277
278	// Ensure that the literal is float.
279	if (!strchr(s: buf, c: '.') && !strchr(s: buf, c: 'e'))
280	strcat(dest: buf, src: ".0");
281	return buf;
282	}
283
284	inline std::string convert_to_string(double t, char locale_radix_point)
285	{
286	// std::to_string for floating point values is broken.
287	// Fallback to something more sane.
288	char buf[64];
289	sprintf(s: buf, SPIRV_CROSS_FLT_FMT, t);
290	fixup_radix_point(str: buf, radix_point: locale_radix_point);
291
292	// Ensure that the literal is float.
293	if (!strchr(s: buf, c: '.') && !strchr(s: buf, c: 'e'))
294	strcat(dest: buf, src: ".0");
295	return buf;
296	}
297
298	#if defined(__clang__) || defined(__GNUC__)
299	#pragma GCC diagnostic pop
300	#elif defined(_MSC_VER)
301	#pragma warning(pop)
302	#endif
303
304	class FloatFormatter
305	{
306	public:
307	virtual ~FloatFormatter() = default;
308	virtual std::string format_float(float value) = 0;
309	virtual std::string format_double(double value) = 0;
310	};
311
312	template <typename T>
313	struct ValueSaver
314	{
315	explicit ValueSaver(T &current_)
316	: current(current_)
317	, saved(current_)
318	{
319	}
320
321	void release()
322	{
323	current = saved;
324	}
325
326	~ValueSaver()
327	{
328	release();
329	}
330
331	T &current;
332	T saved;
333	};
334
335	struct Instruction
336	{
337	uint16_t op = 0;
338	uint16_t count = 0;
339	// If offset is 0 (not a valid offset into the instruction stream),
340	// we have an instruction stream which is embedded in the object.
341	uint32_t offset = 0;
342	uint32_t length = 0;
343
344	inline bool is_embedded() const
345	{
346	return offset == 0;
347	}
348	};
349
350	struct EmbeddedInstruction : Instruction
351	{
352	SmallVector<uint32_t> ops;
353	};
354
355	enum Types
356	{
357	TypeNone,
358	TypeType,
359	TypeVariable,
360	TypeConstant,
361	TypeFunction,
362	TypeFunctionPrototype,
363	TypeBlock,
364	TypeExtension,
365	TypeExpression,
366	TypeConstantOp,
367	TypeCombinedImageSampler,
368	TypeAccessChain,
369	TypeUndef,
370	TypeString,
371	TypeCount
372	};
373
374	template <Types type>
375	class TypedID;
376
377	template <>
378	class TypedID<TypeNone>
379	{
380	public:
381	TypedID() = default;
382	TypedID(uint32_t id_)
383	: id(id_)
384	{
385	}
386
387	template <Types U>
388	TypedID(const TypedID<U> &other)
389	{
390	*this = other;
391	}
392
393	template <Types U>
394	TypedID &operator=(const TypedID<U> &other)
395	{
396	id = uint32_t(other);
397	return *this;
398	}
399
400	// Implicit conversion to u32 is desired here.
401	// As long as we block implicit conversion between TypedID<A> and TypedID<B> we're good.
402	operator uint32_t() const
403	{
404	return id;
405	}
406
407	template <Types U>
408	operator TypedID<U>() const
409	{
410	return TypedID<U>(*this);
411	}
412
413	private:
414	uint32_t id = 0;
415	};
416
417	template <Types type>
418	class TypedID
419	{
420	public:
421	TypedID() = default;
422	TypedID(uint32_t id_)
423	: id(id_)
424	{
425	}
426
427	explicit TypedID(const TypedID<TypeNone> &other)
428	: id(uint32_t(other))
429	{
430	}
431
432	operator uint32_t() const
433	{
434	return id;
435	}
436
437	private:
438	uint32_t id = 0;
439	};
440
441	using VariableID = TypedID<TypeVariable>;
442	using TypeID = TypedID<TypeType>;
443	using ConstantID = TypedID<TypeConstant>;
444	using FunctionID = TypedID<TypeFunction>;
445	using BlockID = TypedID<TypeBlock>;
446	using ID = TypedID<TypeNone>;
447
448	// Helper for Variant interface.
449	struct IVariant
450	{
451	virtual ~IVariant() = default;
452	virtual IVariant *clone(ObjectPoolBase *pool) = 0;
453	ID self = 0;
454
455	protected:
456	IVariant() = default;
457	IVariant(const IVariant&) = default;
458	IVariant &operator=(const IVariant&) = default;
459	};
460
461	#define SPIRV_CROSS_DECLARE_CLONE(T) \
462	IVariant *clone(ObjectPoolBase *pool) override \
463	{ \
464	return static_cast<ObjectPool<T> *>(pool)->allocate(*this); \
465	}
466
467	struct SPIRUndef : IVariant
468	{
469	enum
470	{
471	type = TypeUndef
472	};
473
474	explicit SPIRUndef(TypeID basetype_)
475	: basetype(basetype_)
476	{
477	}
478	TypeID basetype;
479
480	SPIRV_CROSS_DECLARE_CLONE(SPIRUndef)
481	};
482
483	struct SPIRString : IVariant
484	{
485	enum
486	{
487	type = TypeString
488	};
489
490	explicit SPIRString(std::string str_)
491	: str(std::move(str_))
492	{
493	}
494
495	std::string str;
496
497	SPIRV_CROSS_DECLARE_CLONE(SPIRString)
498	};
499
500	// This type is only used by backends which need to access the combined image and sampler IDs separately after
501	// the OpSampledImage opcode.
502	struct SPIRCombinedImageSampler : IVariant
503	{
504	enum
505	{
506	type = TypeCombinedImageSampler
507	};
508	SPIRCombinedImageSampler(TypeID type_, VariableID image_, VariableID sampler_)
509	: combined_type(type_)
510	, image(image_)
511	, sampler(sampler_)
512	{
513	}
514	TypeID combined_type;
515	VariableID image;
516	VariableID sampler;
517
518	SPIRV_CROSS_DECLARE_CLONE(SPIRCombinedImageSampler)
519	};
520
521	struct SPIRConstantOp : IVariant
522	{
523	enum
524	{
525	type = TypeConstantOp
526	};
527
528	SPIRConstantOp(TypeID result_type, spv::Op op, const uint32_t *args, uint32_t length)
529	: opcode(op)
530	, basetype(result_type)
531	{
532	arguments.reserve(count: length);
533	for (uint32_t i = 0; i < length; i++)
534	arguments.push_back(t: args[i]);
535	}
536
537	spv::Op opcode;
538	SmallVector<uint32_t> arguments;
539	TypeID basetype;
540
541	SPIRV_CROSS_DECLARE_CLONE(SPIRConstantOp)
542	};
543
544	struct SPIRType : IVariant
545	{
546	enum
547	{
548	type = TypeType
549	};
550
551	spv::Op op = spv::Op::OpNop;
552	explicit SPIRType(spv::Op op_) : op(op_) {}
553
554	enum BaseType
555	{
556	Unknown,
557	Void,
558	Boolean,
559	SByte,
560	UByte,
561	Short,
562	UShort,
563	Int,
564	UInt,
565	Int64,
566	UInt64,
567	AtomicCounter,
568	Half,
569	Float,
570	Double,
571	Struct,
572	Image,
573	SampledImage,
574	Sampler,
575	AccelerationStructure,
576	RayQuery,
577
578	// Keep internal types at the end.
579	ControlPointArray,
580	Interpolant,
581	Char
582	};
583
584	// Scalar/vector/matrix support.
585	BaseType basetype = Unknown;
586	uint32_t width = 0;
587	uint32_t vecsize = 1;
588	uint32_t columns = 1;
589
590	// Arrays, support array of arrays by having a vector of array sizes.
591	SmallVector<uint32_t> array;
592
593	// Array elements can be either specialization constants or specialization ops.
594	// This array determines how to interpret the array size.
595	// If an element is true, the element is a literal,
596	// otherwise, it's an expression, which must be resolved on demand.
597	// The actual size is not really known until runtime.
598	SmallVector<bool> array_size_literal;
599
600	// Pointers
601	// Keep track of how many pointer layers we have.
602	uint32_t pointer_depth = 0;
603	bool pointer = false;
604	bool forward_pointer = false;
605
606	spv::StorageClass storage = spv::StorageClassGeneric;
607
608	SmallVector<TypeID> member_types;
609
610	// If member order has been rewritten to handle certain scenarios with Offset,
611	// allow codegen to rewrite the index.
612	SmallVector<uint32_t> member_type_index_redirection;
613
614	struct ImageType
615	{
616	TypeID type;
617	spv::Dim dim;
618	bool depth;
619	bool arrayed;
620	bool ms;
621	uint32_t sampled;
622	spv::ImageFormat format;
623	spv::AccessQualifier access;
624	} image = {};
625
626	// Structs can be declared multiple times if they are used as part of interface blocks.
627	// We want to detect this so that we only emit the struct definition once.
628	// Since we cannot rely on OpName to be equal, we need to figure out aliases.
629	TypeID type_alias = 0;
630
631	// Denotes the type which this type is based on.
632	// Allows the backend to traverse how a complex type is built up during access chains.
633	TypeID parent_type = 0;
634
635	// Used in backends to avoid emitting members with conflicting names.
636	std::unordered_set<std::string> member_name_cache;
637
638	SPIRV_CROSS_DECLARE_CLONE(SPIRType)
639	};
640
641	struct SPIRExtension : IVariant
642	{
643	enum
644	{
645	type = TypeExtension
646	};
647
648	enum Extension
649	{
650	Unsupported,
651	GLSL,
652	SPV_debug_info,
653	SPV_AMD_shader_ballot,
654	SPV_AMD_shader_explicit_vertex_parameter,
655	SPV_AMD_shader_trinary_minmax,
656	SPV_AMD_gcn_shader,
657	NonSemanticDebugPrintf,
658	NonSemanticShaderDebugInfo,
659	NonSemanticGeneric
660	};
661
662	explicit SPIRExtension(Extension ext_)
663	: ext(ext_)
664	{
665	}
666
667	Extension ext;
668	SPIRV_CROSS_DECLARE_CLONE(SPIRExtension)
669	};
670
671	// SPIREntryPoint is not a variant since its IDs are used to decorate OpFunction,
672	// so in order to avoid conflicts, we can't stick them in the ids array.
673	struct SPIREntryPoint
674	{
675	SPIREntryPoint(FunctionID self_, spv::ExecutionModel execution_model, const std::string &entry_name)
676	: self(self_)
677	, name(entry_name)
678	, orig_name(entry_name)
679	, model(execution_model)
680	{
681	}
682	SPIREntryPoint() = default;
683
684	FunctionID self = 0;
685	std::string name;
686	std::string orig_name;
687	SmallVector<VariableID> interface_variables;
688
689	Bitset flags;
690	struct WorkgroupSize
691	{
692	uint32_t x = 0, y = 0, z = 0;
693	uint32_t id_x = 0, id_y = 0, id_z = 0;
694	uint32_t constant = 0; // Workgroup size can be expressed as a constant/spec-constant instead.
695	} workgroup_size;
696	uint32_t invocations = 0;
697	uint32_t output_vertices = 0;
698	uint32_t output_primitives = 0;
699	spv::ExecutionModel model = spv::ExecutionModelMax;
700	bool geometry_passthrough = false;
701	};
702
703	struct SPIRExpression : IVariant
704	{
705	enum
706	{
707	type = TypeExpression
708	};
709
710	// Only created by the backend target to avoid creating tons of temporaries.
711	SPIRExpression(std::string expr, TypeID expression_type_, bool immutable_)
712	: expression(std::move(expr))
713	, expression_type(expression_type_)
714	, immutable(immutable_)
715	{
716	}
717
718	// If non-zero, prepend expression with to_expression(base_expression).
719	// Used in amortizing multiple calls to to_expression()
720	// where in certain cases that would quickly force a temporary when not needed.
721	ID base_expression = 0;
722
723	std::string expression;
724	TypeID expression_type = 0;
725
726	// If this expression is a forwarded load,
727	// allow us to reference the original variable.
728	ID loaded_from = 0;
729
730	// If this expression will never change, we can avoid lots of temporaries
731	// in high level source.
732	// An expression being immutable can be speculative,
733	// it is assumed that this is true almost always.
734	bool immutable = false;
735
736	// Before use, this expression must be transposed.
737	// This is needed for targets which don't support row_major layouts.
738	bool need_transpose = false;
739
740	// Whether or not this is an access chain expression.
741	bool access_chain = false;
742
743	// Whether or not gl_MeshVerticesEXT[].gl_Position (as a whole or .y) is referenced
744	bool access_meshlet_position_y = false;
745
746	// A list of expressions which this expression depends on.
747	SmallVector<ID> expression_dependencies;
748
749	// By reading this expression, we implicitly read these expressions as well.
750	// Used by access chain Store and Load since we read multiple expressions in this case.
751	SmallVector<ID> implied_read_expressions;
752
753	// The expression was emitted at a certain scope. Lets us track when an expression read means multiple reads.
754	uint32_t emitted_loop_level = 0;
755
756	SPIRV_CROSS_DECLARE_CLONE(SPIRExpression)
757	};
758
759	struct SPIRFunctionPrototype : IVariant
760	{
761	enum
762	{
763	type = TypeFunctionPrototype
764	};
765
766	explicit SPIRFunctionPrototype(TypeID return_type_)
767	: return_type(return_type_)
768	{
769	}
770
771	TypeID return_type;
772	SmallVector<uint32_t> parameter_types;
773
774	SPIRV_CROSS_DECLARE_CLONE(SPIRFunctionPrototype)
775	};
776
777	struct SPIRBlock : IVariant
778	{
779	enum
780	{
781	type = TypeBlock
782	};
783
784	enum Terminator
785	{
786	Unknown,
787	Direct, // Emit next block directly without a particular condition.
788
789	Select, // Block ends with an if/else block.
790	MultiSelect, // Block ends with switch statement.
791
792	Return, // Block ends with return.
793	Unreachable, // Noop
794	Kill, // Discard
795	IgnoreIntersection, // Ray Tracing
796	TerminateRay, // Ray Tracing
797	EmitMeshTasks // Mesh shaders
798	};
799
800	enum Merge
801	{
802	MergeNone,
803	MergeLoop,
804	MergeSelection
805	};
806
807	enum Hints
808	{
809	HintNone,
810	HintUnroll,
811	HintDontUnroll,
812	HintFlatten,
813	HintDontFlatten
814	};
815
816	enum Method
817	{
818	MergeToSelectForLoop,
819	MergeToDirectForLoop,
820	MergeToSelectContinueForLoop
821	};
822
823	enum ContinueBlockType
824	{
825	ContinueNone,
826
827	// Continue block is branchless and has at least one instruction.
828	ForLoop,
829
830	// Noop continue block.
831	WhileLoop,
832
833	// Continue block is conditional.
834	DoWhileLoop,
835
836	// Highly unlikely that anything will use this,
837	// since it is really awkward/impossible to express in GLSL.
838	ComplexLoop
839	};
840
841	enum : uint32_t
842	{
843	NoDominator = 0xffffffffu
844	};
845
846	Terminator terminator = Unknown;
847	Merge merge = MergeNone;
848	Hints hint = HintNone;
849	BlockID next_block = 0;
850	BlockID merge_block = 0;
851	BlockID continue_block = 0;
852
853	ID return_value = 0; // If 0, return nothing (void).
854	ID condition = 0;
855	BlockID true_block = 0;
856	BlockID false_block = 0;
857	BlockID default_block = 0;
858
859	// If terminator is EmitMeshTasksEXT.
860	struct
861	{
862	ID groups[3];
863	ID payload;
864	} mesh = {};
865
866	SmallVector<Instruction> ops;
867
868	struct Phi
869	{
870	ID local_variable; // flush local variable ...
871	BlockID parent; // If we're in from_block and want to branch into this block ...
872	VariableID function_variable; // to this function-global "phi" variable first.
873	};
874
875	// Before entering this block flush out local variables to magical "phi" variables.
876	SmallVector<Phi> phi_variables;
877
878	// Declare these temporaries before beginning the block.
879	// Used for handling complex continue blocks which have side effects.
880	SmallVector<std::pair<TypeID, ID>> declare_temporary;
881
882	// Declare these temporaries, but only conditionally if this block turns out to be
883	// a complex loop header.
884	SmallVector<std::pair<TypeID, ID>> potential_declare_temporary;
885
886	struct Case
887	{
888	uint64_t value;
889	BlockID block;
890	};
891	SmallVector<Case> cases_32bit;
892	SmallVector<Case> cases_64bit;
893
894	// If we have tried to optimize code for this block but failed,
895	// keep track of this.
896	bool disable_block_optimization = false;
897
898	// If the continue block is complex, fallback to "dumb" for loops.
899	bool complex_continue = false;
900
901	// Do we need a ladder variable to defer breaking out of a loop construct after a switch block?
902	bool need_ladder_break = false;
903
904	// If marked, we have explicitly handled Phi from this block, so skip any flushes related to that on a branch.
905	// Used to handle an edge case with switch and case-label fallthrough where fall-through writes to Phi.
906	BlockID ignore_phi_from_block = 0;
907
908	// The dominating block which this block might be within.
909	// Used in continue; blocks to determine if we really need to write continue.
910	BlockID loop_dominator = 0;
911
912	// All access to these variables are dominated by this block,
913	// so before branching anywhere we need to make sure that we declare these variables.
914	SmallVector<VariableID> dominated_variables;
915
916	// These are variables which should be declared in a for loop header, if we
917	// fail to use a classic for-loop,
918	// we remove these variables, and fall back to regular variables outside the loop.
919	SmallVector<VariableID> loop_variables;
920
921	// Some expressions are control-flow dependent, i.e. any instruction which relies on derivatives or
922	// sub-group-like operations.
923	// Make sure that we only use these expressions in the original block.
924	SmallVector<ID> invalidate_expressions;
925
926	SPIRV_CROSS_DECLARE_CLONE(SPIRBlock)
927	};
928
929	struct SPIRFunction : IVariant
930	{
931	enum
932	{
933	type = TypeFunction
934	};
935
936	SPIRFunction(TypeID return_type_, TypeID function_type_)
937	: return_type(return_type_)
938	, function_type(function_type_)
939	{
940	}
941
942	struct Parameter
943	{
944	TypeID type;
945	ID id;
946	uint32_t read_count;
947	uint32_t write_count;
948
949	// Set to true if this parameter aliases a global variable,
950	// used mostly in Metal where global variables
951	// have to be passed down to functions as regular arguments.
952	// However, for this kind of variable, we should not care about
953	// read and write counts as access to the function arguments
954	// is not local to the function in question.
955	bool alias_global_variable;
956	};
957
958	// When calling a function, and we're remapping separate image samplers,
959	// resolve these arguments into combined image samplers and pass them
960	// as additional arguments in this order.
961	// It gets more complicated as functions can pull in their own globals
962	// and combine them with parameters,
963	// so we need to distinguish if something is local parameter index
964	// or a global ID.
965	struct CombinedImageSamplerParameter
966	{
967	VariableID id;
968	VariableID image_id;
969	VariableID sampler_id;
970	bool global_image;
971	bool global_sampler;
972	bool depth;
973	};
974
975	TypeID return_type;
976	TypeID function_type;
977	SmallVector<Parameter> arguments;
978
979	// Can be used by backends to add magic arguments.
980	// Currently used by combined image/sampler implementation.
981
982	SmallVector<Parameter> shadow_arguments;
983	SmallVector<VariableID> local_variables;
984	BlockID entry_block = 0;
985	SmallVector<BlockID> blocks;
986	SmallVector<CombinedImageSamplerParameter> combined_parameters;
987
988	struct EntryLine
989	{
990	uint32_t file_id = 0;
991	uint32_t line_literal = 0;
992	};
993	EntryLine entry_line;
994
995	void add_local_variable(VariableID id)
996	{
997	local_variables.push_back(t: id);
998	}
999
1000	void add_parameter(TypeID parameter_type, ID id, bool alias_global_variable = false)
1001	{
1002	// Arguments are read-only until proven otherwise.
1003	arguments.push_back(t: { .type: parameter_type, .id: id, .read_count: 0u, .write_count: 0u, .alias_global_variable: alias_global_variable });
1004	}
1005
1006	// Hooks to be run when the function returns.
1007	// Mostly used for lowering internal data structures onto flattened structures.
1008	// Need to defer this, because they might rely on things which change during compilation.
1009	// Intentionally not a small vector, this one is rare, and std::function can be large.
1010	Vector<std::function<void()>> fixup_hooks_out;
1011
1012	// Hooks to be run when the function begins.
1013	// Mostly used for populating internal data structures from flattened structures.
1014	// Need to defer this, because they might rely on things which change during compilation.
1015	// Intentionally not a small vector, this one is rare, and std::function can be large.
1016	Vector<std::function<void()>> fixup_hooks_in;
1017
1018	// On function entry, make sure to copy a constant array into thread addr space to work around
1019	// the case where we are passing a constant array by value to a function on backends which do not
1020	// consider arrays value types.
1021	SmallVector<ID> constant_arrays_needed_on_stack;
1022
1023	bool active = false;
1024	bool flush_undeclared = true;
1025	bool do_combined_parameters = true;
1026
1027	SPIRV_CROSS_DECLARE_CLONE(SPIRFunction)
1028	};
1029
1030	struct SPIRAccessChain : IVariant
1031	{
1032	enum
1033	{
1034	type = TypeAccessChain
1035	};
1036
1037	SPIRAccessChain(TypeID basetype_, spv::StorageClass storage_, std::string base_, std::string dynamic_index_,
1038	int32_t static_index_)
1039	: basetype(basetype_)
1040	, storage(storage_)
1041	, base(std::move(base_))
1042	, dynamic_index(std::move(dynamic_index_))
1043	, static_index(static_index_)
1044	{
1045	}
1046
1047	// The access chain represents an offset into a buffer.
1048	// Some backends need more complicated handling of access chains to be able to use buffers, like HLSL
1049	// which has no usable buffer type ala GLSL SSBOs.
1050	// StructuredBuffer is too limited, so our only option is to deal with ByteAddressBuffer which works with raw addresses.
1051
1052	TypeID basetype;
1053	spv::StorageClass storage;
1054	std::string base;
1055	std::string dynamic_index;
1056	int32_t static_index;
1057
1058	VariableID loaded_from = 0;
1059	uint32_t matrix_stride = 0;
1060	uint32_t array_stride = 0;
1061	bool row_major_matrix = false;
1062	bool immutable = false;
1063
1064	// By reading this expression, we implicitly read these expressions as well.
1065	// Used by access chain Store and Load since we read multiple expressions in this case.
1066	SmallVector<ID> implied_read_expressions;
1067
1068	SPIRV_CROSS_DECLARE_CLONE(SPIRAccessChain)
1069	};
1070
1071	struct SPIRVariable : IVariant
1072	{
1073	enum
1074	{
1075	type = TypeVariable
1076	};
1077
1078	SPIRVariable() = default;
1079	SPIRVariable(TypeID basetype_, spv::StorageClass storage_, ID initializer_ = 0, VariableID basevariable_ = 0)
1080	: basetype(basetype_)
1081	, storage(storage_)
1082	, initializer(initializer_)
1083	, basevariable(basevariable_)
1084	{
1085	}
1086
1087	TypeID basetype = 0;
1088	spv::StorageClass storage = spv::StorageClassGeneric;
1089	uint32_t decoration = 0;
1090	ID initializer = 0;
1091	VariableID basevariable = 0;
1092
1093	SmallVector<uint32_t> dereference_chain;
1094	bool compat_builtin = false;
1095
1096	// If a variable is shadowed, we only statically assign to it
1097	// and never actually emit a statement for it.
1098	// When we read the variable as an expression, just forward
1099	// shadowed_id as the expression.
1100	bool statically_assigned = false;
1101	ID static_expression = 0;
1102
1103	// Temporaries which can remain forwarded as long as this variable is not modified.
1104	SmallVector<ID> dependees;
1105
1106	bool deferred_declaration = false;
1107	bool phi_variable = false;
1108
1109	// Used to deal with Phi variable flushes. See flush_phi().
1110	bool allocate_temporary_copy = false;
1111
1112	bool remapped_variable = false;
1113	uint32_t remapped_components = 0;
1114
1115	// The block which dominates all access to this variable.
1116	BlockID dominator = 0;
1117	// If true, this variable is a loop variable, when accessing the variable
1118	// outside a loop,
1119	// we should statically forward it.
1120	bool loop_variable = false;
1121	// Set to true while we're inside the for loop.
1122	bool loop_variable_enable = false;
1123
1124	// Used to find global LUTs
1125	bool is_written_to = false;
1126
1127	SPIRFunction::Parameter *parameter = nullptr;
1128
1129	SPIRV_CROSS_DECLARE_CLONE(SPIRVariable)
1130	};
1131
1132	struct SPIRConstant : IVariant
1133	{
1134	enum
1135	{
1136	type = TypeConstant
1137	};
1138
1139	union Constant
1140	{
1141	uint32_t u32;
1142	int32_t i32;
1143	float f32;
1144
1145	uint64_t u64;
1146	int64_t i64;
1147	double f64;
1148	};
1149
1150	struct ConstantVector
1151	{
1152	Constant r[4];
1153	// If != 0, this element is a specialization constant, and we should keep track of it as such.
1154	ID id[4];
1155	uint32_t vecsize = 1;
1156
1157	ConstantVector()
1158	{
1159	memset(s: r, c: 0, n: sizeof(r));
1160	}
1161	};
1162
1163	struct ConstantMatrix
1164	{
1165	ConstantVector c[4];
1166	// If != 0, this column is a specialization constant, and we should keep track of it as such.
1167	ID id[4];
1168	uint32_t columns = 1;
1169	};
1170
1171	static inline float f16_to_f32(uint16_t u16_value)
1172	{
1173	// Based on the GLM implementation.
1174	int s = (u16_value >> 15) & 0x1;
1175	int e = (u16_value >> 10) & 0x1f;
1176	int m = (u16_value >> 0) & 0x3ff;
1177
1178	union
1179	{
1180	float f32;
1181	uint32_t u32;
1182	} u;
1183
1184	if (e == 0)
1185	{
1186	if (m == 0)
1187	{
1188	u.u32 = uint32_t(s) << 31;
1189	return u.f32;
1190	}
1191	else
1192	{
1193	while ((m & 0x400) == 0)
1194	{
1195	m <<= 1;
1196	e--;
1197	}
1198
1199	e++;
1200	m &= ~0x400;
1201	}
1202	}
1203	else if (e == 31)
1204	{
1205	if (m == 0)
1206	{
1207	u.u32 = (uint32_t(s) << 31) | 0x7f800000u;
1208	return u.f32;
1209	}
1210	else
1211	{
1212	u.u32 = (uint32_t(s) << 31) | 0x7f800000u | (m << 13);
1213	return u.f32;
1214	}
1215	}
1216
1217	e += 127 - 15;
1218	m <<= 13;
1219	u.u32 = (uint32_t(s) << 31) | (e << 23) | m;
1220	return u.f32;
1221	}
1222
1223	inline uint32_t specialization_constant_id(uint32_t col, uint32_t row) const
1224	{
1225	return m.c[col].id[row];
1226	}
1227
1228	inline uint32_t specialization_constant_id(uint32_t col) const
1229	{
1230	return m.id[col];
1231	}
1232
1233	inline uint32_t scalar(uint32_t col = 0, uint32_t row = 0) const
1234	{
1235	return m.c[col].r[row].u32;
1236	}
1237
1238	inline int16_t scalar_i16(uint32_t col = 0, uint32_t row = 0) const
1239	{
1240	return int16_t(m.c[col].r[row].u32 & 0xffffu);
1241	}
1242
1243	inline uint16_t scalar_u16(uint32_t col = 0, uint32_t row = 0) const
1244	{
1245	return uint16_t(m.c[col].r[row].u32 & 0xffffu);
1246	}
1247
1248	inline int8_t scalar_i8(uint32_t col = 0, uint32_t row = 0) const
1249	{
1250	return int8_t(m.c[col].r[row].u32 & 0xffu);
1251	}
1252
1253	inline uint8_t scalar_u8(uint32_t col = 0, uint32_t row = 0) const
1254	{
1255	return uint8_t(m.c[col].r[row].u32 & 0xffu);
1256	}
1257
1258	inline float scalar_f16(uint32_t col = 0, uint32_t row = 0) const
1259	{
1260	return f16_to_f32(u16_value: scalar_u16(col, row));
1261	}
1262
1263	inline float scalar_f32(uint32_t col = 0, uint32_t row = 0) const
1264	{
1265	return m.c[col].r[row].f32;
1266	}
1267
1268	inline int32_t scalar_i32(uint32_t col = 0, uint32_t row = 0) const
1269	{
1270	return m.c[col].r[row].i32;
1271	}
1272
1273	inline double scalar_f64(uint32_t col = 0, uint32_t row = 0) const
1274	{
1275	return m.c[col].r[row].f64;
1276	}
1277
1278	inline int64_t scalar_i64(uint32_t col = 0, uint32_t row = 0) const
1279	{
1280	return m.c[col].r[row].i64;
1281	}
1282
1283	inline uint64_t scalar_u64(uint32_t col = 0, uint32_t row = 0) const
1284	{
1285	return m.c[col].r[row].u64;
1286	}
1287
1288	inline const ConstantVector &vector() const
1289	{
1290	return m.c[0];
1291	}
1292
1293	inline uint32_t vector_size() const
1294	{
1295	return m.c[0].vecsize;
1296	}
1297
1298	inline uint32_t columns() const
1299	{
1300	return m.columns;
1301	}
1302
1303	inline void make_null(const SPIRType &constant_type_)
1304	{
1305	m = {};
1306	m.columns = constant_type_.columns;
1307	for (auto &c : m.c)
1308	c.vecsize = constant_type_.vecsize;
1309	}
1310
1311	inline bool constant_is_null() const
1312	{
1313	if (specialization)
1314	return false;
1315	if (!subconstants.empty())
1316	return false;
1317
1318	for (uint32_t col = 0; col < columns(); col++)
1319	for (uint32_t row = 0; row < vector_size(); row++)
1320	if (scalar_u64(col, row) != 0)
1321	return false;
1322
1323	return true;
1324	}
1325
1326	explicit SPIRConstant(uint32_t constant_type_)
1327	: constant_type(constant_type_)
1328	{
1329	}
1330
1331	SPIRConstant() = default;
1332
1333	SPIRConstant(TypeID constant_type_, const uint32_t *elements, uint32_t num_elements, bool specialized)
1334	: constant_type(constant_type_)
1335	, specialization(specialized)
1336	{
1337	subconstants.reserve(count: num_elements);
1338	for (uint32_t i = 0; i < num_elements; i++)
1339	subconstants.push_back(t: elements[i]);
1340	specialization = specialized;
1341	}
1342
1343	// Construct scalar (32-bit).
1344	SPIRConstant(TypeID constant_type_, uint32_t v0, bool specialized)
1345	: constant_type(constant_type_)
1346	, specialization(specialized)
1347	{
1348	m.c[0].r[0].u32 = v0;
1349	m.c[0].vecsize = 1;
1350	m.columns = 1;
1351	}
1352
1353	// Construct scalar (64-bit).
1354	SPIRConstant(TypeID constant_type_, uint64_t v0, bool specialized)
1355	: constant_type(constant_type_)
1356	, specialization(specialized)
1357	{
1358	m.c[0].r[0].u64 = v0;
1359	m.c[0].vecsize = 1;
1360	m.columns = 1;
1361	}
1362
1363	// Construct vectors and matrices.
1364	SPIRConstant(TypeID constant_type_, const SPIRConstant *const *vector_elements, uint32_t num_elements,
1365	bool specialized)
1366	: constant_type(constant_type_)
1367	, specialization(specialized)
1368	{
1369	bool matrix = vector_elements[0]->m.c[0].vecsize > 1;
1370
1371	if (matrix)
1372	{
1373	m.columns = num_elements;
1374
1375	for (uint32_t i = 0; i < num_elements; i++)
1376	{
1377	m.c[i] = vector_elements[i]->m.c[0];
1378	if (vector_elements[i]->specialization)
1379	m.id[i] = vector_elements[i]->self;
1380	}
1381	}
1382	else
1383	{
1384	m.c[0].vecsize = num_elements;
1385	m.columns = 1;
1386
1387	for (uint32_t i = 0; i < num_elements; i++)
1388	{
1389	m.c[0].r[i] = vector_elements[i]->m.c[0].r[0];
1390	if (vector_elements[i]->specialization)
1391	m.c[0].id[i] = vector_elements[i]->self;
1392	}
1393	}
1394	}
1395
1396	TypeID constant_type = 0;
1397	ConstantMatrix m;
1398
1399	// If this constant is a specialization constant (i.e. created with OpSpecConstant*).
1400	bool specialization = false;
1401	// If this constant is used as an array length which creates specialization restrictions on some backends.
1402	bool is_used_as_array_length = false;
1403
1404	// If true, this is a LUT, and should always be declared in the outer scope.
1405	bool is_used_as_lut = false;
1406
1407	// For composites which are constant arrays, etc.
1408	SmallVector<ConstantID> subconstants;
1409
1410	// Non-Vulkan GLSL, HLSL and sometimes MSL emits defines for each specialization constant,
1411	// and uses them to initialize the constant. This allows the user
1412	// to still be able to specialize the value by supplying corresponding
1413	// preprocessor directives before compiling the shader.
1414	std::string specialization_constant_macro_name;
1415
1416	SPIRV_CROSS_DECLARE_CLONE(SPIRConstant)
1417	};
1418
1419	// Variants have a very specific allocation scheme.
1420	struct ObjectPoolGroup
1421	{
1422	std::unique_ptr<ObjectPoolBase> pools[TypeCount];
1423	};
1424
1425	class Variant
1426	{
1427	public:
1428	explicit Variant(ObjectPoolGroup *group_)
1429	: group(group_)
1430	{
1431	}
1432
1433	~Variant()
1434	{
1435	if (holder)
1436	group->pools[type]->deallocate_opaque(ptr: holder);
1437	}
1438
1439	// Marking custom move constructor as noexcept is important.
1440	Variant(Variant &&other) SPIRV_CROSS_NOEXCEPT
1441	{
1442	*this = std::move(other);
1443	}
1444
1445	// We cannot copy from other variant without our own pool group.
1446	// Have to explicitly copy.
1447	Variant(const Variant &variant) = delete;
1448
1449	// Marking custom move constructor as noexcept is important.
1450	Variant &operator=(Variant &&other) SPIRV_CROSS_NOEXCEPT
1451	{
1452	if (this != &other)
1453	{
1454	if (holder)
1455	group->pools[type]->deallocate_opaque(ptr: holder);
1456	holder = other.holder;
1457	group = other.group;
1458	type = other.type;
1459	allow_type_rewrite = other.allow_type_rewrite;
1460
1461	other.holder = nullptr;
1462	other.type = TypeNone;
1463	}
1464	return *this;
1465	}
1466
1467	// This copy/clone should only be called in the Compiler constructor.
1468	// If this is called inside ::compile(), we invalidate any references we took higher in the stack.
1469	// This should never happen.
1470	Variant &operator=(const Variant &other)
1471	{
1472	//#define SPIRV_CROSS_COPY_CONSTRUCTOR_SANITIZE
1473	#ifdef SPIRV_CROSS_COPY_CONSTRUCTOR_SANITIZE
1474	abort();
1475	#endif
1476	if (this != &other)
1477	{
1478	if (holder)
1479	group->pools[type]->deallocate_opaque(ptr: holder);
1480
1481	if (other.holder)
1482	holder = other.holder->clone(pool: group->pools[other.type].get());
1483	else
1484	holder = nullptr;
1485
1486	type = other.type;
1487	allow_type_rewrite = other.allow_type_rewrite;
1488	}
1489	return *this;
1490	}
1491
1492	void set(IVariant *val, Types new_type)
1493	{
1494	if (holder)
1495	group->pools[type]->deallocate_opaque(ptr: holder);
1496	holder = nullptr;
1497
1498	if (!allow_type_rewrite && type != TypeNone && type != new_type)
1499	{
1500	if (val)
1501	group->pools[new_type]->deallocate_opaque(ptr: val);
1502	SPIRV_CROSS_THROW("Overwriting a variant with new type.");
1503	}
1504
1505	holder = val;
1506	type = new_type;
1507	allow_type_rewrite = false;
1508	}
1509
1510	template <typename T, typename... Ts>
1511	T *allocate_and_set(Types new_type, Ts &&... ts)
1512	{
1513	T *val = static_cast<ObjectPool<T> &>(*group->pools[new_type]).allocate(std::forward<Ts>(ts)...);
1514	set(val, new_type);
1515	return val;
1516	}
1517
1518	template <typename T>
1519	T &get()
1520	{
1521	if (!holder)
1522	SPIRV_CROSS_THROW("nullptr");
1523	if (static_cast<Types>(T::type) != type)
1524	SPIRV_CROSS_THROW("Bad cast");
1525	return *static_cast<T *>(holder);
1526	}
1527
1528	template <typename T>
1529	const T &get() const
1530	{
1531	if (!holder)
1532	SPIRV_CROSS_THROW("nullptr");
1533	if (static_cast<Types>(T::type) != type)
1534	SPIRV_CROSS_THROW("Bad cast");
1535	return *static_cast<const T *>(holder);
1536	}
1537
1538	Types get_type() const
1539	{
1540	return type;
1541	}
1542
1543	ID get_id() const
1544	{
1545	return holder ? holder->self : ID(0);
1546	}
1547
1548	bool empty() const
1549	{
1550	return !holder;
1551	}
1552
1553	void reset()
1554	{
1555	if (holder)
1556	group->pools[type]->deallocate_opaque(ptr: holder);
1557	holder = nullptr;
1558	type = TypeNone;
1559	}
1560
1561	void set_allow_type_rewrite()
1562	{
1563	allow_type_rewrite = true;
1564	}
1565
1566	private:
1567	ObjectPoolGroup *group = nullptr;
1568	IVariant *holder = nullptr;
1569	Types type = TypeNone;
1570	bool allow_type_rewrite = false;
1571	};
1572
1573	template <typename T>
1574	T &variant_get(Variant &var)
1575	{
1576	return var.get<T>();
1577	}
1578
1579	template <typename T>
1580	const T &variant_get(const Variant &var)
1581	{
1582	return var.get<T>();
1583	}
1584
1585	template <typename T, typename... P>
1586	T &variant_set(Variant &var, P &&... args)
1587	{
1588	auto *ptr = var.allocate_and_set<T>(static_cast<Types>(T::type), std::forward<P>(args)...);
1589	return *ptr;
1590	}
1591
1592	struct AccessChainMeta
1593	{
1594	uint32_t storage_physical_type = 0;
1595	bool need_transpose = false;
1596	bool storage_is_packed = false;
1597	bool storage_is_invariant = false;
1598	bool flattened_struct = false;
1599	bool relaxed_precision = false;
1600	bool access_meshlet_position_y = false;
1601	};
1602
1603	enum ExtendedDecorations
1604	{
1605	// Marks if a buffer block is re-packed, i.e. member declaration might be subject to PhysicalTypeID remapping and padding.
1606	SPIRVCrossDecorationBufferBlockRepacked = 0,
1607
1608	// A type in a buffer block might be declared with a different physical type than the logical type.
1609	// If this is not set, PhysicalTypeID == the SPIR-V type as declared.
1610	SPIRVCrossDecorationPhysicalTypeID,
1611
1612	// Marks if the physical type is to be declared with tight packing rules, i.e. packed_floatN on MSL and friends.
1613	// If this is set, PhysicalTypeID might also be set. It can be set to same as logical type if all we're doing
1614	// is converting float3 to packed_float3 for example.
1615	// If this is marked on a struct, it means the struct itself must use only Packed types for all its members.
1616	SPIRVCrossDecorationPhysicalTypePacked,
1617
1618	// The padding in bytes before declaring this struct member.
1619	// If used on a struct type, marks the target size of a struct.
1620	SPIRVCrossDecorationPaddingTarget,
1621
1622	SPIRVCrossDecorationInterfaceMemberIndex,
1623	SPIRVCrossDecorationInterfaceOrigID,
1624	SPIRVCrossDecorationResourceIndexPrimary,
1625	// Used for decorations like resource indices for samplers when part of combined image samplers.
1626	// A variable might need to hold two resource indices in this case.
1627	SPIRVCrossDecorationResourceIndexSecondary,
1628	// Used for resource indices for multiplanar images when part of combined image samplers.
1629	SPIRVCrossDecorationResourceIndexTertiary,
1630	SPIRVCrossDecorationResourceIndexQuaternary,
1631
1632	// Marks a buffer block for using explicit offsets (GLSL/HLSL).
1633	SPIRVCrossDecorationExplicitOffset,
1634
1635	// Apply to a variable in the Input storage class; marks it as holding the base group passed to vkCmdDispatchBase(),
1636	// or the base vertex and instance indices passed to vkCmdDrawIndexed().
1637	// In MSL, this is used to adjust the WorkgroupId and GlobalInvocationId variables in compute shaders,
1638	// and to hold the BaseVertex and BaseInstance variables in vertex shaders.
1639	SPIRVCrossDecorationBuiltInDispatchBase,
1640
1641	// Apply to a variable that is a function parameter; marks it as being a "dynamic"
1642	// combined image-sampler. In MSL, this is used when a function parameter might hold
1643	// either a regular combined image-sampler or one that has an attached sampler
1644	// Y'CbCr conversion.
1645	SPIRVCrossDecorationDynamicImageSampler,
1646
1647	// Apply to a variable in the Input storage class; marks it as holding the size of the stage
1648	// input grid.
1649	// In MSL, this is used to hold the vertex and instance counts in a tessellation pipeline
1650	// vertex shader.
1651	SPIRVCrossDecorationBuiltInStageInputSize,
1652
1653	// Apply to any access chain of a tessellation I/O variable; stores the type of the sub-object
1654	// that was chained to, as recorded in the input variable itself. This is used in case the pointer
1655	// is itself used as the base of an access chain, to calculate the original type of the sub-object
1656	// chained to, in case a swizzle needs to be applied. This should not happen normally with valid
1657	// SPIR-V, but the MSL backend can change the type of input variables, necessitating the
1658	// addition of swizzles to keep the generated code compiling.
1659	SPIRVCrossDecorationTessIOOriginalInputTypeID,
1660
1661	// Apply to any access chain of an interface variable used with pull-model interpolation, where the variable is a
1662	// vector but the resulting pointer is a scalar; stores the component index that is to be accessed by the chain.
1663	// This is used when emitting calls to interpolation functions on the chain in MSL: in this case, the component
1664	// must be applied to the result, since pull-model interpolants in MSL cannot be swizzled directly, but the
1665	// results of interpolation can.
1666	SPIRVCrossDecorationInterpolantComponentExpr,
1667
1668	// Apply to any struct type that is used in the Workgroup storage class.
1669	// This causes matrices in MSL prior to Metal 3.0 to be emitted using a special
1670	// class that is convertible to the standard matrix type, to work around the
1671	// lack of constructors in the 'threadgroup' address space.
1672	SPIRVCrossDecorationWorkgroupStruct,
1673
1674	SPIRVCrossDecorationOverlappingBinding,
1675
1676	SPIRVCrossDecorationCount
1677	};
1678
1679	struct Meta
1680	{
1681	struct Decoration
1682	{
1683	std::string alias;
1684	std::string qualified_alias;
1685	std::string hlsl_semantic;
1686	std::string user_type;
1687	Bitset decoration_flags;
1688	spv::BuiltIn builtin_type = spv::BuiltInMax;
1689	uint32_t location = 0;
1690	uint32_t component = 0;
1691	uint32_t set = 0;
1692	uint32_t binding = 0;
1693	uint32_t offset = 0;
1694	uint32_t xfb_buffer = 0;
1695	uint32_t xfb_stride = 0;
1696	uint32_t stream = 0;
1697	uint32_t array_stride = 0;
1698	uint32_t matrix_stride = 0;
1699	uint32_t input_attachment = 0;
1700	uint32_t spec_id = 0;
1701	uint32_t index = 0;
1702	spv::FPRoundingMode fp_rounding_mode = spv::FPRoundingModeMax;
1703	bool builtin = false;
1704	bool qualified_alias_explicit_override = false;
1705
1706	struct Extended
1707	{
1708	Extended()
1709	{
1710	// MSVC 2013 workaround to init like this.
1711	for (auto &v : values)
1712	v = 0;
1713	}
1714
1715	Bitset flags;
1716	uint32_t values[SPIRVCrossDecorationCount];
1717	} extended;
1718	};
1719
1720	Decoration decoration;
1721
1722	// Intentionally not a SmallVector. Decoration is large and somewhat rare.
1723	Vector<Decoration> members;
1724
1725	std::unordered_map<uint32_t, uint32_t> decoration_word_offset;
1726
1727	// For SPV_GOOGLE_hlsl_functionality1.
1728	bool hlsl_is_magic_counter_buffer = false;
1729	// ID for the sibling counter buffer.
1730	uint32_t hlsl_magic_counter_buffer = 0;
1731	};
1732
1733	// A user callback that remaps the type of any variable.
1734	// var_name is the declared name of the variable.
1735	// name_of_type is the textual name of the type which will be used in the code unless written to by the callback.
1736	using VariableTypeRemapCallback =
1737	std::function<void(const SPIRType &type, const std::string &var_name, std::string &name_of_type)>;
1738
1739	class Hasher
1740	{
1741	public:
1742	inline void u32(uint32_t value)
1743	{
1744	h = (h * 0x100000001b3ull) ^ value;
1745	}
1746
1747	inline uint64_t get() const
1748	{
1749	return h;
1750	}
1751
1752	private:
1753	uint64_t h = 0xcbf29ce484222325ull;
1754	};
1755
1756	static inline bool type_is_floating_point(const SPIRType &type)
1757	{
1758	return type.basetype == SPIRType::Half || type.basetype == SPIRType::Float || type.basetype == SPIRType::Double;
1759	}
1760
1761	static inline bool type_is_integral(const SPIRType &type)
1762	{
1763	return type.basetype == SPIRType::SByte || type.basetype == SPIRType::UByte || type.basetype == SPIRType::Short ||
1764	type.basetype == SPIRType::UShort || type.basetype == SPIRType::Int || type.basetype == SPIRType::UInt ||
1765	type.basetype == SPIRType::Int64 || type.basetype == SPIRType::UInt64;
1766	}
1767
1768	static inline SPIRType::BaseType to_signed_basetype(uint32_t width)
1769	{
1770	switch (width)
1771	{
1772	case 8:
1773	return SPIRType::SByte;
1774	case 16:
1775	return SPIRType::Short;
1776	case 32:
1777	return SPIRType::Int;
1778	case 64:
1779	return SPIRType::Int64;
1780	default:
1781	SPIRV_CROSS_THROW("Invalid bit width.");
1782	}
1783	}
1784
1785	static inline SPIRType::BaseType to_unsigned_basetype(uint32_t width)
1786	{
1787	switch (width)
1788	{
1789	case 8:
1790	return SPIRType::UByte;
1791	case 16:
1792	return SPIRType::UShort;
1793	case 32:
1794	return SPIRType::UInt;
1795	case 64:
1796	return SPIRType::UInt64;
1797	default:
1798	SPIRV_CROSS_THROW("Invalid bit width.");
1799	}
1800	}
1801
1802	// Returns true if an arithmetic operation does not change behavior depending on signedness.
1803	static inline bool opcode_is_sign_invariant(spv::Op opcode)
1804	{
1805	switch (opcode)
1806	{
1807	case spv::OpIEqual:
1808	case spv::OpINotEqual:
1809	case spv::OpISub:
1810	case spv::OpIAdd:
1811	case spv::OpIMul:
1812	case spv::OpShiftLeftLogical:
1813	case spv::OpBitwiseOr:
1814	case spv::OpBitwiseXor:
1815	case spv::OpBitwiseAnd:
1816	return true;
1817
1818	default:
1819	return false;
1820	}
1821	}
1822
1823	static inline bool opcode_can_promote_integer_implicitly(spv::Op opcode)
1824	{
1825	switch (opcode)
1826	{
1827	case spv::OpSNegate:
1828	case spv::OpNot:
1829	case spv::OpBitwiseAnd:
1830	case spv::OpBitwiseOr:
1831	case spv::OpBitwiseXor:
1832	case spv::OpShiftLeftLogical:
1833	case spv::OpShiftRightLogical:
1834	case spv::OpShiftRightArithmetic:
1835	case spv::OpIAdd:
1836	case spv::OpISub:
1837	case spv::OpIMul:
1838	case spv::OpSDiv:
1839	case spv::OpUDiv:
1840	case spv::OpSRem:
1841	case spv::OpUMod:
1842	case spv::OpSMod:
1843	return true;
1844
1845	default:
1846	return false;
1847	}
1848	}
1849
1850	struct SetBindingPair
1851	{
1852	uint32_t desc_set;
1853	uint32_t binding;
1854
1855	inline bool operator==(const SetBindingPair &other) const
1856	{
1857	return desc_set == other.desc_set && binding == other.binding;
1858	}
1859
1860	inline bool operator<(const SetBindingPair &other) const
1861	{
1862	return desc_set < other.desc_set || (desc_set == other.desc_set && binding < other.binding);
1863	}
1864	};
1865
1866	struct LocationComponentPair
1867	{
1868	uint32_t location;
1869	uint32_t component;
1870
1871	inline bool operator==(const LocationComponentPair &other) const
1872	{
1873	return location == other.location && component == other.component;
1874	}
1875
1876	inline bool operator<(const LocationComponentPair &other) const
1877	{
1878	return location < other.location || (location == other.location && component < other.component);
1879	}
1880	};
1881
1882	struct StageSetBinding
1883	{
1884	spv::ExecutionModel model;
1885	uint32_t desc_set;
1886	uint32_t binding;
1887
1888	inline bool operator==(const StageSetBinding &other) const
1889	{
1890	return model == other.model && desc_set == other.desc_set && binding == other.binding;
1891	}
1892	};
1893
1894	struct InternalHasher
1895	{
1896	inline size_t operator()(const SetBindingPair &value) const
1897	{
1898	// Quality of hash doesn't really matter here.
1899	auto hash_set = std::hash<uint32_t>()(value.desc_set);
1900	auto hash_binding = std::hash<uint32_t>()(value.binding);
1901	return (hash_set * 0x10001b31) ^ hash_binding;
1902	}
1903
1904	inline size_t operator()(const LocationComponentPair &value) const
1905	{
1906	// Quality of hash doesn't really matter here.
1907	auto hash_set = std::hash<uint32_t>()(value.location);
1908	auto hash_binding = std::hash<uint32_t>()(value.component);
1909	return (hash_set * 0x10001b31) ^ hash_binding;
1910	}
1911
1912	inline size_t operator()(const StageSetBinding &value) const
1913	{
1914	// Quality of hash doesn't really matter here.
1915	auto hash_model = std::hash<uint32_t>()(value.model);
1916	auto hash_set = std::hash<uint32_t>()(value.desc_set);
1917	auto tmp_hash = (hash_model * 0x10001b31) ^ hash_set;
1918	return (tmp_hash * 0x10001b31) ^ value.binding;
1919	}
1920	};
1921
1922	// Special constant used in a {MSL,HLSL}ResourceBinding desc_set
1923	// element to indicate the bindings for the push constants.
1924	static const uint32_t ResourceBindingPushConstantDescriptorSet = ~(0u);
1925
1926	// Special constant used in a {MSL,HLSL}ResourceBinding binding
1927	// element to indicate the bindings for the push constants.
1928	static const uint32_t ResourceBindingPushConstantBinding = 0;
1929	} // namespace SPIRV_CROSS_NAMESPACE
1930
1931	namespace std
1932	{
1933	template <SPIRV_CROSS_NAMESPACE::Types type>
1934	struct hash<SPIRV_CROSS_NAMESPACE::TypedID<type>>
1935	{
1936	size_t operator()(const SPIRV_CROSS_NAMESPACE::TypedID<type> &value) const
1937	{
1938	return std::hash<uint32_t>()(value);
1939	}
1940	};
1941	} // namespace std
1942
1943	#endif
1944