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