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
298#if defined(__clang__) || defined(__GNUC__)
299#pragma GCC diagnostic pop
300#elif defined(_MSC_VER)
301#pragma warning(pop)
302#endif
303
304class FloatFormatter
305{
306public:
307 virtual ~FloatFormatter() = default;
308 virtual std::string format_float(float value) = 0;
309 virtual std::string format_double(double value) = 0;
310};
311
312template <typename T>
313struct 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
335struct 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
350struct EmbeddedInstruction : Instruction
351{
352 SmallVector<uint32_t> ops;
353};
354
355enum 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
374template <Types type>
375class TypedID;
376
377template <>
378class TypedID<TypeNone>
379{
380public:
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
413private:
414 uint32_t id = 0;
415};
416
417template <Types type>
418class TypedID
419{
420public:
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
437private:
438 uint32_t id = 0;
439};
440
441using VariableID = TypedID<TypeVariable>;
442using TypeID = TypedID<TypeType>;
443using ConstantID = TypedID<TypeConstant>;
444using FunctionID = TypedID<TypeFunction>;
445using BlockID = TypedID<TypeBlock>;
446using ID = TypedID<TypeNone>;
447
448// Helper for Variant interface.
449struct IVariant
450{
451 virtual ~IVariant() = default;
452 virtual IVariant *clone(ObjectPoolBase *pool) = 0;
453 ID self = 0;
454
455protected:
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
467struct 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
483struct 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.
502struct 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
521struct 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
544struct 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
641struct 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.
673struct 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
703struct 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
759struct 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
777struct 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
929struct 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
1030struct 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
1071struct 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
1132struct 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.
1420struct ObjectPoolGroup
1421{
1422 std::unique_ptr<ObjectPoolBase> pools[TypeCount];
1423};
1424
1425class Variant
1426{
1427public:
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
1566private:
1567 ObjectPoolGroup *group = nullptr;
1568 IVariant *holder = nullptr;
1569 Types type = TypeNone;
1570 bool allow_type_rewrite = false;
1571};
1572
1573template <typename T>
1574T &variant_get(Variant &var)
1575{
1576 return var.get<T>();
1577}
1578
1579template <typename T>
1580const T &variant_get(const Variant &var)
1581{
1582 return var.get<T>();
1583}
1584
1585template <typename T, typename... P>
1586T &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
1592struct 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
1603enum 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
1679struct 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.
1736using VariableTypeRemapCallback =
1737 std::function<void(const SPIRType &type, const std::string &var_name, std::string &name_of_type)>;
1738
1739class Hasher
1740{
1741public:
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
1752private:
1753 uint64_t h = 0xcbf29ce484222325ull;
1754};
1755
1756static 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
1761static 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
1768static 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
1785static 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.
1803static 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
1823static 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
1850struct 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
1866struct 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
1882struct 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
1894struct 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.
1924static 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.
1928static const uint32_t ResourceBindingPushConstantBinding = 0;
1929} // namespace SPIRV_CROSS_NAMESPACE
1930
1931namespace std
1932{
1933template <SPIRV_CROSS_NAMESPACE::Types type>
1934struct 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

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