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

source code of qtshadertools/src/3rdparty/SPIRV-Cross/spirv_common.hpp