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