1 //
2 // Copyright 2020 The ANGLE Project Authors. All rights reserved.
3 // Use of this source code is governed by a BSD-style license that can be
4 // found in the LICENSE file.
5 //
6 
7 #include <cctype>
8 
9 #include "compiler/translator/InfoSink.h"
10 #include "compiler/translator/Symbol.h"
11 #include "compiler/translator/TranslatorMetalDirect/AstHelpers.h"
12 #include "compiler/translator/TranslatorMetalDirect/Name.h"
13 #include "compiler/translator/TranslatorMetalDirect/ProgramPrelude.h"
14 #include "compiler/translator/tree_util/IntermTraverse.h"
15 
16 using namespace sh;
17 
18 ////////////////////////////////////////////////////////////////////////////////
19 
20 namespace
21 {
22 
23 class ProgramPrelude : public TIntermTraverser
24 {
25     using LineTag       = unsigned;
26     using FuncEmitter   = void (*)(ProgramPrelude &, const TFunction &);
27     using FuncToEmitter = std::map<Name, FuncEmitter>;
28 
29   public:
ProgramPrelude(TInfoSinkBase & out,const ProgramPreludeConfig & ppc)30     ProgramPrelude(TInfoSinkBase &out, const ProgramPreludeConfig &ppc)
31         : TIntermTraverser(true, false, false), mOut(out)
32     {
33         ALWAYS_INLINE();
34         int_clamp();
35         if (ppc.hasStructEq)
36         {
37             equalVector();
38             equalMatrix();
39         }
40 
41         switch (ppc.shaderType)
42         {
43             case MetalShaderType::None:
44                 ASSERT(0 && "ppc.shaderType should not be ShaderTypeNone");
45                 break;
46             case MetalShaderType::Vertex:
47                 transform_feedback_guard();
48                 break;
49             case MetalShaderType::Fragment:
50                 writeSampleMask();
51                 break;
52             case MetalShaderType::Compute:
53                 ASSERT(0 && "compute shaders not currently supported");
54                 break;
55             default:
56                 break;
57         }
58 
59 #if 1
60         mOut << "#define ANGLE_tensor metal::array\n";
61         mOut << "#pragma clang diagnostic ignored \"-Wunused-value\"\n";
62 #else
63         tensor();
64 #endif
65     }
66 
67   private:
emitGuard(LineTag lineTag)68     bool emitGuard(LineTag lineTag)
69     {
70         if (mEmitted.find(lineTag) != mEmitted.end())
71         {
72             return false;
73         }
74         mEmitted.insert(lineTag);
75         return true;
76     }
77 
78     static FuncToEmitter BuildFuncToEmitter();
79 
80     void visitOperator(TOperator op,
81                        const TFunction *func,
82                        const TType *argType0,
83                        const TType *argType1 = nullptr);
84     void visitVariable(const Name &name, const TType &type);
85     void visitVariable(const TVariable &var);
86     void visitStructure(const TStructure &s);
87 
88     bool visitBinary(Visit, TIntermBinary *node) override;
89     bool visitUnary(Visit, TIntermUnary *node) override;
90     bool visitAggregate(Visit, TIntermAggregate *node) override;
91     bool visitDeclaration(Visit, TIntermDeclaration *node) override;
92     void visitSymbol(TIntermSymbol *node) override;
93 
94   private:
95     void ALWAYS_INLINE();
96 
97     void include_metal_atomic();
98     void include_metal_common();
99     void include_metal_geometric();
100     void include_metal_graphics();
101     void include_metal_math();
102     void include_metal_matrix();
103     void include_metal_pack();
104     void include_metal_relational();
105 
106     void transform_feedback_guard();
107 
108     void enable_if();
109     void scalar_of();
110     void is_scalar();
111     void is_vector();
112     void is_matrix();
113     void addressof();
114     void distance();
115     void length();
116     void dot();
117     void normalize();
118     void faceforward();
119     void reflect();
120     void refract();
121     void degrees();
122     void radians();
123     void mod();
124     void postIncrementMatrix();
125     void preIncrementMatrix();
126     void postDecrementMatrix();
127     void preDecrementMatrix();
128     void negateMatrix();
129     void matmulAssign();
130     void atan();
131     void int_clamp();
132     void addMatrixScalarAssign();
133     void subMatrixScalarAssign();
134     void addMatrixScalar();
135     void subMatrixScalar();
136     void divMatrixScalar();
137     void divMatrixScalarFast();
138     void divMatrixScalarAssign();
139     void divMatrixScalarAssignFast();
140     void tensor();
141     void componentWiseDivide();
142     void componentWiseDivideAssign();
143     void componentWiseMultiply();
144     void outerProduct();
145     void inverse2();
146     void inverse3();
147     void inverse4();
148     void equalVector();
149     void equalMatrix();
150     void notEqualVector();
151     void notEqualStruct();
152     void notEqualStructArray();
153     void notEqualMatrix();
154     void equalArray();
155     void equalStructArray();
156     void notEqualArray();
157     void sign();
158     void pack_half_2x16();
159     void unpack_half_2x16();
160     void vectorElemRef();
161     void swizzleRef();
162     void out();
163     void inout();
164     void flattenArray();
165     void castVector();
166     void castMatrix();
167     void functionConstants();
168     void gradient();
169     void writeSampleMask();
170     void textureEnv();
171     void texelFetch();
172     void texelFetchOffset();
173     void texture();
174     void texture_generic_float2();
175     void texture_generic_float2_float();
176     void texture_generic_float3();
177     void texture_generic_float3_float();
178     void texture_depth2d_float3();
179     void texture_depth2d_float3_float();
180     void texture_depth2darray_float4();
181     void texture_depth2darray_float4_float();
182     void texture_depthcube_float4();
183     void texture_depthcube_float4_float();
184     void texture_texture2darray_float3();
185     void texture_texture2darray_float3_float();
186     void texture_texture2darray_float4();
187     void texture_texture2darray_float4_float();
188     void texture1DLod();
189     void texture1DProj();
190     void texture1DProjLod();
191     void texture2D();
192     void texture2DLod();
193     void texture2DProj();
194     void texture2DProjLod();
195     void texture2DRect();
196     void texture2DRectProj();
197     void texture3DLod();
198     void texture3DProj();
199     void texture3DProjLod();
200     void textureCube();
201     void textureCubeLod();
202     void textureCubeProj();
203     void textureCubeProjLod();
204     void textureGrad();
205     void textureGrad_generic_floatN_floatN_floatN();
206     void textureGrad_generic_float3_float2_float2();
207     void textureGrad_generic_float4_float2_float2();
208     void textureGrad_depth2d_float3_float2_float2();
209     void textureGrad_depth2darray_float4_float2_float2();
210     void textureGrad_depthcube_float4_float3_float3();
211     void textureGrad_texturecube_float3_float3_float3();
212     void textureGradOffset();
213     void textureGradOffset_generic_floatN_floatN_floatN_intN();
214     void textureGradOffset_generic_float3_float2_float2_int2();
215     void textureGradOffset_generic_float4_float2_float2_int2();
216     void textureGradOffset_depth2d_float3_float2_float2_int2();
217     void textureGradOffset_depth2darray_float4_float2_float2_int2();
218     void textureGradOffset_depthcube_float4_float3_float3_int3();
219     void textureGradOffset_texturecube_float3_float3_float3_int3();
220     void textureLod();
221     void textureLod_generic_float2();
222     void textureLod_generic_float3();
223     void textureLod_depth2d_float3();
224     void textureLod_texture2darray_float3();
225     void textureLod_texture2darray_float4();
226     void textureLodOffset();
227     void textureOffset();
228     void textureProj();
229     void textureProjGrad();
230     void textureProjGrad_generic_float3_float2_float2();
231     void textureProjGrad_generic_float4_float2_float2();
232     void textureProjGrad_depth2d_float4_float2_float2();
233     void textureProjGrad_texture3d_float4_float3_float3();
234     void textureProjGradOffset();
235     void textureProjGradOffset_generic_float3_float2_float2_int2();
236     void textureProjGradOffset_generic_float4_float2_float2_int2();
237     void textureProjGradOffset_depth2d_float4_float2_float2_int2();
238     void textureProjGradOffset_texture3d_float4_float3_float3_int3();
239     void textureProjLod();
240     void textureProjLod_generic_float3();
241     void textureProjLod_generic_float4();
242     void textureProjLod_depth2d_float4();
243     void textureProjLod_texture3d_float4();
244     void textureProjLodOffset();
245     void textureProjOffset();
246     void textureSize();
247 
248   private:
249     TInfoSinkBase &mOut;
250     std::unordered_set<LineTag> mEmitted;
251     std::unordered_set<const TSymbol *> mHandled;
252     const FuncToEmitter mFuncToEmitter = BuildFuncToEmitter();
253 };
254 
255 }  // anonymous namespace
256 
257 ////////////////////////////////////////////////////////////////////////////////
258 
259 #define PROGRAM_PRELUDE_INCLUDE(header)             \
260     void ProgramPrelude::include_##header()         \
261     {                                               \
262         if (emitGuard(__LINE__))                    \
263         {                                           \
264             mOut << ("#include <" #header ">\n\n"); \
265         }                                           \
266     }
267 
268 #define PROGRAM_PRELUDE_DECLARE(name, code, ...)                \
269     void ProgramPrelude::name()                                 \
270     {                                                           \
271         ASSERT(code[0] == '\n');                                \
272         if (emitGuard(__LINE__))                                \
273         {                                                       \
274             __VA_ARGS__; /* dependencies */                     \
275             mOut << (static_cast<const char *>(code "\n") + 1); \
276         }                                                       \
277     }
278 
279 ////////////////////////////////////////////////////////////////////////////////
280 
281 PROGRAM_PRELUDE_INCLUDE(metal_atomic)
PROGRAM_PRELUDE_INCLUDE(metal_common)282 PROGRAM_PRELUDE_INCLUDE(metal_common)
283 PROGRAM_PRELUDE_INCLUDE(metal_geometric)
284 PROGRAM_PRELUDE_INCLUDE(metal_graphics)
285 PROGRAM_PRELUDE_INCLUDE(metal_math)
286 PROGRAM_PRELUDE_INCLUDE(metal_matrix)
287 PROGRAM_PRELUDE_INCLUDE(metal_pack)
288 PROGRAM_PRELUDE_INCLUDE(metal_relational)
289 
290 PROGRAM_PRELUDE_DECLARE(transform_feedback_guard, R"(
291 #if TRANSFORM_FEEDBACK_ENABLED
292     #define __VERTEX_OUT(args) void
293 #else
294     #define __VERTEX_OUT(args) args
295 #endif
296 )")
297 
298 PROGRAM_PRELUDE_DECLARE(ALWAYS_INLINE, R"(
299 #define ANGLE_ALWAYS_INLINE __attribute__((always_inline))
300 )")
301 
302 PROGRAM_PRELUDE_DECLARE(enable_if, R"(
303 template <bool B, typename T = void>
304 struct ANGLE_enable_if {};
305 template <typename T>
306 struct ANGLE_enable_if<true, T>
307 {
308     using type = T;
309 };
310 template <bool B>
311 using ANGLE_enable_if_t = typename ANGLE_enable_if<B>::type;
312 )")
313 
314 PROGRAM_PRELUDE_DECLARE(scalar_of, R"(
315 template <typename T>
316 struct ANGLE_scalar_of
317 {
318     using type = T;
319 };
320 template <typename T>
321 using ANGLE_scalar_of_t = typename ANGLE_scalar_of<T>::type;
322 )")
323 
324 PROGRAM_PRELUDE_DECLARE(is_scalar, R"(
325 template <typename T>
326 struct ANGLE_is_scalar {};
327 #define ANGLE_DEFINE_SCALAR(scalar) \
328     template <> struct ANGLE_is_scalar<scalar> { enum { value = true }; }
329 ANGLE_DEFINE_SCALAR(bool);
330 ANGLE_DEFINE_SCALAR(char);
331 ANGLE_DEFINE_SCALAR(short);
332 ANGLE_DEFINE_SCALAR(int);
333 ANGLE_DEFINE_SCALAR(long);
334 ANGLE_DEFINE_SCALAR(uchar);
335 ANGLE_DEFINE_SCALAR(ushort);
336 ANGLE_DEFINE_SCALAR(uint);
337 ANGLE_DEFINE_SCALAR(ulong);
338 ANGLE_DEFINE_SCALAR(half);
339 ANGLE_DEFINE_SCALAR(float);
340 )")
341 
342 PROGRAM_PRELUDE_DECLARE(is_vector,
343                         R"(
344 template <typename T>
345 struct ANGLE_is_vector
346 {
347     enum { value = false };
348 };
349 #define ANGLE_DEFINE_VECTOR(scalar) \
350     template <> struct ANGLE_is_vector<metal::scalar ## 2> { enum { value = true }; }; \
351     template <> struct ANGLE_is_vector<metal::scalar ## 3> { enum { value = true }; }; \
352     template <> struct ANGLE_is_vector<metal::scalar ## 4> { enum { value = true }; }; \
353     template <> struct ANGLE_scalar_of<metal::scalar ## 2> { using type = scalar; }; \
354     template <> struct ANGLE_scalar_of<metal::scalar ## 3> { using type = scalar; }; \
355     template <> struct ANGLE_scalar_of<metal::scalar ## 4> { using type = scalar; }
356 ANGLE_DEFINE_VECTOR(bool);
357 ANGLE_DEFINE_VECTOR(char);
358 ANGLE_DEFINE_VECTOR(short);
359 ANGLE_DEFINE_VECTOR(int);
360 ANGLE_DEFINE_VECTOR(long);
361 ANGLE_DEFINE_VECTOR(uchar);
362 ANGLE_DEFINE_VECTOR(ushort);
363 ANGLE_DEFINE_VECTOR(uint);
364 ANGLE_DEFINE_VECTOR(ulong);
365 ANGLE_DEFINE_VECTOR(half);
366 ANGLE_DEFINE_VECTOR(float);
367 )",
368                         scalar_of())
369 
370 PROGRAM_PRELUDE_DECLARE(is_matrix,
371                         R"(
372 template <typename T>
373 struct ANGLE_is_matrix
374 {
375     enum { value = false };
376 };
377 #define ANGLE_DEFINE_MATRIX(scalar) \
378     template <> struct ANGLE_is_matrix<metal::scalar ## 2x2> { enum { value = true }; }; \
379     template <> struct ANGLE_is_matrix<metal::scalar ## 2x3> { enum { value = true }; }; \
380     template <> struct ANGLE_is_matrix<metal::scalar ## 2x4> { enum { value = true }; }; \
381     template <> struct ANGLE_is_matrix<metal::scalar ## 3x2> { enum { value = true }; }; \
382     template <> struct ANGLE_is_matrix<metal::scalar ## 3x3> { enum { value = true }; }; \
383     template <> struct ANGLE_is_matrix<metal::scalar ## 3x4> { enum { value = true }; }; \
384     template <> struct ANGLE_is_matrix<metal::scalar ## 4x2> { enum { value = true }; }; \
385     template <> struct ANGLE_is_matrix<metal::scalar ## 4x3> { enum { value = true }; }; \
386     template <> struct ANGLE_is_matrix<metal::scalar ## 4x4> { enum { value = true }; }; \
387     template <> struct ANGLE_scalar_of<metal::scalar ## 2x2> { using type = scalar; }; \
388     template <> struct ANGLE_scalar_of<metal::scalar ## 2x3> { using type = scalar; }; \
389     template <> struct ANGLE_scalar_of<metal::scalar ## 2x4> { using type = scalar; }; \
390     template <> struct ANGLE_scalar_of<metal::scalar ## 3x2> { using type = scalar; }; \
391     template <> struct ANGLE_scalar_of<metal::scalar ## 3x3> { using type = scalar; }; \
392     template <> struct ANGLE_scalar_of<metal::scalar ## 3x4> { using type = scalar; }; \
393     template <> struct ANGLE_scalar_of<metal::scalar ## 4x2> { using type = scalar; }; \
394     template <> struct ANGLE_scalar_of<metal::scalar ## 4x3> { using type = scalar; }; \
395     template <> struct ANGLE_scalar_of<metal::scalar ## 4x4> { using type = scalar; }
396 ANGLE_DEFINE_MATRIX(half);
397 ANGLE_DEFINE_MATRIX(float);
398 )",
399                         scalar_of())
400 
401 PROGRAM_PRELUDE_DECLARE(addressof,
402                         R"(
403 template <typename T>
404 ANGLE_ALWAYS_INLINE thread T * ANGLE_addressof(thread T &ref)
405 {
406     return &ref;
407 }
408 )")
409 
410 PROGRAM_PRELUDE_DECLARE(distance,
411                         R"(
412 template <typename T, typename Enable = void>
413 struct ANGLE_distance_impl
414 {
415     static ANGLE_ALWAYS_INLINE ANGLE_scalar_of_t<T> exec(T x, T y)
416     {
417         return metal::distance(x, y);
418     }
419 };
420 template <typename T>
421 struct ANGLE_distance_impl<T, ANGLE_enable_if_t<ANGLE_is_scalar<T>::value>>
422 {
423     static ANGLE_ALWAYS_INLINE T exec(T x, T y)
424     {
425         return metal::abs(x - y);
426     }
427 };
428 template <typename T>
429 ANGLE_ALWAYS_INLINE ANGLE_scalar_of_t<T> ANGLE_distance(T x, T y)
430 {
431     return ANGLE_distance_impl<T>::exec(x, y);
432 };
433 )",
434                         include_metal_geometric(),
435                         include_metal_math(),
436                         enable_if(),
437                         is_scalar(),
438                         is_vector(),
439                         is_matrix())
440 
441 PROGRAM_PRELUDE_DECLARE(length,
442                         R"(
443 template <typename T, typename Enable = void>
444 struct ANGLE_length_impl
445 {
446     static ANGLE_ALWAYS_INLINE ANGLE_scalar_of_t<T> exec(T x)
447     {
448         return metal::length(x);
449     }
450 };
451 template <typename T>
452 struct ANGLE_length_impl<T, ANGLE_enable_if_t<ANGLE_is_scalar<T>::value>>
453 {
454     static ANGLE_ALWAYS_INLINE T exec(T x)
455     {
456         return metal::abs(x);
457     }
458 };
459 template <typename T>
460 ANGLE_ALWAYS_INLINE ANGLE_scalar_of_t<T> ANGLE_length(T x)
461 {
462     return ANGLE_length_impl<T>::exec(x);
463 };
464 )",
465                         include_metal_geometric(),
466                         include_metal_math(),
467                         enable_if(),
468                         is_scalar(),
469                         is_vector(),
470                         is_matrix())
471 
472 PROGRAM_PRELUDE_DECLARE(dot,
473                         R"(
474 template <typename T, typename Enable = void>
475 struct ANGLE_dot_impl
476 {
477     static ANGLE_ALWAYS_INLINE ANGLE_scalar_of_t<T> exec(T x, T y)
478     {
479         return metal::dot(x, y);
480     }
481 };
482 template <typename T>
483 struct ANGLE_dot_impl<T, ANGLE_enable_if_t<ANGLE_is_scalar<T>::value>>
484 {
485     static ANGLE_ALWAYS_INLINE T exec(T x, T y)
486     {
487         return x * y;
488     }
489 };
490 template <typename T>
491 ANGLE_ALWAYS_INLINE ANGLE_scalar_of_t<T> ANGLE_dot(T x, T y)
492 {
493     return ANGLE_dot_impl<T>::exec(x, y);
494 };
495 )",
496                         include_metal_geometric(),
497                         enable_if(),
498                         is_scalar(),
499                         is_vector(),
500                         is_matrix())
501 
502 PROGRAM_PRELUDE_DECLARE(normalize,
503                         R"(
504 template <typename T, typename Enable = void>
505 struct ANGLE_normalize_impl
506 {
507     static ANGLE_ALWAYS_INLINE T exec(T x)
508     {
509         return metal::normalize(x);
510     }
511 };
512 template <typename T>
513 struct ANGLE_normalize_impl<T, ANGLE_enable_if_t<ANGLE_is_scalar<T>::value>>
514 {
515     static ANGLE_ALWAYS_INLINE T exec(T x)
516     {
517         return ANGLE_sign(x);
518     }
519 };
520 template <typename T>
521 ANGLE_ALWAYS_INLINE T ANGLE_normalize(T x)
522 {
523     return ANGLE_normalize_impl<T>::exec(x);
524 };
525 )",
526                         include_metal_common(),
527                         include_metal_geometric(),
528                         enable_if(),
529                         is_scalar(),
530                         is_vector(),
531                         is_matrix(),
532                         sign())
533 
534 PROGRAM_PRELUDE_DECLARE(faceforward,
535                         R"(
536 template <typename T, typename Enable = void>
537 struct ANGLE_faceforward_impl
538 {
539     static ANGLE_ALWAYS_INLINE T exec(T n, T i, T nref)
540     {
541         return metal::faceforward(n, i, nref);
542     }
543 };
544 template <typename T>
545 struct ANGLE_faceforward_impl<T, ANGLE_enable_if_t<ANGLE_is_scalar<T>::value>>
546 {
547     static ANGLE_ALWAYS_INLINE T exec(T n, T i, T nref)
548     {
549         return ANGLE_dot(nref, i) < T(0) ? n : -n;
550     }
551 };
552 template <typename T>
553 ANGLE_ALWAYS_INLINE T ANGLE_faceforward(T n, T i, T nref)
554 {
555     return ANGLE_faceforward_impl<T>::exec(n, i, nref);
556 };
557 )",
558                         include_metal_geometric(),
559                         enable_if(),
560                         is_scalar(),
561                         is_vector(),
562                         is_matrix(),
563                         dot())
564 
565 PROGRAM_PRELUDE_DECLARE(reflect,
566                         R"(
567 template <typename T, typename Enable = void>
568 struct ANGLE_reflect_impl
569 {
570     static ANGLE_ALWAYS_INLINE T exec(T i, T n)
571     {
572         return metal::reflect(i, n);
573     }
574 };
575 template <typename T>
576 struct ANGLE_reflect_impl<T, ANGLE_enable_if_t<ANGLE_is_scalar<T>::value>>
577 {
578     static ANGLE_ALWAYS_INLINE T exec(T i, T n)
579     {
580         return i - T(2) * ANGLE_dot(n, i) * n;
581     }
582 };
583 template <typename T>
584 ANGLE_ALWAYS_INLINE T ANGLE_reflect(T i, T n)
585 {
586     return ANGLE_reflect_impl<T>::exec(i, n);
587 };
588 )",
589                         include_metal_geometric(),
590                         enable_if(),
591                         is_scalar(),
592                         is_vector(),
593                         is_matrix(),
594                         dot())
595 
596 PROGRAM_PRELUDE_DECLARE(refract,
597                         R"(
598 template <typename T, typename Enable = void>
599 struct ANGLE_refract_impl
600 {
601     static ANGLE_ALWAYS_INLINE T exec(T i, T n, ANGLE_scalar_of_t<T> eta)
602     {
603         return metal::refract(i, n, eta);
604     }
605 };
606 template <typename T>
607 struct ANGLE_refract_impl<T, ANGLE_enable_if_t<ANGLE_is_scalar<T>::value>>
608 {
609     static ANGLE_ALWAYS_INLINE T exec(T i, T n, T eta)
610     {
611         auto dotNI = n * i;
612         auto k = T(1) - eta * eta * (T(1) - dotNI * dotNI);
613         if (k < T(0))
614         {
615             return T(0);
616         }
617         else
618         {
619             return eta * i - (eta * dotNI + metal::sqrt(k)) * n;
620         }
621     }
622 };
623 template <typename T>
624 ANGLE_ALWAYS_INLINE T ANGLE_refract(T i, T n, ANGLE_scalar_of_t<T> eta)
625 {
626     return ANGLE_refract_impl<T>::exec(i, n, eta);
627 };
628 )",
629                         include_metal_math(),
630                         include_metal_geometric(),
631                         enable_if(),
632                         is_scalar(),
633                         is_vector(),
634                         is_matrix())
635 
636 PROGRAM_PRELUDE_DECLARE(sign,
637                         R"(
638 template <typename T, typename Enable = void>
639 struct ANGLE_sign_impl
640 {
641     static ANGLE_ALWAYS_INLINE T exec(T x)
642     {
643         return metal::sign(x);
644     }
645 };
646 template <>
647 struct ANGLE_sign_impl<int>
648 {
649     static ANGLE_ALWAYS_INLINE int exec(int x)
650     {
651         return (0 < x) - (x < 0);
652     }
653 };
654 template <int N>
655 struct ANGLE_sign_impl<metal::vec<int, N>>
656 {
657     static ANGLE_ALWAYS_INLINE metal::vec<int, N> exec(metal::vec<int, N> x)
658     {
659         metal::vec<int, N> s;
660         for (int i = 0; i < N; ++i)
661         {
662             s[i] = ANGLE_sign_impl<int>::exec(x[i]);
663         }
664         return s;
665     }
666 };
667 template <typename T>
668 ANGLE_ALWAYS_INLINE T ANGLE_sign(T x)
669 {
670     return ANGLE_sign_impl<T>::exec(x);
671 };
672 )",
673                         include_metal_common())
674 
675 PROGRAM_PRELUDE_DECLARE(int_clamp,
676                         R"(
677 ANGLE_ALWAYS_INLINE int ANGLE_int_clamp(int value, int minValue, int maxValue)
678 {
679     return ((value < minValue) ?  minValue : ((value > maxValue) ? maxValue : value));
680 };
681 )")
682 
683 PROGRAM_PRELUDE_DECLARE(atan,
684                         R"(
685 template <typename T>
686 ANGLE_ALWAYS_INLINE T ANGLE_atan(T yOverX)
687 {
688     return metal::atan(yOverX);
689 }
690 template <typename T>
691 ANGLE_ALWAYS_INLINE T ANGLE_atan(T y, T x)
692 {
693     return metal::atan2(y, x);
694 }
695 )",
696                         include_metal_math())
697 
698 PROGRAM_PRELUDE_DECLARE(degrees, R"(
699 template <typename T>
700 ANGLE_ALWAYS_INLINE T ANGLE_degrees(T x)
701 {
702     return static_cast<T>(57.29577951308232) * x;
703 }
704 )")
705 
706 PROGRAM_PRELUDE_DECLARE(radians, R"(
707 template <typename T>
708 ANGLE_ALWAYS_INLINE T ANGLE_radians(T x)
709 {
710     return static_cast<T>(1.7453292519943295e-2) * x;
711 }
712 )")
713 
714 PROGRAM_PRELUDE_DECLARE(mod,
715                         R"(
716 template <typename X, typename Y>
717 ANGLE_ALWAYS_INLINE X ANGLE_mod(X x, Y y)
718 {
719     return x - y * metal::floor(x / y);
720 }
721 )",
722                         include_metal_math())
723 
724 PROGRAM_PRELUDE_DECLARE(pack_half_2x16,
725                         R"(
726 ANGLE_ALWAYS_INLINE uint ANGLE_pack_half_2x16(float2 v)
727 {
728     return as_type<uint>(half2(v));
729 }
730 )", )
731 
732 PROGRAM_PRELUDE_DECLARE(unpack_half_2x16,
733                         R"(
734 ANGLE_ALWAYS_INLINE float2 ANGLE_unpack_half_2x16(uint x)
735 {
736     return float2(as_type<half2>(x));
737 }
738 )", )
739 
740 PROGRAM_PRELUDE_DECLARE(matmulAssign, R"(
741 template <typename T, int Cols, int Rows>
742 ANGLE_ALWAYS_INLINE thread metal::matrix<T, Cols, Rows> &operator*=(thread metal::matrix<T, Cols, Rows> &a, metal::matrix<T, Cols, Cols> b)
743 {
744     a = a * b;
745     return a;
746 }
747 )")
748 
749 PROGRAM_PRELUDE_DECLARE(postIncrementMatrix,
750                         R"(
751 template <typename T, int Cols, int Rows>
752 ANGLE_ALWAYS_INLINE metal::matrix<T, Cols, Rows> operator++(thread metal::matrix<T, Cols, Rows> &a, int)
753 {
754     auto b = a;
755     a += T(1);
756     return b;
757 }
758 )",
759                         addMatrixScalarAssign())
760 
761 PROGRAM_PRELUDE_DECLARE(preIncrementMatrix,
762                         R"(
763 template <typename T, int Cols, int Rows>
764 ANGLE_ALWAYS_INLINE thread metal::matrix<T, Cols, Rows> &operator++(thread metal::matrix<T, Cols, Rows> &a)
765 {
766     a += T(1);
767     return a;
768 }
769 )",
770                         addMatrixScalarAssign())
771 
772 PROGRAM_PRELUDE_DECLARE(postDecrementMatrix,
773                         R"(
774 template <typename T, int Cols, int Rows>
775 ANGLE_ALWAYS_INLINE metal::matrix<T, Cols, Rows> operator--(thread metal::matrix<T, Cols, Rows> &a, int)
776 {
777     auto b = a;
778     a -= T(1);
779     return b;
780 }
781 )",
782                         subMatrixScalarAssign())
783 
784 PROGRAM_PRELUDE_DECLARE(preDecrementMatrix,
785                         R"(
786 template <typename T, int Cols, int Rows>
787 ANGLE_ALWAYS_INLINE thread metal::matrix<T, Cols, Rows> &operator--(thread metal::matrix<T, Cols, Rows> &a)
788 {
789     a -= T(1);
790     return a;
791 }
792 )",
793                         subMatrixScalarAssign())
794 
795 PROGRAM_PRELUDE_DECLARE(negateMatrix,
796                         R"(
797 template <typename T, int Cols, int Rows>
798 ANGLE_ALWAYS_INLINE metal::matrix<T, Cols, Rows> operator-(metal::matrix<T, Cols, Rows> m)
799 {
800     for (size_t col = 0; col < Cols; ++col)
801     {
802         thread auto &mCol = m[col];
803         mCol = -mCol;
804     }
805     return m;
806 }
807 )", )
808 
809 PROGRAM_PRELUDE_DECLARE(addMatrixScalarAssign, R"(
810 template <typename T, int Cols, int Rows>
811 ANGLE_ALWAYS_INLINE thread metal::matrix<T, Cols, Rows> &operator+=(thread metal::matrix<T, Cols, Rows> &m, T x)
812 {
813     for (size_t col = 0; col < Cols; ++col)
814     {
815         m[col] += x;
816     }
817     return m;
818 }
819 )")
820 
821 PROGRAM_PRELUDE_DECLARE(addMatrixScalar,
822                         R"(
823 template <typename T, int Cols, int Rows>
824 ANGLE_ALWAYS_INLINE metal::matrix<T, Cols, Rows> operator+(metal::matrix<T, Cols, Rows> m, T x)
825 {
826     m += x;
827     return m;
828 }
829 )",
830                         addMatrixScalarAssign())
831 
832 PROGRAM_PRELUDE_DECLARE(subMatrixScalarAssign,
833                         R"(
834 template <typename T, int Cols, int Rows>
835 ANGLE_ALWAYS_INLINE thread metal::matrix<T, Cols, Rows> &operator-=(thread metal::matrix<T, Cols, Rows> &m, T x)
836 {
837     for (size_t col = 0; col < Cols; ++col)
838     {
839         m[col] -= x;
840     }
841     return m;
842 }
843 )", )
844 
845 PROGRAM_PRELUDE_DECLARE(subMatrixScalar,
846                         R"(
847 template <typename T, int Cols, int Rows>
848 ANGLE_ALWAYS_INLINE metal::matrix<T, Cols, Rows> operator-(metal::matrix<T, Cols, Rows> m, T x)
849 {
850     m -= x;
851     return m;
852 }
853 )",
854                         subMatrixScalarAssign())
855 
856 PROGRAM_PRELUDE_DECLARE(divMatrixScalarAssignFast,
857                         R"(
858 template <typename T, int Cols, int Rows>
859 ANGLE_ALWAYS_INLINE thread metal::matrix<T, Cols, Rows> &operator/=(thread metal::matrix<T, Cols, Rows> &m, T x)
860 {
861     x = T(1) / x;
862     for (size_t col = 0; col < Cols; ++col)
863     {
864         m[col] *= x;
865     }
866     return m;
867 }
868 )", )
869 
870 PROGRAM_PRELUDE_DECLARE(divMatrixScalarAssign,
871                         R"(
872 template <typename T, int Cols, int Rows>
873 ANGLE_ALWAYS_INLINE thread metal::matrix<T, Cols, Rows> &operator/=(thread metal::matrix<T, Cols, Rows> &m, T x)
874 {
875     for (size_t col = 0; col < Cols; ++col)
876     {
877         m[col] /= x;
878     }
879     return m;
880 }
881 )", )
882 
883 PROGRAM_PRELUDE_DECLARE(divMatrixScalarFast,
884                         R"(
885 #if __METAL_VERSION__ <= 220
886 template <typename T, int Cols, int Rows>
887 ANGLE_ALWAYS_INLINE metal::matrix<T, Cols, Rows> operator/(metal::matrix<T, Cols, Rows> m, T x)
888 {
889     m /= x;
890     return m;
891 }
892 #endif
893 )",
894                         divMatrixScalarAssignFast())
895 
896 PROGRAM_PRELUDE_DECLARE(divMatrixScalar,
897                         R"(
898 #if __METAL_VERSION__ <= 220
899 template <typename T, int Cols, int Rows>
900 ANGLE_ALWAYS_INLINE metal::matrix<T, Cols, Rows> operator/(metal::matrix<T, Cols, Rows> m, T x)
901 {
902     m /= x;
903     return m;
904 }
905 #endif
906 )",
907                         divMatrixScalarAssign())
908 
909 PROGRAM_PRELUDE_DECLARE(componentWiseDivide, R"(
910 template <typename T, int Cols, int Rows>
911 ANGLE_ALWAYS_INLINE metal::matrix<T, Cols, Rows> operator/(metal::matrix<T, Cols, Rows> a, metal::matrix<T, Cols, Rows> b)
912 {
913     for (size_t col = 0; col < Cols; ++col)
914     {
915         a[col] /= b[col];
916     }
917     return a;
918 }
919 )")
920 
921 PROGRAM_PRELUDE_DECLARE(componentWiseDivideAssign,
922                         R"(
923 template <typename T, int Cols, int Rows>
924 ANGLE_ALWAYS_INLINE thread metal::matrix<T, Cols, Rows> &operator/=(thread metal::matrix<T, Cols, Rows> &a, metal::matrix<T, Cols, Rows> b)
925 {
926     a = a / b;
927     return a;
928 }
929 )",
930                         componentWiseDivide())
931 
932 PROGRAM_PRELUDE_DECLARE(componentWiseMultiply, R"(
933 template <typename T, int Cols, int Rows>
934 ANGLE_ALWAYS_INLINE metal::matrix<T, Cols, Rows> ANGLE_componentWiseMultiply(metal::matrix<T, Cols, Rows> a, metal::matrix<T, Cols, Rows> b)
935 {
936     for (size_t col = 0; col < Cols; ++col)
937     {
938         a[col] *= b[col];
939     }
940     return a;
941 }
942 )")
943 
944 PROGRAM_PRELUDE_DECLARE(outerProduct, R"(
945 template <typename T, int M, int N>
946 ANGLE_ALWAYS_INLINE metal::matrix<T, N, M> ANGLE_outerProduct(metal::vec<T, M> u, metal::vec<T, N> v)
947 {
948     metal::matrix<T, N, M> o;
949     for (size_t n = 0; n < N; ++n)
950     {
951         o[n] = u * v[n];
952     }
953     return o;
954 }
955 )")
956 
957 PROGRAM_PRELUDE_DECLARE(inverse2, R"(
958 template <typename T>
959 ANGLE_ALWAYS_INLINE metal::matrix<T, 2, 2> ANGLE_inverse(metal::matrix<T, 2, 2> m)
960 {
961     metal::matrix<T, 2, 2> adj;
962     adj[0][0] =  m[1][1];
963     adj[0][1] = -m[0][1];
964     adj[1][0] = -m[1][0];
965     adj[1][1] =  m[0][0];
966     T det = (adj[0][0] * m[0][0]) + (adj[0][1] * m[1][0]);
967     return adj * (T(1) / det);
968 }
969 )")
970 
971 PROGRAM_PRELUDE_DECLARE(inverse3, R"(
972 template <typename T>
973 ANGLE_ALWAYS_INLINE metal::matrix<T, 3, 3> ANGLE_inverse(metal::matrix<T, 3, 3> m)
974 {
975     T a = m[1][1] * m[2][2] - m[2][1] * m[1][2];
976     T b = m[1][0] * m[2][2];
977     T c = m[1][2] * m[2][0];
978     T d = m[1][0] * m[2][1];
979     T det = m[0][0] * (a) -
980             m[0][1] * (b - c) +
981             m[0][2] * (d - m[1][1] * m[2][0]);
982     det = T(1) / det;
983     metal::matrix<T, 3, 3> minv;
984     minv[0][0] = (a) * det;
985     minv[0][1] = (m[0][2] * m[2][1] - m[0][1] * m[2][2]) * det;
986     minv[0][2] = (m[0][1] * m[1][2] - m[0][2] * m[1][1]) * det;
987     minv[1][0] = (c - b) * det;
988     minv[1][1] = (m[0][0] * m[2][2] - m[0][2] * m[2][0]) * det;
989     minv[1][2] = (m[1][0] * m[0][2] - m[0][0] * m[1][2]) * det;
990     minv[2][0] = (d - m[2][0] * m[1][1]) * det;
991     minv[2][1] = (m[2][0] * m[0][1] - m[0][0] * m[2][1]) * det;
992     minv[2][2] = (m[0][0] * m[1][1] - m[1][0] * m[0][1]) * det;
993     return minv;
994 }
995 )")
996 
997 PROGRAM_PRELUDE_DECLARE(inverse4, R"(
998 template <typename T>
999 ANGLE_ALWAYS_INLINE metal::matrix<T, 4, 4> ANGLE_inverse(metal::matrix<T, 4, 4> m)
1000 {
1001     T A2323 = m[2][2] * m[3][3] - m[2][3] * m[3][2];
1002     T A1323 = m[2][1] * m[3][3] - m[2][3] * m[3][1];
1003     T A1223 = m[2][1] * m[3][2] - m[2][2] * m[3][1];
1004     T A0323 = m[2][0] * m[3][3] - m[2][3] * m[3][0];
1005     T A0223 = m[2][0] * m[3][2] - m[2][2] * m[3][0];
1006     T A0123 = m[2][0] * m[3][1] - m[2][1] * m[3][0];
1007     T A2313 = m[1][2] * m[3][3] - m[1][3] * m[3][2];
1008     T A1313 = m[1][1] * m[3][3] - m[1][3] * m[3][1];
1009     T A1213 = m[1][1] * m[3][2] - m[1][2] * m[3][1];
1010     T A2312 = m[1][2] * m[2][3] - m[1][3] * m[2][2];
1011     T A1312 = m[1][1] * m[2][3] - m[1][3] * m[2][1];
1012     T A1212 = m[1][1] * m[2][2] - m[1][2] * m[2][1];
1013     T A0313 = m[1][0] * m[3][3] - m[1][3] * m[3][0];
1014     T A0213 = m[1][0] * m[3][2] - m[1][2] * m[3][0];
1015     T A0312 = m[1][0] * m[2][3] - m[1][3] * m[2][0];
1016     T A0212 = m[1][0] * m[2][2] - m[1][2] * m[2][0];
1017     T A0113 = m[1][0] * m[3][1] - m[1][1] * m[3][0];
1018     T A0112 = m[1][0] * m[2][1] - m[1][1] * m[2][0];
1019     T a = m[1][1] * A2323 - m[1][2] * A1323 + m[1][3] * A1223;
1020     T b = m[1][0] * A2323 - m[1][2] * A0323 + m[1][3] * A0223;
1021     T c = m[1][0] * A1323 - m[1][1] * A0323 + m[1][3] * A0123;
1022     T d = m[1][0] * A1223 - m[1][1] * A0223 + m[1][2] * A0123;
1023     T det = m[0][0] * ( a )
1024           - m[0][1] * ( b )
1025           + m[0][2] * ( c )
1026           - m[0][3] * ( d );
1027     det = T(1) / det;
1028     metal::matrix<T, 4, 4> im;
1029     im[0][0] = det *   ( a );
1030     im[0][1] = det * - ( m[0][1] * A2323 - m[0][2] * A1323 + m[0][3] * A1223 );
1031     im[0][2] = det *   ( m[0][1] * A2313 - m[0][2] * A1313 + m[0][3] * A1213 );
1032     im[0][3] = det * - ( m[0][1] * A2312 - m[0][2] * A1312 + m[0][3] * A1212 );
1033     im[1][0] = det * - ( b );
1034     im[1][1] = det *   ( m[0][0] * A2323 - m[0][2] * A0323 + m[0][3] * A0223 );
1035     im[1][2] = det * - ( m[0][0] * A2313 - m[0][2] * A0313 + m[0][3] * A0213 );
1036     im[1][3] = det *   ( m[0][0] * A2312 - m[0][2] * A0312 + m[0][3] * A0212 );
1037     im[2][0] = det *   ( c );
1038     im[2][1] = det * - ( m[0][0] * A1323 - m[0][1] * A0323 + m[0][3] * A0123 );
1039     im[2][2] = det *   ( m[0][0] * A1313 - m[0][1] * A0313 + m[0][3] * A0113 );
1040     im[2][3] = det * - ( m[0][0] * A1312 - m[0][1] * A0312 + m[0][3] * A0112 );
1041     im[3][0] = det * - ( d );
1042     im[3][1] = det *   ( m[0][0] * A1223 - m[0][1] * A0223 + m[0][2] * A0123 );
1043     im[3][2] = det * - ( m[0][0] * A1213 - m[0][1] * A0213 + m[0][2] * A0113 );
1044     im[3][3] = det *   ( m[0][0] * A1212 - m[0][1] * A0212 + m[0][2] * A0112 );
1045     return im;
1046 }
1047 )")
1048 
1049 PROGRAM_PRELUDE_DECLARE(equalArray,
1050                         R"(
1051 template <typename T, size_t N>
1052 ANGLE_ALWAYS_INLINE bool ANGLE_equal(metal::array<T, N> u, metal::array<T, N> v)
1053 {
1054     for(size_t i = 0; i < N; i++)
1055         if (u[i] != v[i]) return false;
1056     return true;
1057 }
1058 )")
1059 
1060 PROGRAM_PRELUDE_DECLARE(equalStructArray,
1061                         R"(
1062 template <typename T, size_t N>
1063 ANGLE_ALWAYS_INLINE bool ANGLE_equalStructArray(metal::array<T, N> u, metal::array<T, N> v)
1064 {
1065     for(size_t i = 0; i < N; i++)
1066     {
1067         if (ANGLE_equal(u[i], v[i]) == false)
1068             return false;
1069     }
1070     return true;
1071 }
1072 )")
1073 
1074 PROGRAM_PRELUDE_DECLARE(notEqualArray,
1075                         R"(
1076 template <typename T, size_t N>
1077 ANGLE_ALWAYS_INLINE bool ANGLE_notEqual(metal::array<T, N> u, metal::array<T, N> v)
1078 {
1079     return !ANGLE_equal(u,v);
1080 }
1081 )",
1082                         equalArray())
1083 
1084 PROGRAM_PRELUDE_DECLARE(equalVector,
1085                         R"(
1086 template <typename T, int N>
1087 ANGLE_ALWAYS_INLINE bool ANGLE_equal(metal::vec<T, N> u, metal::vec<T, N> v)
1088 {
1089     return metal::all(u == v);
1090 }
1091 )",
1092                         include_metal_math())
1093 
1094 PROGRAM_PRELUDE_DECLARE(equalMatrix,
1095                         R"(
1096 template <typename T, int C, int R>
1097 ANGLE_ALWAYS_INLINE bool ANGLE_equal(metal::matrix<T, C, R> a, metal::matrix<T, C, R> b)
1098 {
1099     for (int c = 0; c < C; ++c)
1100     {
1101         if (!ANGLE_equal(a[c], b[c]))
1102         {
1103             return false;
1104         }
1105     }
1106     return true;
1107 }
1108 )",
1109                         equalVector())
1110 
1111 PROGRAM_PRELUDE_DECLARE(notEqualMatrix,
1112                         R"(
1113 template <typename T, int C, int R>
1114 ANGLE_ALWAYS_INLINE bool ANGLE_notEqual(metal::matrix<T, C, R> u, metal::matrix<T, C, R> v)
1115 {
1116     return !ANGLE_equal(u, v);
1117 }
1118 )",
1119                         equalMatrix())
1120 
1121 PROGRAM_PRELUDE_DECLARE(notEqualVector,
1122                         R"(
1123 template <typename T, int N>
1124 ANGLE_ALWAYS_INLINE bool ANGLE_notEqual(metal::vec<T, N> u, metal::vec<T, N> v)
1125 {
1126     return !ANGLE_equal(u, v);
1127 }
1128 )",
1129                         equalVector())
1130 
1131 PROGRAM_PRELUDE_DECLARE(notEqualStruct,
1132                         R"(
1133 template <typename T>
1134 ANGLE_ALWAYS_INLINE bool ANGLE_notEqualStruct(thread const T &a, thread const T &b)
1135 {
1136     return !ANGLE_equal(a, b);
1137 }
1138 )",
1139                         equalVector(),
1140                         equalMatrix())
1141 
1142 PROGRAM_PRELUDE_DECLARE(notEqualStructArray,
1143                         R"(
1144 template <typename T, size_t N>
1145 ANGLE_ALWAYS_INLINE bool ANGLE_notEqualStructArray(metal::array<T, N> u, metal::array<T, N> v)
1146 {
1147     for(size_t i = 0; i < N; i++)
1148     {
1149         if (ANGLE_notEqualStruct(u[i], v[i]))
1150             return true;
1151     }
1152     return false;
1153 }
1154 )",
1155                         notEqualStruct())
1156 
1157 PROGRAM_PRELUDE_DECLARE(vectorElemRef, R"(
1158 template <typename T, int N>
1159 struct ANGLE_VectorElemRef
1160 {
1161     thread metal::vec<T, N> &mVec;
1162     T mRef;
1163     const int mIndex;
1164     ~ANGLE_VectorElemRef() { mVec[mIndex] = mRef; }
1165     ANGLE_VectorElemRef(thread metal::vec<T, N> &vec, int index)
1166         : mVec(vec), mRef(vec[index]), mIndex(index)
1167     {}
1168     operator thread T &() { return mRef; }
1169 };
1170 template <typename T, int N>
1171 ANGLE_ALWAYS_INLINE ANGLE_VectorElemRef<T, N> ANGLE_elem_ref(thread metal::vec<T, N> &vec, int index)
1172 {
1173     return ANGLE_VectorElemRef<T, N>(vec, index);
1174 }
1175 )")
1176 
1177 PROGRAM_PRELUDE_DECLARE(swizzleRef,
1178                         R"(
1179 template <typename T, int VN, int SN>
1180 struct ANGLE_SwizzleRef
1181 {
1182     thread metal::vec<T, VN> &mVec;
1183     metal::vec<T, SN> mRef;
1184     int mIndices[SN];
1185     ~ANGLE_SwizzleRef()
1186     {
1187         for (int i = 0; i < SN; ++i)
1188         {
1189             const int j = mIndices[i];
1190             mVec[j] = mRef[i];
1191         }
1192     }
1193     ANGLE_SwizzleRef(thread metal::vec<T, VN> &vec, thread const int *indices)
1194         : mVec(vec)
1195     {
1196         for (int i = 0; i < SN; ++i)
1197         {
1198             const int j = indices[i];
1199             mIndices[i] = j;
1200             mRef[i] = mVec[j];
1201         }
1202     }
1203     operator thread metal::vec<T, SN> &() { return mRef; }
1204 };
1205 template <typename T, int N>
1206 ANGLE_ALWAYS_INLINE ANGLE_VectorElemRef<T, N> ANGLE_swizzle_ref(thread metal::vec<T, N> &vec, int i0)
1207 {
1208     const int is[] = { i0 };
1209     return ANGLE_VectorElemRef<T, N>(vec, is);
1210 }
1211 template <typename T, int N>
1212 ANGLE_ALWAYS_INLINE ANGLE_SwizzleRef<T, N, 2> ANGLE_swizzle_ref(thread metal::vec<T, N> &vec, int i0, int i1)
1213 {
1214     const int is[] = { i0, i1 };
1215     return ANGLE_SwizzleRef<T, N, 2>(vec, is);
1216 }
1217 template <typename T, int N>
1218 ANGLE_ALWAYS_INLINE ANGLE_SwizzleRef<T, N, 3> ANGLE_swizzle_ref(thread metal::vec<T, N> &vec, int i0, int i1, int i2)
1219 {
1220     const int is[] = { i0, i1, i2 };
1221     return ANGLE_SwizzleRef<T, N, 3>(vec, is);
1222 }
1223 template <typename T, int N>
1224 ANGLE_ALWAYS_INLINE ANGLE_SwizzleRef<T, N, 4> ANGLE_swizzle_ref(thread metal::vec<T, N> &vec, int i0, int i1, int i2, int i3)
1225 {
1226     const int is[] = { i0, i1, i2, i3 };
1227     return ANGLE_SwizzleRef<T, N, 4>(vec, is);
1228 }
1229 )",
1230                         vectorElemRef())
1231 
1232 PROGRAM_PRELUDE_DECLARE(out, R"(
1233 template <typename T>
1234 struct ANGLE_Out
1235 {
1236     T mTemp;
1237     thread T &mDest;
1238     ~ANGLE_Out() { mDest = mTemp; }
1239     ANGLE_Out(thread T &dest)
1240         : mTemp(dest), mDest(dest)
1241     {}
1242     operator thread T &() { return mTemp; }
1243 };
1244 template <typename T>
1245 ANGLE_ALWAYS_INLINE ANGLE_Out<T> ANGLE_out(thread T &dest)
1246 {
1247     return ANGLE_Out<T>(dest);
1248 }
1249 )")
1250 
1251 PROGRAM_PRELUDE_DECLARE(inout, R"(
1252 template <typename T>
1253 struct ANGLE_InOut
1254 {
1255     T mTemp;
1256     thread T &mDest;
1257     ~ANGLE_InOut() { mDest = mTemp; }
1258     ANGLE_InOut(thread T &dest)
1259         : mTemp(dest), mDest(dest)
1260     {}
1261     operator thread T &() { return mTemp; }
1262 };
1263 template <typename T>
1264 ANGLE_ALWAYS_INLINE ANGLE_InOut<T> ANGLE_inout(thread T &dest)
1265 {
1266     return ANGLE_InOut<T>(dest);
1267 }
1268 )")
1269 
1270 PROGRAM_PRELUDE_DECLARE(flattenArray, R"(
1271 template <typename T>
1272 struct ANGLE_flatten_impl
1273 {
1274     static ANGLE_ALWAYS_INLINE thread T *exec(thread T &x)
1275     {
1276         return &x;
1277     }
1278 };
1279 template <typename T, size_t N>
1280 struct ANGLE_flatten_impl<metal::array<T, N>>
1281 {
1282     static ANGLE_ALWAYS_INLINE auto exec(thread metal::array<T, N> &arr) -> T
1283     {
1284         return ANGLE_flatten_impl<T>::exec(arr[0]);
1285     }
1286 };
1287 template <typename T, size_t N>
1288 ANGLE_ALWAYS_INLINE auto ANGLE_flatten(thread metal::array<T, N> &arr) -> T
1289 {
1290     return ANGLE_flatten_impl<T>::exec(arr[0]);
1291 }
1292 )")
1293 
1294 PROGRAM_PRELUDE_DECLARE(castVector, R"(
1295 template <typename T, int N1, int N2>
1296 struct ANGLE_castVector {};
1297 template <typename T, int N>
1298 struct ANGLE_castVector<T, N, N>
1299 {
1300     static ANGLE_ALWAYS_INLINE metal::vec<T, N> exec(thread metal::vec<T, N> const &v)
1301     {
1302         return v;
1303     }
1304 };
1305 template <typename T>
1306 struct ANGLE_castVector<T, 2, 3>
1307 {
1308     static ANGLE_ALWAYS_INLINE metal::vec<T, 2> exec(thread metal::vec<T, 3> const &v)
1309     {
1310         return v.xy;
1311     }
1312 };
1313 template <typename T>
1314 struct ANGLE_castVector<T, 2, 4>
1315 {
1316     static ANGLE_ALWAYS_INLINE metal::vec<T, 2> exec(thread metal::vec<T, 4> const &v)
1317     {
1318         return v.xy;
1319     }
1320 };
1321 template <typename T>
1322 struct ANGLE_castVector<T, 3, 4>
1323 {
1324     static ANGLE_ALWAYS_INLINE metal::vec<T, 3> exec(thread metal::vec<T, 4> const &v)
1325     {
1326         return as_type<metal::vec<T, 3>>(v);
1327     }
1328 };
1329 template <int N1, int N2, typename T>
1330 ANGLE_ALWAYS_INLINE metal::vec<T, N1> ANGLE_cast(thread metal::vec<T, N2> const &v)
1331 {
1332     return ANGLE_castVector<T, N1, N2>::exec(v);
1333 }
1334 )")
1335 
1336 PROGRAM_PRELUDE_DECLARE(castMatrix,
1337                         R"(
1338 template <typename T, int C1, int R1, int C2, int R2, typename Enable = void>
1339 struct ANGLE_castMatrix
1340 {
1341     static ANGLE_ALWAYS_INLINE metal::matrix<T, C1, R1> exec(thread metal::matrix<T, C2, R2> const &m2)
1342     {
1343         metal::matrix<T, C1, R1> m1;
1344         const int MinC = C1 <= C2 ? C1 : C2;
1345         const int MinR = R1 <= R2 ? R1 : R2;
1346         for (int c = 0; c < MinC; ++c)
1347         {
1348             for (int r = 0; r < MinR; ++r)
1349             {
1350                 m1[c][r] = m2[c][r];
1351             }
1352             for (int r = R2; r < R1; ++r)
1353             {
1354                 m1[c][r] = c == r ? T(1) : T(0);
1355             }
1356         }
1357         for (int c = C2; c < C1; ++c)
1358         {
1359             for (int r = 0; r < R1; ++r)
1360             {
1361                 m1[c][r] = c == r ? T(1) : T(0);
1362             }
1363         }
1364         return m1;
1365     }
1366 };
1367 template <typename T, int C1, int R1, int C2, int R2>
1368 struct ANGLE_castMatrix<T, C1, R1, C2, R2, ANGLE_enable_if_t<(C1 <= C2 && R1 <= R2)>>
1369 {
1370     static ANGLE_ALWAYS_INLINE metal::matrix<T, C1, R1> exec(thread metal::matrix<T, C2, R2> const &m2)
1371     {
1372         metal::matrix<T, C1, R1> m1;
1373         for (size_t c = 0; c < C1; ++c)
1374         {
1375             m1[c] = ANGLE_cast<R1>(m2[c]);
1376         }
1377         return m1;
1378     }
1379 };
1380 template <int C1, int R1, int C2, int R2, typename T>
1381 ANGLE_ALWAYS_INLINE metal::matrix<T, C1, R1> ANGLE_cast(thread metal::matrix<T, C2, R2> const &m)
1382 {
1383     return ANGLE_castMatrix<T, C1, R1, C2, R2>::exec(m);
1384 };
1385 )",
1386                         enable_if(),
1387                         castVector())
1388 
1389 PROGRAM_PRELUDE_DECLARE(tensor, R"(
1390 template <typename T, size_t... DS>
1391 struct ANGLE_tensor_traits;
1392 template <typename T, size_t D>
1393 struct ANGLE_tensor_traits<T, D>
1394 {
1395     enum : size_t { outer_dim = D };
1396     using inner_type = T;
1397     using outer_type = inner_type[D];
1398 };
1399 template <typename T, size_t D, size_t... DS>
1400 struct ANGLE_tensor_traits<T, D, DS...>
1401 {
1402     enum : size_t { outer_dim = D };
1403     using inner_type = typename ANGLE_tensor_traits<T, DS...>::outer_type;
1404     using outer_type = inner_type[D];
1405 };
1406 template <size_t D, typename value_type_, typename inner_type_>
1407 struct ANGLE_tensor_impl
1408 {
1409     enum : size_t { outer_dim = D };
1410     using value_type = value_type_;
1411     using inner_type = inner_type_;
1412     using outer_type = inner_type[D];
1413     outer_type _data;
1414     ANGLE_ALWAYS_INLINE size_t size() const { return outer_dim; }
1415     ANGLE_ALWAYS_INLINE inner_type &operator[](size_t i) { return _data[i]; }
1416     ANGLE_ALWAYS_INLINE const inner_type &operator[](size_t i) const { return _data[i]; }
1417 };
1418 template <typename T, size_t... DS>
1419 using ANGLE_tensor = ANGLE_tensor_impl<
1420     ANGLE_tensor_traits<T, DS...>::outer_dim,
1421     T,
1422     typename ANGLE_tensor_traits<T, DS...>::inner_type>;
1423 )")
1424 
1425 PROGRAM_PRELUDE_DECLARE(gradient,
1426                         R"(
1427 template <int N>
1428 struct ANGLE_gradient_traits;
1429 template <>
1430 struct ANGLE_gradient_traits<2> { using type = metal::gradient2d; };
1431 template <>
1432 struct ANGLE_gradient_traits<3> { using type = metal::gradient3d; };
1433 
1434 template <int N>
1435 using ANGLE_gradient = typename ANGLE_gradient_traits<N>::type;
1436 )")
1437 
1438 PROGRAM_PRELUDE_DECLARE(writeSampleMask,
1439                         R"(
1440 ANGLE_ALWAYS_INLINE void ANGLE_writeSampleMask(const uint mask,
1441                                                thread uint& gl_SampleMask)
1442 {
1443     if (ANGLECoverageMaskEnabled)
1444     {
1445         gl_SampleMask = as_type<int>(mask);
1446     }
1447 }
1448 )",
1449                         functionConstants())
1450 
1451 PROGRAM_PRELUDE_DECLARE(textureEnv,
1452                         R"(
1453 template <typename T>
1454 struct ANGLE_TextureEnv
1455 {
1456     thread T *texture;
1457     thread metal::sampler *sampler;
1458 };
1459 )")
1460 
1461 // Note: for the time being, names must match those in TranslatorMetal.
1462 PROGRAM_PRELUDE_DECLARE(functionConstants,
1463                         R"(
1464 #define ANGLE_SAMPLE_COMPARE_GRADIENT_INDEX 0
1465 #define ANGLE_SAMPLE_COMPARE_LOD_INDEX      1
1466 #define ANGLE_RASTERIZATION_DISCARD_INDEX   2
1467 #define ANGLE_COVERAGE_MASK_ENABLED_INDEX   3
1468 
1469 constant bool ANGLEUseSampleCompareGradient [[function_constant(ANGLE_SAMPLE_COMPARE_GRADIENT_INDEX)]];
1470 constant bool ANGLEUseSampleCompareLod      [[function_constant(ANGLE_SAMPLE_COMPARE_LOD_INDEX)]];
1471 constant bool ANGLERasterizerDisabled       [[function_constant(ANGLE_RASTERIZATION_DISCARD_INDEX)]];
1472 constant bool ANGLECoverageMaskEnabled      [[function_constant(ANGLE_COVERAGE_MASK_ENABLED_INDEX)]];
1473 )")
1474 
1475 PROGRAM_PRELUDE_DECLARE(texelFetch,
1476                         R"(
1477 #define ANGLE_texelFetch(env, ...) ANGLE_texelFetch_impl(*env.texture, __VA_ARGS__)
1478 
1479 template <typename Texture>
1480 ANGLE_ALWAYS_INLINE auto ANGLE_texelFetch_impl(
1481     thread Texture &texture,
1482     metal::int2 const coord,
1483     uint level)
1484 {
1485     return texture.read(uint2(coord), level);
1486 }
1487 
1488 template <typename Texture>
1489 ANGLE_ALWAYS_INLINE auto ANGLE_texelFetch_impl(
1490     thread Texture &texture,
1491     metal::int3 const coord,
1492     uint level)
1493 {
1494     return texture.read(uint3(coord), level);
1495 }
1496 
1497 template <typename T>
1498 ANGLE_ALWAYS_INLINE auto ANGLE_texelFetch_impl(
1499     thread metal::texture2d_array<T> &texture,
1500     metal::int3 const coord,
1501     uint level)
1502 {
1503     return texture.read(uint2(coord.xy), uint(coord.z), level);
1504 }
1505 )",
1506                         textureEnv())
1507 
1508 PROGRAM_PRELUDE_DECLARE(texelFetchOffset,
1509                         R"(
1510 #define ANGLE_texelFetchOffset(env, ...) ANGLE_texelFetchOffset_impl(*env.texture, __VA_ARGS__)
1511 
1512 template <typename Texture>
1513 ANGLE_ALWAYS_INLINE auto ANGLE_texelFetchOffset_impl(
1514     thread Texture &texture,
1515     metal::int2 const coord,
1516     uint level,
1517     metal::int2 const offset)
1518 {
1519     return texture.read(uint2(coord + offset), level);
1520 }
1521 
1522 template <typename Texture>
1523 ANGLE_ALWAYS_INLINE auto ANGLE_texelFetchOffset_impl(
1524     thread Texture &texture,
1525     metal::int3 const coord,
1526     uint level,
1527     metal::int3 const offset)
1528 {
1529     return texture.read(uint3(coord + offset), level);
1530 }
1531 
1532 template <typename T>
1533 ANGLE_ALWAYS_INLINE auto ANGLE_texelFetchOffset_impl(
1534     thread metal::texture2d_array<T> &texture,
1535     metal::int3 const coord,
1536     uint level,
1537     metal::int2 const offset)
1538 {
1539     return texture.read(uint2(coord.xy + offset), uint(coord.z), level);
1540 }
1541 )",
1542                         textureEnv())
1543 
1544 PROGRAM_PRELUDE_DECLARE(texture,
1545                         R"(
1546 #define ANGLE_texture(env, ...) ANGLE_texture_impl(*env.texture, *env.sampler, __VA_ARGS__)
1547 )",
1548                         textureEnv())
1549 
1550 PROGRAM_PRELUDE_DECLARE(texture_generic_float2,
1551                         R"(
1552 template <typename Texture>
1553 ANGLE_ALWAYS_INLINE auto ANGLE_texture_impl(
1554     thread Texture &texture,
1555     thread metal::sampler const &sampler,
1556     metal::float2 const coord)
1557 {
1558     return texture.sample(sampler, coord);
1559 }
1560 )",
1561                         texture())
1562 
1563 PROGRAM_PRELUDE_DECLARE(texture_generic_float2_float,
1564                         R"(
1565 template <typename Texture>
1566 ANGLE_ALWAYS_INLINE auto ANGLE_texture_impl(
1567     thread Texture &texture,
1568     thread metal::sampler const &sampler,
1569     metal::float2 const coord,
1570     float bias)
1571 {
1572     return texture.sample(sampler, coord, metal::bias(bias));
1573 }
1574 )",
1575                         texture())
1576 
1577 PROGRAM_PRELUDE_DECLARE(texture_generic_float3,
1578                         R"(
1579 template <typename Texture>
1580 ANGLE_ALWAYS_INLINE auto ANGLE_texture_impl(
1581     thread Texture &texture,
1582     thread metal::sampler const &sampler,
1583     metal::float3 const coord)
1584 {
1585     return texture.sample(sampler, coord);
1586 }
1587 )",
1588                         texture())
1589 
1590 PROGRAM_PRELUDE_DECLARE(texture_generic_float3_float,
1591                         R"(
1592 template <typename Texture>
1593 ANGLE_ALWAYS_INLINE auto ANGLE_texture_impl(
1594     thread Texture &texture,
1595     thread metal::sampler const &sampler,
1596     metal::float3 const coord,
1597     float bias)
1598 {
1599     return texture.sample(sampler, coord, metal::bias(bias));
1600 }
1601 )",
1602                         texture())
1603 
1604 PROGRAM_PRELUDE_DECLARE(texture_depth2d_float3,
1605                         R"(
1606 template <typename T>
1607 ANGLE_ALWAYS_INLINE auto ANGLE_texture_impl(
1608     thread metal::depth2d<T> &texture,
1609     thread metal::sampler const &sampler,
1610     metal::float3 const coord)
1611 {
1612     return texture.sample_compare(sampler, coord.xy, coord.z);
1613 }
1614 )",
1615                         texture())
1616 
1617 PROGRAM_PRELUDE_DECLARE(texture_depth2d_float3_float,
1618                         R"(
1619 template <typename T>
1620 ANGLE_ALWAYS_INLINE auto ANGLE_texture_impl(
1621     thread metal::depth2d<T> &texture,
1622     thread metal::sampler const &sampler,
1623     metal::float3 const coord,
1624     float bias)
1625 {
1626     return texture.sample_compare(sampler, coord.xy, coord.z, metal::bias(bias));
1627 }
1628 )",
1629                         texture())
1630 
1631 PROGRAM_PRELUDE_DECLARE(texture_depth2darray_float4,
1632                         R"(
1633 template <typename T>
1634 ANGLE_ALWAYS_INLINE auto ANGLE_texture_impl(
1635     thread metal::depth2d_array<T> &texture,
1636     thread metal::sampler const &sampler,
1637     metal::float4 const coord)
1638 {
1639     return texture.sample_compare(sampler, coord.xy, uint(metal::round(coord.z)), coord.w);
1640 }
1641 )",
1642                         texture())
1643 
1644 PROGRAM_PRELUDE_DECLARE(texture_depth2darray_float4_float,
1645                         R"(
1646 template <typename T>
1647 ANGLE_ALWAYS_INLINE auto ANGLE_texture_impl(
1648     thread metal::depth2d_array<T> &texture,
1649     thread metal::sampler const &sampler,
1650     metal::float4 const coord,
1651     float compare)
1652 {
1653     return texture.sample_compare(sampler, coord.xyz, uint(metal::round(coord.w)), compare);
1654 }
1655 )",
1656                         texture())
1657 
1658 PROGRAM_PRELUDE_DECLARE(texture_depthcube_float4,
1659                         R"(
1660 template <typename T>
1661 ANGLE_ALWAYS_INLINE auto ANGLE_texture_impl(
1662     thread metal::depthcube<T> &texture,
1663     thread metal::sampler const &sampler,
1664     metal::float4 const coord)
1665 {
1666     return texture.sample_compare(sampler, coord.xyz, coord.w);
1667 }
1668 )",
1669                         texture())
1670 
1671 PROGRAM_PRELUDE_DECLARE(texture_depthcube_float4_float,
1672                         R"(
1673 template <typename T>
1674 ANGLE_ALWAYS_INLINE auto ANGLE_texture_impl(
1675     thread metal::depthcube<T> &texture,
1676     thread metal::sampler const &sampler,
1677     metal::float4 const coord,
1678     float bias)
1679 {
1680     return texture.sample_compare(sampler, coord.xyz, coord.w, metal::bias(bias));
1681 }
1682 )",
1683                         texture())
1684 
1685 PROGRAM_PRELUDE_DECLARE(texture_texture2darray_float3,
1686                         R"(
1687 template <typename T>
1688 ANGLE_ALWAYS_INLINE auto ANGLE_texture_impl(
1689     thread metal::texture2d_array<T> &texture,
1690     thread metal::sampler const &sampler,
1691     metal::float3 const coord)
1692 {
1693     return texture.sample(sampler, coord.xy, uint(metal::round(coord.z)));
1694 }
1695 )",
1696                         texture())
1697 
1698 PROGRAM_PRELUDE_DECLARE(texture_texture2darray_float3_float,
1699                         R"(
1700 template <typename T>
1701 ANGLE_ALWAYS_INLINE auto ANGLE_texture_impl(
1702     thread metal::texture2d_array<T> &texture,
1703     thread metal::sampler const &sampler,
1704     metal::float3 const coord,
1705     float bias)
1706 {
1707     return texture.sample(sampler, coord.xy, uint(metal::round(coord.z)), metal::bias(bias));
1708 }
1709 )",
1710                         texture())
1711 
1712 PROGRAM_PRELUDE_DECLARE(texture_texture2darray_float4,
1713                         R"(
1714 template <typename T>
1715 ANGLE_ALWAYS_INLINE auto ANGLE_texture_impl(
1716     thread metal::texture2d_array<T> &texture,
1717     thread metal::sampler const &sampler,
1718     metal::float4 const coord)
1719 {
1720     return texture.sample(sampler, coord.xyz, uint(metal::round(coord.w)));
1721 }
1722 )",
1723                         texture())
1724 
1725 PROGRAM_PRELUDE_DECLARE(texture_texture2darray_float4_float,
1726                         R"(
1727 template <typename T>
1728 ANGLE_ALWAYS_INLINE auto ANGLE_texture_impl(
1729     thread metal::texture2d_array<T> &texture,
1730     thread metal::sampler const &sampler,
1731     metal::float4 const coord,
1732     float bias)
1733 {
1734     return texture.sample(sampler, coord.xyz, uint(metal::round(coord.w)), metal::bias(bias));
1735 }
1736 )",
1737                         texture())
1738 
1739 PROGRAM_PRELUDE_DECLARE(texture1DLod,
1740                         R"(
1741 #define ANGLE_texture1DLod(env, ...) ANGLE_texture1DLod_impl(*env.texture, *env.sampler, __VA_ARGS__)
1742 
1743 template <typename Texture>
1744 ANGLE_ALWAYS_INLINE auto ANGLE_texture1DLod_impl(
1745     thread Texture &texture,
1746     thread metal::sampler const &sampler,
1747     float const coord,
1748     float level)
1749 {
1750     return texture.sample(sampler, coord, metal::level(level));
1751 }
1752 )",
1753                         textureEnv())
1754 
1755 PROGRAM_PRELUDE_DECLARE(texture1DProj,
1756                         R"(
1757 #define ANGLE_texture1DProj(env, ...) ANGLE_texture1DProj_impl(*env.texture, *env.sampler, __VA_ARGS__)
1758 
1759 template <typename Texture>
1760 ANGLE_ALWAYS_INLINE auto ANGLE_texture1DProj_impl(
1761     thread Texture &texture,
1762     thread metal::sampler const &sampler,
1763     metal::float2 const coord,
1764     float bias = 0)
1765 {
1766     return texture.sample(sampler, coord.x/coord.y, metal::bias(bias));
1767 }
1768 
1769 template <typename Texture>
1770 ANGLE_ALWAYS_INLINE auto ANGLE_texture1DProj_impl(
1771     thread Texture &texture,
1772     thread metal::sampler const &sampler,
1773     metal::float4 const coord,
1774     float bias = 0)
1775 {
1776     return texture.sample(sampler, coord.x/coord.w, metal::bias(bias));
1777 }
1778 )",
1779                         textureEnv())
1780 
1781 PROGRAM_PRELUDE_DECLARE(texture1DProjLod,
1782                         R"(
1783 #define ANGLE_texture1DProjLod(env, ...) ANGLE_texture1DProjLod_impl(*env.texture, *env.sampler, __VA_ARGS__)
1784 
1785 template <typename Texture>
1786 ANGLE_ALWAYS_INLINE auto ANGLE_texture1DProjLod_impl(
1787     thread Texture &texture,
1788     thread metal::sampler const &sampler,
1789     metal::float2 const coord,
1790     float level)
1791 {
1792     return texture.sample(sampler, coord.x/coord.y, metal::level(level));
1793 }
1794 
1795 template <typename Texture>
1796 ANGLE_ALWAYS_INLINE auto ANGLE_texture1DProjLod_impl(
1797     thread Texture &texture,
1798     thread metal::sampler const &sampler,
1799     metal::float4 const coord,
1800     float level)
1801 {
1802     return texture.sample(sampler, coord.x/coord.w, metal::level(level));
1803 }
1804 )",
1805                         textureEnv())
1806 
1807 PROGRAM_PRELUDE_DECLARE(texture2D,
1808                         R"(
1809 #define ANGLE_texture2D(env, ...) ANGLE_texture2D_impl(*env.texture, *env.sampler, __VA_ARGS__)
1810 
1811 template <typename Texture>
1812 ANGLE_ALWAYS_INLINE auto ANGLE_texture2D_impl(
1813     thread Texture &texture,
1814     thread metal::sampler const &sampler,
1815     metal::float2 const coord)
1816 {
1817     return texture.sample(sampler, coord);
1818 }
1819 
1820 template <typename Texture>
1821 ANGLE_ALWAYS_INLINE auto ANGLE_texture2D_impl(
1822     thread Texture &texture,
1823     thread metal::sampler const &sampler,
1824     metal::float2 const coord,
1825     float bias)
1826 {
1827     return texture.sample(sampler, coord, metal::bias(bias));
1828 }
1829 
1830 template <typename Texture>
1831 ANGLE_ALWAYS_INLINE auto ANGLE_texture2D_impl(
1832     thread Texture &texture,
1833     thread metal::sampler const &sampler,
1834     metal::float3 const coord)
1835 {
1836     return texture.sample(sampler, coord);
1837 }
1838 
1839 template <typename Texture>
1840 ANGLE_ALWAYS_INLINE auto ANGLE_texture2D_impl(
1841     thread Texture &texture,
1842     thread metal::sampler const &sampler,
1843     metal::float3 const coord,
1844     float bias)
1845 {
1846     return texture.sample(sampler, coord, metal::bias(bias));
1847 }
1848 )",
1849                         textureEnv())
1850 
1851 PROGRAM_PRELUDE_DECLARE(texture2DRect,
1852                         R"(
1853 #define ANGLE_texture2DRect(env, ...) ANGLE_texture2DRect_impl(*env.texture, *env.sampler, __VA_ARGS__)
1854 template <typename Texture>
1855 ANGLE_ALWAYS_INLINE auto ANGLE_texture2DRect_impl(
1856     thread Texture &texture,
1857     thread metal::sampler const &sampler,
1858     metal::float2 const coord)
1859 {
1860     return texture.sample(sampler, coord);
1861 }
1862 )",
1863                         textureEnv())
1864 
1865 PROGRAM_PRELUDE_DECLARE(texture2DLod,
1866                         R"(
1867 #define ANGLE_texture2DLod(env, ...) ANGLE_texture2DLod_impl(*env.texture, *env.sampler, __VA_ARGS__)
1868 
1869 template <typename Texture>
1870 ANGLE_ALWAYS_INLINE auto ANGLE_texture2DLod_impl(
1871     thread Texture &texture,
1872     thread metal::sampler const &sampler,
1873     metal::float2 const coord,
1874     float level)
1875 {
1876     return texture.sample(sampler, coord, metal::level(level));
1877 }
1878 )",
1879                         textureEnv())
1880 
1881 PROGRAM_PRELUDE_DECLARE(texture2DProj,
1882                         R"(
1883 #define ANGLE_texture2DProj(env, ...) ANGLE_texture2DProj_impl(*env.texture, *env.sampler, __VA_ARGS__)
1884 
1885 template <typename Texture>
1886 ANGLE_ALWAYS_INLINE auto ANGLE_texture2DProj_impl(
1887     thread Texture &texture,
1888     thread metal::sampler const &sampler,
1889     metal::float3 const coord,
1890     float bias = 0)
1891 {
1892     return texture.sample(sampler, coord.xy/coord.z, metal::bias(bias));
1893 }
1894 
1895 template <typename Texture>
1896 ANGLE_ALWAYS_INLINE auto ANGLE_texture2DProj_impl(
1897     thread Texture &texture,
1898     thread metal::sampler const &sampler,
1899     metal::float4 const coord,
1900     float bias = 0)
1901 {
1902     return texture.sample(sampler, coord.xy/coord.w, metal::bias(bias));
1903 }
1904 )",
1905                         textureEnv())
1906 
1907 PROGRAM_PRELUDE_DECLARE(texture2DRectProj,
1908                         R"(
1909 #define ANGLE_texture2DRectProj(env, ...) ANGLE_texture2DRectProj_impl(*env.texture, *env.sampler, __VA_ARGS__)
1910 
1911 template <typename Texture>
1912 ANGLE_ALWAYS_INLINE auto ANGLE_texture2DRectProj_impl(
1913     thread Texture &texture,
1914     thread metal::sampler const &sampler,
1915     metal::float3 const coord)
1916 {
1917     return texture.sample(sampler, coord.xy/coord.z);
1918 }
1919 template <typename Texture>
1920 ANGLE_ALWAYS_INLINE auto ANGLE_texture2DRectProj_impl(
1921     thread Texture &texture,
1922     thread metal::sampler const &sampler,
1923     metal::float4 const coord)
1924 {
1925     return texture.sample(sampler, coord.xy/coord.w);
1926 }
1927 )",
1928                         textureEnv())
1929 
1930 PROGRAM_PRELUDE_DECLARE(texture2DProjLod,
1931                         R"(
1932 #define ANGLE_texture2DProjLod(env, ...) ANGLE_texture2DProjLod_impl(*env.texture, *env.sampler, __VA_ARGS__)
1933 
1934 template <typename Texture>
1935 ANGLE_ALWAYS_INLINE auto ANGLE_texture2DProjLod_impl(
1936     thread Texture &texture,
1937     thread metal::sampler const &sampler,
1938     metal::float3 const coord,
1939     float level)
1940 {
1941     return texture.sample(sampler, coord.xy/coord.z, metal::level(level));
1942 }
1943 
1944 template <typename Texture>
1945 ANGLE_ALWAYS_INLINE auto ANGLE_texture2DProjLod_impl(
1946     thread Texture &texture,
1947     thread metal::sampler const &sampler,
1948     metal::float4 const coord,
1949     float level)
1950 {
1951     return texture.sample(sampler, coord.xy/coord.w, metal::level(level));
1952 }
1953 )",
1954                         textureEnv())
1955 
1956 PROGRAM_PRELUDE_DECLARE(texture3DLod,
1957                         R"(
1958 #define ANGLE_texture3DLod(env, ...) ANGLE_texture3DLod_impl(*env.texture, *env.sampler, __VA_ARGS__)
1959 
1960 template <typename Texture>
1961 ANGLE_ALWAYS_INLINE auto ANGLE_texture3DLod_impl(
1962     thread Texture &texture,
1963     thread metal::sampler const &sampler,
1964     metal::float3 const coord,
1965     float level)
1966 {
1967     return texture.sample(sampler, coord, metal::level(level));
1968 }
1969 )",
1970                         textureEnv())
1971 
1972 PROGRAM_PRELUDE_DECLARE(texture3DProj,
1973                         R"(
1974 #define ANGLE_texture3DProj(env, ...) ANGLE_texture3DProj_impl(*env.texture, *env.sampler, __VA_ARGS__)
1975 
1976 template <typename Texture>
1977 ANGLE_ALWAYS_INLINE auto ANGLE_texture3DProj_impl(
1978     thread Texture &texture,
1979     thread metal::sampler const &sampler,
1980     metal::float4 const coord,
1981     float bias = 0)
1982 {
1983     return texture.sample(sampler, coord.xyz/coord.w, metal::bias(bias));
1984 }
1985 )",
1986                         textureEnv())
1987 
1988 PROGRAM_PRELUDE_DECLARE(texture3DProjLod,
1989                         R"(
1990 #define ANGLE_texture3DProjLod(env, ...) ANGLE_texture3DProjLod_impl(*env.texture, *env.sampler, __VA_ARGS__)
1991 
1992 template <typename Texture>
1993 ANGLE_ALWAYS_INLINE auto ANGLE_texture3DProjLod_impl(
1994     thread Texture &texture,
1995     thread metal::sampler const &sampler,
1996     metal::float4 const coord,
1997     float level)
1998 {
1999     return texture.sample(sampler, coord.xyz/coord.w, metal::level(level));
2000 }
2001 )",
2002                         textureEnv())
2003 
2004 PROGRAM_PRELUDE_DECLARE(textureCube,
2005                         R"(
2006 #define ANGLE_textureCube(env, ...) ANGLE_textureCube_impl(*env.texture, *env.sampler, __VA_ARGS__)
2007 
2008 template <typename Texture>
2009 ANGLE_ALWAYS_INLINE auto ANGLE_textureCube_impl(
2010     thread Texture &texture,
2011     thread metal::sampler const &sampler,
2012     metal::float3 const coord)
2013 {
2014     return texture.sample(sampler, coord);
2015 }
2016 
2017 template <typename Texture>
2018 ANGLE_ALWAYS_INLINE auto ANGLE_textureCube_impl(
2019     thread Texture &texture,
2020     thread metal::sampler const &sampler,
2021     metal::float3 const coord,
2022     float bias)
2023 {
2024     return texture.sample(sampler, coord, metal::bias(bias));
2025 }
2026 )",
2027                         textureEnv())
2028 
2029 PROGRAM_PRELUDE_DECLARE(textureCubeLod,
2030                         R"(
2031 #define ANGLE_textureCubeLod(env, ...) ANGLE_textureCubeLod_impl(*env.texture, *env.sampler, __VA_ARGS__)
2032 
2033 template <typename Texture>
2034 ANGLE_ALWAYS_INLINE auto ANGLE_textureCubeLod_impl(
2035     thread Texture &texture,
2036     thread metal::sampler const &sampler,
2037     metal::float3 const coord,
2038     float level)
2039 {
2040     return texture.sample(sampler, coord, metal::level(level));
2041 }
2042 )",
2043                         textureEnv())
2044 
2045 PROGRAM_PRELUDE_DECLARE(textureCubeProj,
2046                         R"(
2047 #define ANGLE_textureCubeProj(env, ...) ANGLE_textureCubeProj_impl(*env.texture, *env.sampler, __VA_ARGS__)
2048 
2049 template <typename Texture>
2050 ANGLE_ALWAYS_INLINE auto ANGLE_textureCubeProj_impl(
2051     thread Texture &texture,
2052     thread metal::sampler const &sampler,
2053     metal::float4 const coord,
2054     float bias = 0)
2055 {
2056     return texture.sample(sampler, coord.xyz/coord.w, metal::bias(bias));
2057 }
2058 )",
2059                         textureEnv())
2060 
2061 PROGRAM_PRELUDE_DECLARE(textureCubeProjLod,
2062                         R"(
2063 #define ANGLE_textureCubeProjLod(env, ...) ANGLE_textureCubeProjLod_impl(*env.texture, *env.sampler, __VA_ARGS__)
2064 
2065 template <typename Texture>
2066 ANGLE_ALWAYS_INLINE auto ANGLE_textureCubeProjLod_impl(
2067     thread Texture &texture,
2068     thread metal::sampler const &sampler,
2069     metal::float4 const coord,
2070     float level)
2071 {
2072     return texture.sample(sampler, coord.xyz/coord.w, metal::level(level));
2073 }
2074 )",
2075                         textureEnv())
2076 
2077 PROGRAM_PRELUDE_DECLARE(textureGrad,
2078                         R"(
2079 #define ANGLE_textureGrad(env, ...) ANGLE_textureGrad_impl(*env.texture, *env.sampler, __VA_ARGS__)
2080 )",
2081                         textureEnv())
2082 
2083 PROGRAM_PRELUDE_DECLARE(textureGrad_generic_floatN_floatN_floatN,
2084                         R"(
2085 template <typename Texture, int N>
2086 ANGLE_ALWAYS_INLINE auto ANGLE_textureGrad_impl(
2087     thread Texture &texture,
2088     thread metal::sampler const &sampler,
2089     metal::vec<float, N> const coord,
2090     metal::vec<float, N> const dPdx,
2091     metal::vec<float, N> const dPdy)
2092 {
2093     return texture.sample(sampler, coord, ANGLE_gradient<N>(dPdx, dPdy));
2094 }
2095 )",
2096                         gradient(),
2097                         textureGrad())
2098 
2099 PROGRAM_PRELUDE_DECLARE(textureGrad_generic_float3_float2_float2,
2100                         R"(
2101 template <typename Texture>
2102 ANGLE_ALWAYS_INLINE auto ANGLE_textureGrad_impl(
2103     thread Texture &texture,
2104     thread metal::sampler const &sampler,
2105     metal::float3 const coord,
2106     metal::float2 const dPdx,
2107     metal::float2 const dPdy)
2108 {
2109     return texture.sample(sampler, coord.xy, uint(metal::round(coord.z)), metal::gradient2d(dPdx, dPdy));
2110 }
2111 )",
2112                         textureGrad())
2113 
2114 PROGRAM_PRELUDE_DECLARE(textureGrad_generic_float4_float2_float2,
2115                         R"(
2116 template <typename Texture>
2117 ANGLE_ALWAYS_INLINE auto ANGLE_textureGrad_impl(
2118     thread Texture &texture,
2119     thread metal::sampler const &sampler,
2120     metal::float4 const coord,
2121     metal::float2 const dPdx,
2122     metal::float2 const dPdy)
2123 {
2124     return texture.sample(sampler, coord.xy, uint(metal::round(coord.z)), metal::gradient2d(dPdx, dPdy));
2125 }
2126 )",
2127                         textureGrad())
2128 
2129 PROGRAM_PRELUDE_DECLARE(textureGrad_depth2d_float3_float2_float2,
2130                         R"(
2131 template <typename T>
2132 ANGLE_ALWAYS_INLINE auto ANGLE_textureGrad_impl(
2133     thread metal::depth2d<T> &texture,
2134     thread metal::sampler const &sampler,
2135     metal::float3 const coord,
2136     metal::float2 const dPdx,
2137     metal::float2 const dPdy)
2138 {
2139     if (ANGLEUseSampleCompareGradient)
2140     {
2141         return static_cast<T>(texture.sample_compare(sampler, coord.xy, coord.z, metal::gradient2d(dPdx, dPdy)));
2142     }
2143     else
2144     {
2145         return static_cast<T>(texture.sample(sampler, coord.xy, metal::gradient2d(dPdx, dPdy)) > coord.z);
2146     }
2147 }
2148 )",
2149                         functionConstants(),
2150                         textureGrad())
2151 
2152 PROGRAM_PRELUDE_DECLARE(textureGrad_depth2darray_float4_float2_float2,
2153                         R"(
2154 template <typename T>
2155 ANGLE_ALWAYS_INLINE auto ANGLE_textureGrad_impl(
2156     thread metal::depth2d_array<T> &texture,
2157     thread metal::sampler const &sampler,
2158     metal::float4 const coord,
2159     metal::float2 const dPdx,
2160     metal::float2 const dPdy)
2161 {
2162     if (ANGLEUseSampleCompareGradient)
2163     {
2164         return static_cast<T>(texture.sample_compare(sampler, coord.xy, uint(metal::round(coord.z)), coord.w, metal::gradient2d(dPdx, dPdy)));
2165     }
2166     else
2167     {
2168         return static_cast<T>(texture.sample(sampler, coord.xy, uint(metal::round(coord.z)), metal::gradient2d(dPdx, dPdy)) > coord.w);
2169     }
2170 }
2171 )",
2172                         functionConstants(),
2173                         textureGrad())
2174 
2175 PROGRAM_PRELUDE_DECLARE(textureGrad_depthcube_float4_float3_float3,
2176                         R"(
2177 template <typename T>
2178 ANGLE_ALWAYS_INLINE auto ANGLE_textureGrad_impl(
2179     thread metal::depthcube<T> &texture,
2180     thread metal::sampler const &sampler,
2181     metal::float4 const coord,
2182     metal::float3 const dPdx,
2183     metal::float3 const dPdy)
2184 {
2185     if (ANGLEUseSampleCompareGradient)
2186     {
2187         return static_cast<T>(texture.sample_compare(sampler, coord.xyz, coord.w, metal::gradientcube(dPdx, dPdy)));
2188     }
2189     else
2190     {
2191         return static_cast<T>(texture.sample(sampler, coord.xyz, metal::gradientcube(dPdx, dPdy)) > coord.w);
2192     }
2193 }
2194 )",
2195                         functionConstants(),
2196                         textureGrad())
2197 
2198 PROGRAM_PRELUDE_DECLARE(textureGrad_texturecube_float3_float3_float3,
2199                         R"(
2200 template <typename T>
2201 ANGLE_ALWAYS_INLINE auto ANGLE_textureGrad_impl(
2202     thread metal::texturecube<T> &texture,
2203     thread metal::sampler const &sampler,
2204     metal::float3 const coord,
2205     metal::float3 const dPdx,
2206     metal::float3 const dPdy)
2207 {
2208     return texture.sample(sampler, coord, metal::gradientcube(dPdx, dPdy));
2209 }
2210 )",
2211                         textureGrad())
2212 
2213 PROGRAM_PRELUDE_DECLARE(textureGradOffset,
2214                         R"(
2215 #define ANGLE_textureGradOffset(env, ...) ANGLE_textureGradOffset_impl(*env.texture, *env.sampler, __VA_ARGS__)
2216 )",
2217                         textureEnv())
2218 
2219 PROGRAM_PRELUDE_DECLARE(textureGradOffset_generic_floatN_floatN_floatN_intN,
2220                         R"(
2221 template <typename Texture, int N>
2222 ANGLE_ALWAYS_INLINE auto ANGLE_textureGradOffset_impl(
2223     thread Texture &texture,
2224     thread metal::sampler const &sampler,
2225     metal::vec<float, N> const coord,
2226     metal::vec<float, N> const dPdx,
2227     metal::vec<float, N> const dPdy,
2228     metal::vec<int, N> const offset)
2229 {
2230     return texture.sample(sampler, coord, ANGLE_gradient<N>(dPdx, dPdy), offset);
2231 }
2232 )",
2233                         gradient(),
2234                         textureGradOffset())
2235 
2236 PROGRAM_PRELUDE_DECLARE(textureGradOffset_generic_float3_float2_float2_int2,
2237                         R"(
2238 template <typename Texture>
2239 ANGLE_ALWAYS_INLINE auto ANGLE_textureGradOffset_impl(
2240     thread Texture &texture,
2241     thread metal::sampler const &sampler,
2242     metal::float3 const coord,
2243     metal::float2 const dPdx,
2244     metal::float2 const dPdy,
2245     metal::int2 const offset)
2246 {
2247     return texture.sample(sampler, coord.xy, uint(metal::round(coord.z)), metal::gradient2d(dPdx, dPdy), offset);
2248 }
2249 )",
2250                         textureGradOffset())
2251 
2252 PROGRAM_PRELUDE_DECLARE(textureGradOffset_generic_float4_float2_float2_int2,
2253                         R"(
2254 template <typename Texture>
2255 ANGLE_ALWAYS_INLINE auto ANGLE_textureGradOffset_impl(
2256     thread Texture &texture,
2257     thread metal::sampler const &sampler,
2258     metal::float4 const coord,
2259     metal::float2 const dPdx,
2260     metal::float2 const dPdy,
2261     metal::int2 const offset)
2262 {
2263     return texture.sample(sampler, coord.xy, uint(metal::round(coord.z)), metal::gradient2d(dPdx, dPdy), offset);
2264 }
2265 )",
2266                         textureGradOffset())
2267 
2268 PROGRAM_PRELUDE_DECLARE(textureGradOffset_depth2d_float3_float2_float2_int2,
2269                         R"(
2270 template <typename T>
2271 ANGLE_ALWAYS_INLINE auto ANGLE_textureGradOffset_impl(
2272     thread metal::depth2d<T> &texture,
2273     thread metal::sampler const &sampler,
2274     metal::float3 const coord,
2275     metal::float2 const dPdx,
2276     metal::float2 const dPdy,
2277     metal::int2 const offset)
2278 {
2279     if (ANGLEUseSampleCompareGradient)
2280     {
2281         return static_cast<T>(texture.sample_compare(sampler, coord.xy, coord.z, metal::gradient2d(dPdx, dPdy), offset));
2282     }
2283     else
2284     {
2285         return static_cast<T>(texture.sample(sampler, coord.xy, metal::gradient2d(dPdx, dPdy), offset) > coord.z);
2286     }
2287 }
2288 )",
2289                         functionConstants(),
2290                         textureGradOffset())
2291 
2292 PROGRAM_PRELUDE_DECLARE(textureGradOffset_depth2darray_float4_float2_float2_int2,
2293                         R"(
2294 template <typename T>
2295 ANGLE_ALWAYS_INLINE auto ANGLE_textureGradOffset_impl(
2296     thread metal::depth2d_array<T> &texture,
2297     thread metal::sampler const &sampler,
2298     metal::float4 const coord,
2299     metal::float2 const dPdx,
2300     metal::float2 const dPdy,
2301     metal::int2 const offset)
2302 {
2303     if (ANGLEUseSampleCompareGradient)
2304     {
2305         return static_cast<T>(texture.sample_compare(sampler, coord.xy, uint(metal::round(coord.z)), coord.w, metal::gradient2d(dPdx, dPdy), offset));
2306     }
2307     else
2308     {
2309         return static_cast<T>(texture.sample(sampler, coord.xy, uint(metal::round(coord.z)), metal::gradient2d(dPdx, dPdy), offset) > coord.w);
2310     }
2311 }
2312 )",
2313                         functionConstants(),
2314                         textureGradOffset())
2315 
2316 PROGRAM_PRELUDE_DECLARE(textureGradOffset_depthcube_float4_float3_float3_int3,
2317                         R"(
2318 template <typename T>
2319 ANGLE_ALWAYS_INLINE auto ANGLE_textureGradOffset_impl(
2320     thread metal::depthcube<T> &texture,
2321     thread metal::sampler const &sampler,
2322     metal::float4 const coord,
2323     metal::float3 const dPdx,
2324     metal::float3 const dPdy,
2325     metal::int3 const offset)
2326 {
2327     return texture.sample_compare(sampler, coord.xyz, coord.w, metal::gradientcube(dPdx, dPdy), offset);
2328 }
2329 )",
2330                         textureGradOffset())
2331 
2332 PROGRAM_PRELUDE_DECLARE(textureGradOffset_texturecube_float3_float3_float3_int3,
2333                         R"(
2334 template <typename T>
2335 ANGLE_ALWAYS_INLINE auto ANGLE_textureGradOffset_impl(
2336     thread metal::texturecube<T> &texture,
2337     thread metal::sampler const &sampler,
2338     metal::float3 const coord,
2339     metal::float3 const dPdx,
2340     metal::float3 const dPdy,
2341     metal::int3 const offset)
2342 {
2343     return texture.sample(sampler, coord, metal::gradientcube(dPdx, dPdy), offset);
2344 }
2345 )",
2346                         textureGradOffset())
2347 
2348 PROGRAM_PRELUDE_DECLARE(textureLod,
2349                         R"(
2350 #define ANGLE_textureLod(env, ...) ANGLE_textureLod_impl(*env.texture, *env.sampler, __VA_ARGS__)
2351 )",
2352                         textureEnv())
2353 
2354 PROGRAM_PRELUDE_DECLARE(textureLod_generic_float2,
2355                         R"(
2356 template <typename Texture>
2357 ANGLE_ALWAYS_INLINE auto ANGLE_textureLod_impl(
2358     thread Texture &texture,
2359     thread metal::sampler const &sampler,
2360     metal::float2 const coord,
2361     float level)
2362 {
2363     return texture.sample(sampler, coord, metal::level(level));
2364 }
2365 )",
2366                         textureLod())
2367 
2368 PROGRAM_PRELUDE_DECLARE(textureLod_generic_float3,
2369                         R"(
2370 template <typename Texture>
2371 ANGLE_ALWAYS_INLINE auto ANGLE_textureLod_impl(
2372     thread Texture &texture,
2373     thread metal::sampler const &sampler,
2374     metal::float3 const coord,
2375     float level)
2376 {
2377     return texture.sample(sampler, coord, metal::level(level));
2378 }
2379 )",
2380                         textureLod())
2381 
2382 PROGRAM_PRELUDE_DECLARE(textureLod_depth2d_float3,
2383                         R"(
2384 template <typename T>
2385 ANGLE_ALWAYS_INLINE auto ANGLE_textureLod_impl(
2386     thread metal::depth2d<T> &texture,
2387     thread metal::sampler const &sampler,
2388     metal::float3 const coord,
2389     float level)
2390 {
2391     if (ANGLEUseSampleCompareLod)
2392     {
2393         return static_cast<T>(texture.sample_compare(sampler, coord.xy, coord.z, metal::level(level)));
2394     }
2395     else
2396     {
2397         return static_cast<T>(texture.sample(sampler, coord.xy, metal::level(level)) > coord.z);
2398     }
2399 }
2400 )",
2401                         functionConstants(),
2402                         textureLod())
2403 
2404 PROGRAM_PRELUDE_DECLARE(textureLod_texture2darray_float3,
2405                         R"(
2406 template <typename T>
2407 ANGLE_ALWAYS_INLINE auto ANGLE_textureLod_impl(
2408     thread metal::texture2d_array<T> &texture,
2409     thread metal::sampler const &sampler,
2410     metal::float3 const coord,
2411     float level)
2412 {
2413     return texture.sample(sampler, coord.xy, uint(metal::round(coord.z)), metal::level(level));
2414 }
2415 )",
2416                         textureLod())
2417 
2418 PROGRAM_PRELUDE_DECLARE(textureLod_texture2darray_float4,
2419                         R"(
2420 template <typename T>
2421 ANGLE_ALWAYS_INLINE auto ANGLE_textureLod_impl(
2422     thread metal::texture2d_array<T> &texture,
2423     thread metal::sampler const &sampler,
2424     metal::float4 const coord,
2425     float level)
2426 {
2427     return texture.sample(sampler, coord.xyz, uint(metal::round(coord.w)), metal::level(level));
2428 }
2429 )",
2430                         textureLod())
2431 
2432 PROGRAM_PRELUDE_DECLARE(textureLodOffset,
2433                         R"(
2434 #define ANGLE_textureLodOffset(env, ...) ANGLE_textureLodOffset_impl(*env.texture, *env.sampler, __VA_ARGS__)
2435 
2436 template <typename Texture>
2437 ANGLE_ALWAYS_INLINE auto ANGLE_textureLodOffset_impl(
2438     thread Texture &texture,
2439     thread metal::sampler const &sampler,
2440     metal::float2 const coord,
2441     float level,
2442     metal::int2 const offset)
2443 {
2444     return texture.sample(sampler, coord, metal::level(level), offset);
2445 }
2446 
2447 template <typename Texture>
2448 ANGLE_ALWAYS_INLINE auto ANGLE_textureLodOffset_impl(
2449     thread Texture &texture,
2450     thread metal::sampler const &sampler,
2451     metal::float3 const coord,
2452     float level,
2453     metal::int3 const offset)
2454 {
2455     return texture.sample(sampler, coord, metal::level(level), offset);
2456 }
2457 
2458 template <typename T>
2459 ANGLE_ALWAYS_INLINE auto ANGLE_textureLodOffset_impl(
2460     thread metal::depth2d<T> &texture,
2461     thread metal::sampler const &sampler,
2462     metal::float3 const coord,
2463     float level,
2464     int2 const offset)
2465 {
2466     if (ANGLEUseSampleCompareLod)
2467     {
2468         return static_cast<T>(texture.sample_compare(sampler, coord.xy, coord.z, metal::level(level), offset));
2469     }
2470     else
2471     {
2472         return static_cast<T>(texture.sample(sampler, coord.xy, metal::level(level), offset) > coord.z);
2473     }
2474 }
2475 
2476 template <typename T>
2477 ANGLE_ALWAYS_INLINE auto ANGLE_textureLodOffset_impl(
2478     thread metal::texture2d_array<T> &texture,
2479     thread metal::sampler const &sampler,
2480     metal::float3 const coord,
2481     float level,
2482     metal::int2 const offset)
2483 {
2484     return texture.sample(sampler, coord.xy, uint(metal::round(coord.z)), metal::level(level), offset);
2485 }
2486 
2487 template <typename T>
2488 ANGLE_ALWAYS_INLINE auto ANGLE_textureLodOffset_impl(
2489     thread metal::texture2d_array<T> &texture,
2490     thread metal::sampler const &sampler,
2491     metal::float4 const coord,
2492     float level,
2493     metal::int3 const offset)
2494 {
2495     return texture.sample(sampler, coord.xyz, uint(metal::round(coord.w)), metal::level(level), offset);
2496 }
2497 )",
2498                         functionConstants(),
2499                         textureEnv())
2500 
2501 PROGRAM_PRELUDE_DECLARE(textureOffset,
2502                         R"(
2503 #define ANGLE_textureOffset(env, ...) ANGLE_textureOffset_impl(*env.texture, *env.sampler, __VA_ARGS__)
2504 
2505 template <typename Texture>
2506 ANGLE_ALWAYS_INLINE auto ANGLE_textureOffset_impl(
2507     thread Texture &texture,
2508     thread metal::sampler const &sampler,
2509     metal::float2 const coord,
2510     metal::int2 const offset)
2511 {
2512     return texture.sample(sampler, coord, offset);
2513 }
2514 
2515 template <typename Texture>
2516 ANGLE_ALWAYS_INLINE auto ANGLE_textureOffset_impl(
2517     thread Texture &texture,
2518     thread metal::sampler const &sampler,
2519     metal::float2 const coord,
2520     metal::int2 const offset,
2521     float bias)
2522 {
2523     return texture.sample(sampler, coord, metal::bias(bias), offset);
2524 }
2525 
2526 template <typename Texture>
2527 ANGLE_ALWAYS_INLINE auto ANGLE_textureOffset_impl(
2528     thread Texture &texture,
2529     thread metal::sampler const &sampler,
2530     metal::float3 const coord,
2531     metal::int2 const offset)
2532 {
2533     return texture.sample(sampler, coord.xy, uint(metal::round(coord.z)), offset);
2534 }
2535 
2536 template <typename Texture>
2537 ANGLE_ALWAYS_INLINE auto ANGLE_textureOffset_impl(
2538     thread Texture &texture,
2539     thread metal::sampler const &sampler,
2540     metal::float3 const coord,
2541     metal::int2 const offset,
2542     float bias)
2543 {
2544     return texture.sample(sampler, coord.xy, uint(metal::round(coord.z)), metal::bias(bias), offset);
2545 }
2546 
2547 template <typename Texture>
2548 ANGLE_ALWAYS_INLINE auto ANGLE_textureOffset_impl(
2549     thread Texture &texture,
2550     thread metal::sampler const &sampler,
2551     metal::float3 const coord,
2552     metal::int3 const offset)
2553 {
2554     return texture.sample(sampler, coord, offset);
2555 }
2556 
2557 template <typename Texture>
2558 ANGLE_ALWAYS_INLINE auto ANGLE_textureOffset_impl(
2559     thread Texture &texture,
2560     thread metal::sampler const &sampler,
2561     metal::float3 const coord,
2562     metal::int3 const offset,
2563     float bias)
2564 {
2565     return texture.sample(sampler, coord, metal::bias(bias), offset);
2566 }
2567 
2568 template <typename T>
2569 ANGLE_ALWAYS_INLINE auto ANGLE_textureOffset_impl(
2570     thread metal::depth2d<T> &texture,
2571     thread metal::sampler const &sampler,
2572     metal::float3 const coord,
2573     metal::int2 const offset)
2574 {
2575     return texture.sample_compare(sampler, coord.xy, coord.z, offset);
2576 }
2577 
2578 template <typename T>
2579 ANGLE_ALWAYS_INLINE auto ANGLE_textureOffset_impl(
2580     thread metal::depth2d<T> &texture,
2581     thread metal::sampler const &sampler,
2582     metal::float3 const coord,
2583     metal::int2 const offset,
2584     float bias)
2585 {
2586     return texture.sample_compare(sampler, coord.xy, coord.z, metal::bias(bias), offset);
2587 }
2588 )",
2589                         textureEnv())
2590 
2591 PROGRAM_PRELUDE_DECLARE(textureProj,
2592                         R"(
2593 #define ANGLE_textureProj(env, ...) ANGLE_textureProj_impl(*env.texture, *env.sampler, __VA_ARGS__)
2594 
2595 template <typename Texture>
2596 ANGLE_ALWAYS_INLINE auto ANGLE_textureProj_impl(
2597     thread Texture &texture,
2598     thread metal::sampler const &sampler,
2599     metal::float3 const coord,
2600     float bias = 0)
2601 {
2602     return texture.sample(sampler, coord.xy/coord.z, metal::bias(bias));
2603 }
2604 
2605 template <typename Texture>
2606 ANGLE_ALWAYS_INLINE auto ANGLE_textureProj_impl(
2607     thread Texture &texture,
2608     thread metal::sampler const &sampler,
2609     metal::float4 const coord,
2610     float bias = 0)
2611 {
2612     return texture.sample(sampler, coord.xy/coord.w, metal::bias(bias));
2613 }
2614 
2615 template <typename T>
2616 ANGLE_ALWAYS_INLINE auto ANGLE_textureProj_impl(
2617     thread metal::texture3d<T> &texture,
2618     thread metal::sampler const &sampler,
2619     metal::float4 const coord,
2620     float bias = 0)
2621 {
2622     return texture.sample(sampler, coord.xyz/coord.w, metal::bias(bias));
2623 }
2624 )",
2625                         textureEnv())
2626 
2627 PROGRAM_PRELUDE_DECLARE(textureProjGrad,
2628                         R"(
2629 #define ANGLE_textureProjGrad(env, ...) ANGLE_textureProjGrad_impl(*env.texture, *env.sampler, __VA_ARGS__)
2630 )",
2631                         textureEnv())
2632 
2633 PROGRAM_PRELUDE_DECLARE(textureProjGrad_generic_float3_float2_float2,
2634                         R"(
2635 template <typename Texture>
2636 ANGLE_ALWAYS_INLINE auto ANGLE_textureProjGrad_impl(
2637     thread Texture &texture,
2638     thread metal::sampler const &sampler,
2639     metal::float3 const coord,
2640     metal::float2 const dPdx,
2641     metal::float2 const dPdy)
2642 {
2643     return texture.sample(sampler, coord.xy/coord.z, metal::gradient2d(dPdx, dPdy));
2644 }
2645 )",
2646                         textureProjGrad())
2647 
2648 PROGRAM_PRELUDE_DECLARE(textureProjGrad_generic_float4_float2_float2,
2649                         R"(
2650 template <typename Texture>
2651 ANGLE_ALWAYS_INLINE auto ANGLE_textureProjGrad_impl(
2652     thread Texture &texture,
2653     thread metal::sampler const &sampler,
2654     metal::float4 const coord,
2655     metal::float2 const dPdx,
2656     metal::float2 const dPdy)
2657 {
2658     return texture.sample(sampler, coord.xy/coord.w, metal::gradient2d(dPdx, dPdy));
2659 }
2660 )",
2661                         textureProjGrad())
2662 
2663 PROGRAM_PRELUDE_DECLARE(textureProjGrad_depth2d_float4_float2_float2,
2664                         R"(
2665 template <typename T>
2666 ANGLE_ALWAYS_INLINE auto ANGLE_textureProjGrad_impl(
2667     thread metal::depth2d<T> &texture,
2668     thread metal::sampler const &sampler,
2669     metal::float4 const coord,
2670     metal::float2 const dPdx,
2671     metal::float2 const dPdy)
2672 {
2673     if (ANGLEUseSampleCompareGradient)
2674     {
2675         return static_cast<T>(texture.sample_compare(sampler, coord.xy/coord.w, coord.z/coord.w, metal::gradient2d(dPdx, dPdy)));
2676     }
2677     else
2678     {
2679         return static_cast<T>(texture.sample(sampler, coord.xy/coord.w, metal::gradient2d(dPdx, dPdy)) > coord.z/coord.w);
2680     }
2681 }
2682 )",
2683                         functionConstants(),
2684                         textureProjGrad())
2685 
2686 PROGRAM_PRELUDE_DECLARE(textureProjGrad_texture3d_float4_float3_float3,
2687                         R"(
2688 template <typename T>
2689 ANGLE_ALWAYS_INLINE auto ANGLE_textureProjGrad_impl(
2690     thread metal::texture3d<T> &texture,
2691     thread metal::sampler const &sampler,
2692     metal::float4 const coord,
2693     metal::float3 const dPdx,
2694     metal::float3 const dPdy)
2695 {
2696     return texture.sample(sampler, coord.xyz/coord.w, metal::gradient3d(dPdx, dPdy));
2697 }
2698 )",
2699                         textureProjGrad())
2700 
2701 PROGRAM_PRELUDE_DECLARE(textureProjGradOffset,
2702                         R"(
2703 #define ANGLE_textureProjGradOffset(env, ...) ANGLE_textureProjGradOffset_impl(*env.texture, *env.sampler, __VA_ARGS__)
2704 )",
2705                         textureEnv())
2706 
2707 PROGRAM_PRELUDE_DECLARE(textureProjGradOffset_generic_float3_float2_float2_int2,
2708                         R"(
2709 template <typename Texture>
2710 ANGLE_ALWAYS_INLINE auto ANGLE_textureProjGradOffset_impl(
2711     thread Texture &texture,
2712     thread metal::sampler const &sampler,
2713     metal::float3 const coord,
2714     metal::float2 const dPdx,
2715     metal::float2 const dPdy,
2716     int2 const offset)
2717 {
2718     return texture.sample(sampler, coord.xy/coord.z, metal::gradient2d(dPdx, dPdy), offset);
2719 }
2720 )",
2721                         textureProjGradOffset())
2722 
2723 PROGRAM_PRELUDE_DECLARE(textureProjGradOffset_generic_float4_float2_float2_int2,
2724                         R"(
2725 template <typename Texture>
2726 ANGLE_ALWAYS_INLINE auto ANGLE_textureProjGradOffset_impl(
2727     thread Texture &texture,
2728     thread metal::sampler const &sampler,
2729     metal::float4 const coord,
2730     metal::float2 const dPdx,
2731     metal::float2 const dPdy,
2732     int2 const offset)
2733 {
2734     return texture.sample(sampler, coord.xy/coord.w, metal::gradient2d(dPdx, dPdy), offset);
2735 }
2736 )",
2737                         textureProjGradOffset())
2738 
2739 PROGRAM_PRELUDE_DECLARE(textureProjGradOffset_depth2d_float4_float2_float2_int2,
2740                         R"(
2741 template <typename T>
2742 ANGLE_ALWAYS_INLINE auto ANGLE_textureProjGradOffset_impl(
2743     thread metal::depth2d<T> &texture,
2744     thread metal::sampler const &sampler,
2745     metal::float4 const coord,
2746     metal::float2 const dPdx,
2747     metal::float2 const dPdy,
2748     int2 const offset)
2749 {
2750     if (ANGLEUseSampleCompareGradient)
2751     {
2752         return static_cast<T>(texture.sample_compare(sampler, coord.xy/coord.w, coord.z/coord.w, metal::gradient2d(dPdx, dPdy), offset));
2753     }
2754     else
2755     {
2756         return static_cast<T>(texture.sample(sampler, coord.xy/coord.w, metal::gradient2d(dPdx, dPdy), offset) > coord.z/coord.w);
2757     }
2758 }
2759 )",
2760                         functionConstants(),
2761                         textureProjGradOffset())
2762 
2763 PROGRAM_PRELUDE_DECLARE(textureProjGradOffset_texture3d_float4_float3_float3_int3,
2764                         R"(
2765 template <typename T>
2766 ANGLE_ALWAYS_INLINE auto ANGLE_textureProjGradOffset_impl(
2767     thread metal::texture3d<T> &texture,
2768     thread metal::sampler const &sampler,
2769     metal::float4 const coord,
2770     metal::float3 const dPdx,
2771     metal::float3 const dPdy,
2772     int3 const offset)
2773 {
2774     return texture.sample(sampler, coord.xyz/coord.w, metal::gradient3d(dPdx, dPdy), offset);
2775 }
2776 )",
2777                         textureProjGradOffset())
2778 
2779 PROGRAM_PRELUDE_DECLARE(textureProjLod,
2780                         R"(
2781 #define ANGLE_textureProjLod(env, ...) ANGLE_textureProjLod_impl(*env.texture, *env.sampler, __VA_ARGS__)
2782 )",
2783                         textureEnv())
2784 
2785 PROGRAM_PRELUDE_DECLARE(textureProjLod_generic_float3,
2786                         R"(
2787 template <typename Texture>
2788 ANGLE_ALWAYS_INLINE auto ANGLE_textureProjLod_impl(
2789     thread Texture &texture,
2790     thread metal::sampler const &sampler,
2791     metal::float3 const coord,
2792     float level)
2793 {
2794     return texture.sample(sampler, coord.xy/coord.z, metal::level(level));
2795 }
2796 )",
2797                         textureProjLod())
2798 
2799 PROGRAM_PRELUDE_DECLARE(textureProjLod_generic_float4,
2800                         R"(
2801 template <typename Texture>
2802 ANGLE_ALWAYS_INLINE auto ANGLE_textureProjLod_impl(
2803     thread Texture &texture,
2804     thread metal::sampler const &sampler,
2805     metal::float4 const coord,
2806     float level)
2807 {
2808     return texture.sample(sampler, coord.xy/coord.w, metal::level(level));
2809 }
2810 )",
2811                         textureProjLod())
2812 
2813 PROGRAM_PRELUDE_DECLARE(textureProjLod_depth2d_float4,
2814                         R"(
2815 template <typename T>
2816 ANGLE_ALWAYS_INLINE auto ANGLE_textureProjLod_impl(
2817     thread metal::depth2d<T> &texture,
2818     thread metal::sampler const &sampler,
2819     metal::float4 const coord,
2820     float level)
2821 {
2822     if (ANGLEUseSampleCompareLod)
2823     {
2824         return static_cast<T>(texture.sample_compare(sampler, coord.xy/coord.w, coord.z/coord.w, metal::level(level)));
2825     }
2826     else
2827     {
2828         return static_cast<T>(texture.sample(sampler, coord.xy/coord.w, metal::level(level)) > coord.z/coord.w);
2829     }
2830 }
2831 )",
2832                         functionConstants(),
2833                         textureProjLod())
2834 
2835 PROGRAM_PRELUDE_DECLARE(textureProjLod_texture3d_float4,
2836                         R"(
2837 template <typename T>
2838 ANGLE_ALWAYS_INLINE auto ANGLE_textureProjLod_impl(
2839     thread metal::texture3d<T> &texture,
2840     thread metal::sampler const &sampler,
2841     metal::float4 const coord,
2842     float level)
2843 {
2844     return texture.sample(sampler, coord.xyz/coord.w, metal::level(level));
2845 }
2846 )",
2847                         textureProjLod())
2848 
2849 PROGRAM_PRELUDE_DECLARE(textureProjLodOffset,
2850                         R"(
2851 #define ANGLE_textureProjLodOffset(env, ...) ANGLE_textureProjLodOffset_impl(*env.texture, *env.sampler, __VA_ARGS__)
2852 
2853 template <typename Texture>
2854 ANGLE_ALWAYS_INLINE auto ANGLE_textureProjLodOffset_impl(
2855     thread Texture &texture,
2856     thread metal::sampler const &sampler,
2857     metal::float3 const coord,
2858     float level,
2859     int2 const offset)
2860 {
2861     return texture.sample(sampler, coord.xy/coord.z, metal::level(level), offset);
2862 }
2863 
2864 template <typename Texture>
2865 ANGLE_ALWAYS_INLINE auto ANGLE_textureProjLodOffset_impl(
2866     thread Texture &texture,
2867     thread metal::sampler const &sampler,
2868     metal::float4 const coord,
2869     float level,
2870     int2 const offset)
2871 {
2872     return texture.sample(sampler, coord.xy/coord.w, metal::level(level), offset);
2873 }
2874 
2875 template <typename T>
2876 ANGLE_ALWAYS_INLINE auto ANGLE_textureProjLodOffset_impl(
2877     thread metal::depth2d<T> &texture,
2878     thread metal::sampler const &sampler,
2879     metal::float4 const coord,
2880     float level,
2881     int2 const offset)
2882 {
2883     if (ANGLEUseSampleCompareLod)
2884     {
2885         return static_cast<T>(texture.sample_compare(sampler, coord.xy/coord.w, coord.z/coord.w, metal::level(level), offset));
2886     }
2887     else
2888     {
2889         return static_cast<T>(texture.sample(sampler, coord.xy/coord.w, metal::level(level), offset) > coord.z/coord.w);
2890     }
2891 }
2892 
2893 template <typename T>
2894 ANGLE_ALWAYS_INLINE auto ANGLE_textureProjLodOffset_impl(
2895     thread metal::texture3d<T> &texture,
2896     thread metal::sampler const &sampler,
2897     metal::float4 const coord,
2898     float level,
2899     int3 const offset)
2900 {
2901     return texture.sample(sampler, coord.xyz/coord.w, metal::level(level), offset);
2902 }
2903 )",
2904                         functionConstants(),
2905                         textureEnv())
2906 
2907 PROGRAM_PRELUDE_DECLARE(textureProjOffset,
2908                         R"(
2909 #define ANGLE_textureProjOffset(env, ...) ANGLE_textureProjOffset_impl(*env.texture, *env.sampler, __VA_ARGS__)
2910 
2911 template <typename Texture>
2912 ANGLE_ALWAYS_INLINE auto ANGLE_textureProjOffset_impl(
2913     thread Texture &texture,
2914     thread metal::sampler const &sampler,
2915     metal::float3 const coord,
2916     int2 const offset,
2917     float bias = 0)
2918 {
2919     return texture.sample(sampler, coord.xy/coord.z, metal::bias(bias), offset);
2920 }
2921 
2922 template <typename Texture>
2923 ANGLE_ALWAYS_INLINE auto ANGLE_textureProjOffset_impl(
2924     thread Texture &texture,
2925     thread metal::sampler const &sampler,
2926     metal::float4 const coord,
2927     int2 const offset,
2928     float bias = 0)
2929 {
2930     return texture.sample(sampler, coord.xy/coord.w, metal::bias(bias), offset);
2931 }
2932 
2933 template <typename T>
2934 ANGLE_ALWAYS_INLINE auto ANGLE_textureProjOffset_impl(
2935     thread metal::texture3d<T> &texture,
2936     thread metal::sampler const &sampler,
2937     metal::float4 const coord,
2938     int3 const offset,
2939     float bias = 0)
2940 {
2941     return texture.sample(sampler, coord.xyz/coord.w, metal::bias(bias), offset);
2942 }
2943 )",
2944                         textureEnv())
2945 
2946 PROGRAM_PRELUDE_DECLARE(textureSize,
2947                         R"(
2948 #define ANGLE_textureSize(env, ...) ANGLE_textureSize_impl(*env.texture, __VA_ARGS__)
2949 
2950 template <typename Texture>
2951 ANGLE_ALWAYS_INLINE auto ANGLE_textureSize_impl(
2952     thread Texture &texture,
2953     int level)
2954 {
2955     return int2(texture.get_width(uint(level)), texture.get_height(uint(level)));
2956 }
2957 
2958 template <typename T>
2959 ANGLE_ALWAYS_INLINE auto ANGLE_textureSize_impl(
2960     thread metal::texture3d<T> &texture,
2961     int level)
2962 {
2963     return int3(texture.get_width(uint(level)), texture.get_height(uint(level)), texture.get_depth(uint(level)));
2964 }
2965 
2966 template <typename T>
2967 ANGLE_ALWAYS_INLINE auto ANGLE_textureSize_impl(
2968     thread metal::depth2d_array<T> &texture,
2969     int level)
2970 {
2971     return int3(texture.get_width(uint(level)), texture.get_height(uint(level)), texture.get_array_size());
2972 }
2973 
2974 template <typename T>
2975 ANGLE_ALWAYS_INLINE auto ANGLE_textureSize_impl(
2976     thread metal::texture2d_array<T> &texture,
2977     int level)
2978 {
2979     return int3(texture.get_width(uint(level)), texture.get_height(uint(level)), texture.get_array_size());
2980 }
2981 )",
2982                         textureEnv())
2983 
2984 ////////////////////////////////////////////////////////////////////////////////
2985 
2986 // Returned Name is valid for as long as `buffer` is still alive.
2987 // Returns false if no template args exist.
2988 // Returns false if buffer is not large enough.
2989 //
2990 // Example:
2991 //  "foo<1,2>" --> "foo<>"
2992 static std::pair<Name, bool> MaskTemplateArgs(const Name &name, size_t bufferSize, char *buffer)
2993 {
2994     const char *begin = name.rawName().data();
2995     const char *end   = strchr(begin, '<');
2996     if (!end)
2997     {
2998         return {{}, false};
2999     }
3000     size_t n = end - begin;
3001     if (n + 3 > bufferSize)
3002     {
3003         return {{}, false};
3004     }
3005     for (size_t i = 0; i < n; ++i)
3006     {
3007         buffer[i] = begin[i];
3008     }
3009     buffer[n + 0] = '<';
3010     buffer[n + 1] = '>';
3011     buffer[n + 2] = '\0';
3012     return {Name(buffer, name.symbolType()), true};
3013 }
3014 
BuildFuncToEmitter()3015 ProgramPrelude::FuncToEmitter ProgramPrelude::BuildFuncToEmitter()
3016 {
3017 #define EMIT_METHOD(method) \
3018     [](ProgramPrelude &pp, const TFunction &) -> void { return pp.method(); }
3019     FuncToEmitter map;
3020 
3021     auto put = [&](Name name, FuncEmitter emitter) {
3022         FuncEmitter &dest = map[name];
3023         ASSERT(!dest);
3024         dest = emitter;
3025     };
3026 
3027     auto putAngle = [&](const char *nameStr, FuncEmitter emitter) {
3028         Name name(nameStr, SymbolType::AngleInternal);
3029         put(name, emitter);
3030     };
3031 
3032     auto putBuiltIn = [&](const char *nameStr, FuncEmitter emitter) {
3033         Name name(nameStr, SymbolType::BuiltIn);
3034         put(name, emitter);
3035     };
3036 
3037     putAngle("addressof", EMIT_METHOD(addressof));
3038     putAngle("cast<>", EMIT_METHOD(castMatrix));
3039     putAngle("elem_ref", EMIT_METHOD(vectorElemRef));
3040     putAngle("flatten", EMIT_METHOD(flattenArray));
3041     putAngle("inout", EMIT_METHOD(inout));
3042     putAngle("out", EMIT_METHOD(out));
3043     putAngle("swizzle_ref", EMIT_METHOD(swizzleRef));
3044 
3045     putBuiltIn("texelFetch", EMIT_METHOD(texelFetch));
3046     putBuiltIn("texelFetchOffset", EMIT_METHOD(texelFetchOffset));
3047     putBuiltIn("texture", [](ProgramPrelude &pp, const TFunction &func) {
3048         const ImmutableString textureName =
3049             GetTextureTypeName(func.getParam(0)->getType().getBasicType()).rawName();
3050         const TType &coord          = func.getParam(1)->getType();
3051         const TBasicType coordBasic = coord.getBasicType();
3052         const int coordN            = coord.getNominalSize();
3053         const bool bias             = func.getParamCount() >= 3;
3054         if (textureName.beginsWith("metal::depth2d<"))
3055         {
3056             if (coordBasic == TBasicType::EbtFloat && coordN == 3)
3057             {
3058                 if (bias)
3059                 {
3060                     return pp.texture_depth2d_float3_float();
3061                 }
3062                 return pp.texture_depth2d_float3();
3063             }
3064         }
3065         if (textureName.beginsWith("metal::depthcube<"))
3066         {
3067             if (coordBasic == TBasicType::EbtFloat && coordN == 4)
3068             {
3069                 if (bias)
3070                 {
3071                     return pp.texture_depthcube_float4_float();
3072                 }
3073                 return pp.texture_depthcube_float4();
3074             }
3075         }
3076         if (textureName.beginsWith("metal::depth2d_array<"))
3077         {
3078             if (coordBasic == TBasicType::EbtFloat && coordN == 4)
3079             {
3080                 if (bias)
3081                 {
3082                     return pp.texture_depth2darray_float4_float();
3083                 }
3084                 return pp.texture_depth2darray_float4();
3085             }
3086         }
3087         if (textureName.beginsWith("metal::texture2d_array<"))
3088         {
3089             if (coordBasic == TBasicType::EbtFloat && coordN == 3)
3090             {
3091                 if (bias)
3092                 {
3093                     return pp.texture_texture2darray_float3_float();
3094                 }
3095                 return pp.texture_texture2darray_float3();
3096             }
3097             if (coordBasic == TBasicType::EbtFloat && coordN == 4)
3098             {
3099                 if (bias)
3100                 {
3101                     return pp.texture_texture2darray_float4_float();
3102                 }
3103                 return pp.texture_texture2darray_float4();
3104             }
3105         }
3106         if (coordBasic == TBasicType::EbtFloat && coordN == 2)
3107         {
3108             if (bias)
3109             {
3110                 return pp.texture_generic_float2_float();
3111             }
3112             return pp.texture_generic_float2();
3113         }
3114         if (coordBasic == TBasicType::EbtFloat && coordN == 3)
3115         {
3116             if (bias)
3117             {
3118                 return pp.texture_generic_float3_float();
3119             }
3120             return pp.texture_generic_float3();
3121         }
3122         UNIMPLEMENTED();
3123     });
3124     putBuiltIn("texture1DLod", EMIT_METHOD(texture1DLod));
3125     putBuiltIn("texture1DProj", EMIT_METHOD(texture1DProj));
3126     putBuiltIn("texture1DProjLod", EMIT_METHOD(texture1DProjLod));
3127     putBuiltIn("texture2D", EMIT_METHOD(texture2D));
3128     putBuiltIn("texture2DLod", EMIT_METHOD(texture2DLod));
3129     putBuiltIn("texture2DProj", EMIT_METHOD(texture2DProj));
3130     putBuiltIn("texture2DRect", EMIT_METHOD(texture2DRect));
3131     putBuiltIn("texture2DRectProj", EMIT_METHOD(texture2DRectProj));
3132     putBuiltIn("texture3DLod", EMIT_METHOD(texture3DLod));
3133     putBuiltIn("texture3DProj", EMIT_METHOD(texture3DProj));
3134     putBuiltIn("texture3DProjLod", EMIT_METHOD(texture3DProjLod));
3135     putBuiltIn("textureCube", EMIT_METHOD(textureCube));
3136     putBuiltIn("textureCubeLod", EMIT_METHOD(textureCubeLod));
3137     putBuiltIn("textureCubeProj", EMIT_METHOD(textureCubeProj));
3138     putBuiltIn("textureCubeProjLod", EMIT_METHOD(textureCubeProjLod));
3139     putBuiltIn("texture2DProjLod", EMIT_METHOD(texture2DProjLod));
3140     putBuiltIn("textureGrad", [](ProgramPrelude &pp, const TFunction &func) {
3141         const ImmutableString textureName =
3142             GetTextureTypeName(func.getParam(0)->getType().getBasicType()).rawName();
3143         const TType &coord          = func.getParam(1)->getType();
3144         const TBasicType coordBasic = coord.getBasicType();
3145         const int coordN            = coord.getNominalSize();
3146         const TType &dPdx           = func.getParam(2)->getType();
3147         const int dPdxN             = dPdx.getNominalSize();
3148         if (textureName.beginsWith("metal::depth2d<"))
3149         {
3150             if (coordBasic == TBasicType::EbtFloat && coordN == 3 && dPdxN == 2)
3151             {
3152                 return pp.textureGrad_depth2d_float3_float2_float2();
3153             }
3154         }
3155         if (textureName.beginsWith("metal::depth2d_array<"))
3156         {
3157             if (coordBasic == TBasicType::EbtFloat && coordN == 4 && dPdxN == 2)
3158             {
3159                 return pp.textureGrad_depth2darray_float4_float2_float2();
3160             }
3161         }
3162         if (textureName.beginsWith("metal::depthcube<"))
3163         {
3164             if (coordBasic == TBasicType::EbtFloat && coordN == 4 && dPdxN == 3)
3165             {
3166                 return pp.textureGrad_depthcube_float4_float3_float3();
3167             }
3168         }
3169         if (textureName.beginsWith("metal::texturecube<"))
3170         {
3171             if (coordBasic == TBasicType::EbtFloat && coordN == 3 && dPdxN == 3)
3172             {
3173                 return pp.textureGrad_texturecube_float3_float3_float3();
3174             }
3175         }
3176         if (coordBasic == TBasicType::EbtFloat && coordN == 3 && dPdxN == 2)
3177         {
3178             return pp.textureGrad_generic_float3_float2_float2();
3179         }
3180         if (coordBasic == TBasicType::EbtFloat && coordN == 4 && dPdxN == 2)
3181         {
3182             return pp.textureGrad_generic_float4_float2_float2();
3183         }
3184         if (coordBasic == TBasicType::EbtFloat && coordN == dPdxN)
3185         {
3186             return pp.textureGrad_generic_floatN_floatN_floatN();
3187         }
3188         UNIMPLEMENTED();
3189     });
3190     putBuiltIn("textureGradOffset", [](ProgramPrelude &pp, const TFunction &func) {
3191         const ImmutableString textureName =
3192             GetTextureTypeName(func.getParam(0)->getType().getBasicType()).rawName();
3193         const TType &coord          = func.getParam(1)->getType();
3194         const TBasicType coordBasic = coord.getBasicType();
3195         const int coordN            = coord.getNominalSize();
3196         const TType &dPdx           = func.getParam(2)->getType();
3197         const int dPdxN             = dPdx.getNominalSize();
3198         if (textureName.beginsWith("metal::depth2d<"))
3199         {
3200             if (coordBasic == TBasicType::EbtFloat && coordN == 3 && dPdxN == 2)
3201             {
3202                 return pp.textureGradOffset_depth2d_float3_float2_float2_int2();
3203             }
3204         }
3205         if (textureName.beginsWith("metal::depth2d_array<"))
3206         {
3207             if (coordBasic == TBasicType::EbtFloat && coordN == 4 && dPdxN == 2)
3208             {
3209                 return pp.textureGradOffset_depth2darray_float4_float2_float2_int2();
3210             }
3211         }
3212         if (textureName.beginsWith("metal::depthcube<"))
3213         {
3214             if (coordBasic == TBasicType::EbtFloat && coordN == 4 && dPdxN == 3)
3215             {
3216                 return pp.textureGradOffset_depthcube_float4_float3_float3_int3();
3217             }
3218         }
3219         if (textureName.beginsWith("metal::texturecube<"))
3220         {
3221             if (coordBasic == TBasicType::EbtFloat && coordN == 3 && dPdxN == 3)
3222             {
3223                 return pp.textureGradOffset_texturecube_float3_float3_float3_int3();
3224             }
3225         }
3226         if (coordBasic == TBasicType::EbtFloat && coordN == 3 && dPdxN == 2)
3227         {
3228             return pp.textureGradOffset_generic_float3_float2_float2_int2();
3229         }
3230         if (coordBasic == TBasicType::EbtFloat && coordN == 4 && dPdxN == 2)
3231         {
3232             return pp.textureGradOffset_generic_float4_float2_float2_int2();
3233         }
3234         if (coordBasic == TBasicType::EbtFloat && coordN == dPdxN)
3235         {
3236             return pp.textureGradOffset_generic_floatN_floatN_floatN_intN();
3237         }
3238         UNIMPLEMENTED();
3239     });
3240     putBuiltIn("textureLod", [](ProgramPrelude &pp, const TFunction &func) {
3241         const ImmutableString textureName =
3242             GetTextureTypeName(func.getParam(0)->getType().getBasicType()).rawName();
3243         const TType &coord          = func.getParam(1)->getType();
3244         const TBasicType coordBasic = coord.getBasicType();
3245         const int coordN            = coord.getNominalSize();
3246         if (textureName.beginsWith("metal::depth2d<"))
3247         {
3248             if (coordBasic == TBasicType::EbtFloat && coordN == 3)
3249             {
3250                 return pp.textureLod_depth2d_float3();
3251             }
3252         }
3253         if (textureName.beginsWith("metal::texture2d_array<"))
3254         {
3255             if (coordBasic == TBasicType::EbtFloat && coordN == 3)
3256             {
3257                 return pp.textureLod_texture2darray_float3();
3258             }
3259             if (coordBasic == TBasicType::EbtFloat && coordN == 4)
3260             {
3261                 return pp.textureLod_texture2darray_float4();
3262             }
3263         }
3264         if (coordBasic == TBasicType::EbtFloat && coordN == 2)
3265         {
3266             return pp.textureLod_generic_float2();
3267         }
3268         if (coordBasic == TBasicType::EbtFloat && coordN == 3)
3269         {
3270             return pp.textureLod_generic_float3();
3271         }
3272         UNIMPLEMENTED();
3273     });
3274     putBuiltIn("textureLodOffset", EMIT_METHOD(textureLodOffset));
3275     putBuiltIn("textureOffset", EMIT_METHOD(textureOffset));
3276     putBuiltIn("textureProj", EMIT_METHOD(textureProj));
3277     putBuiltIn("textureProjGrad", [](ProgramPrelude &pp, const TFunction &func) {
3278         const ImmutableString textureName =
3279             GetTextureTypeName(func.getParam(0)->getType().getBasicType()).rawName();
3280         const TType &coord          = func.getParam(1)->getType();
3281         const TBasicType coordBasic = coord.getBasicType();
3282         const int coordN            = coord.getNominalSize();
3283         const TType &dPdx           = func.getParam(2)->getType();
3284         const int dPdxN             = dPdx.getNominalSize();
3285         if (textureName.beginsWith("metal::depth2d<"))
3286         {
3287             if (coordBasic == TBasicType::EbtFloat && coordN == 4 && dPdxN == 2)
3288             {
3289                 return pp.textureProjGrad_depth2d_float4_float2_float2();
3290             }
3291         }
3292         if (textureName.beginsWith("metal::texture3d<"))
3293         {
3294             if (coordBasic == TBasicType::EbtFloat && coordN == 4 && dPdxN == 3)
3295             {
3296                 return pp.textureProjGrad_texture3d_float4_float3_float3();
3297             }
3298         }
3299         if (coordBasic == TBasicType::EbtFloat && coordN == 3 && dPdxN == 2)
3300         {
3301             return pp.textureProjGrad_generic_float3_float2_float2();
3302         }
3303         if (coordBasic == TBasicType::EbtFloat && coordN == 4 && dPdxN == 2)
3304         {
3305             return pp.textureProjGrad_generic_float4_float2_float2();
3306         }
3307         UNIMPLEMENTED();
3308     });
3309     putBuiltIn("textureProjGradOffset", [](ProgramPrelude &pp, const TFunction &func) {
3310         const ImmutableString textureName =
3311             GetTextureTypeName(func.getParam(0)->getType().getBasicType()).rawName();
3312         const TType &coord          = func.getParam(1)->getType();
3313         const TBasicType coordBasic = coord.getBasicType();
3314         const int coordN            = coord.getNominalSize();
3315         const TType &dPdx           = func.getParam(2)->getType();
3316         const int dPdxN             = dPdx.getNominalSize();
3317         if (textureName.beginsWith("metal::depth2d<"))
3318         {
3319             if (coordBasic == TBasicType::EbtFloat && coordN == 4 && dPdxN == 2)
3320             {
3321                 return pp.textureProjGradOffset_depth2d_float4_float2_float2_int2();
3322             }
3323         }
3324         if (textureName.beginsWith("metal::texture3d<"))
3325         {
3326             if (coordBasic == TBasicType::EbtFloat && coordN == 4 && dPdxN == 3)
3327             {
3328                 return pp.textureProjGradOffset_texture3d_float4_float3_float3_int3();
3329             }
3330         }
3331         if (coordBasic == TBasicType::EbtFloat && coordN == 3 && dPdxN == 2)
3332         {
3333             return pp.textureProjGradOffset_generic_float3_float2_float2_int2();
3334         }
3335         if (coordBasic == TBasicType::EbtFloat && coordN == 4 && dPdxN == 2)
3336         {
3337             return pp.textureProjGradOffset_generic_float4_float2_float2_int2();
3338         }
3339         UNIMPLEMENTED();
3340     });
3341     putBuiltIn("textureProjLod", [](ProgramPrelude &pp, const TFunction &func) {
3342         const ImmutableString textureName =
3343             GetTextureTypeName(func.getParam(0)->getType().getBasicType()).rawName();
3344         const TType &coord          = func.getParam(1)->getType();
3345         const TBasicType coordBasic = coord.getBasicType();
3346         const int coordN            = coord.getNominalSize();
3347         if (textureName.beginsWith("metal::depth2d<"))
3348         {
3349             if (coordBasic == TBasicType::EbtFloat && coordN == 4)
3350             {
3351                 return pp.textureProjLod_depth2d_float4();
3352             }
3353         }
3354         if (textureName.beginsWith("metal::texture3d<"))
3355         {
3356             if (coordBasic == TBasicType::EbtFloat && coordN == 4)
3357             {
3358                 return pp.textureProjLod_texture3d_float4();
3359             }
3360         }
3361         if (coordBasic == TBasicType::EbtFloat && coordN == 3)
3362         {
3363             return pp.textureProjLod_generic_float3();
3364         }
3365         if (coordBasic == TBasicType::EbtFloat && coordN == 4)
3366         {
3367             return pp.textureProjLod_generic_float4();
3368         }
3369         UNIMPLEMENTED();
3370     });
3371     putBuiltIn("textureProjLodOffset", EMIT_METHOD(textureProjLodOffset));
3372     putBuiltIn("textureProjOffset", EMIT_METHOD(textureProjOffset));
3373     putBuiltIn("textureSize", EMIT_METHOD(textureSize));
3374 
3375     return map;
3376 
3377 #undef EMIT_METHOD
3378 }
3379 
visitOperator(TOperator op,const TFunction * func,const TType * argType0,const TType * argType1)3380 void ProgramPrelude::visitOperator(TOperator op,
3381                                    const TFunction *func,
3382                                    const TType *argType0,
3383                                    const TType *argType1)
3384 {
3385     switch (op)
3386     {
3387         case TOperator::EOpRadians:
3388             radians();
3389             break;
3390         case TOperator::EOpDegrees:
3391             degrees();
3392             break;
3393         case TOperator::EOpAtan:
3394             atan();
3395             break;
3396         case TOperator::EOpMod:
3397             mod();
3398             break;
3399         case TOperator::EOpRefract:
3400             refract();
3401             break;
3402         case TOperator::EOpDistance:
3403             distance();
3404             break;
3405         case TOperator::EOpLength:
3406             length();
3407             break;
3408         case TOperator::EOpDot:
3409             dot();
3410             break;
3411         case TOperator::EOpNormalize:
3412             normalize();
3413             break;
3414         case TOperator::EOpFaceforward:
3415             faceforward();
3416             break;
3417         case TOperator::EOpReflect:
3418             reflect();
3419             break;
3420 
3421         case TOperator::EOpSin:
3422         case TOperator::EOpCos:
3423         case TOperator::EOpTan:
3424         case TOperator::EOpAsin:
3425         case TOperator::EOpAcos:
3426         case TOperator::EOpSinh:
3427         case TOperator::EOpCosh:
3428         case TOperator::EOpTanh:
3429         case TOperator::EOpAsinh:
3430         case TOperator::EOpAcosh:
3431         case TOperator::EOpAtanh:
3432         case TOperator::EOpAbs:
3433         case TOperator::EOpFma:
3434         case TOperator::EOpPow:
3435         case TOperator::EOpExp:
3436         case TOperator::EOpExp2:
3437         case TOperator::EOpLog:
3438         case TOperator::EOpLog2:
3439         case TOperator::EOpSqrt:
3440         case TOperator::EOpFloor:
3441         case TOperator::EOpTrunc:
3442         case TOperator::EOpCeil:
3443         case TOperator::EOpFract:
3444         case TOperator::EOpRound:
3445         case TOperator::EOpRoundEven:
3446         case TOperator::EOpModf:
3447         case TOperator::EOpLdexp:
3448         case TOperator::EOpFrexp:
3449         case TOperator::EOpInversesqrt:
3450             include_metal_math();
3451             break;
3452 
3453         case TOperator::EOpEqual:
3454             if (argType0->isVector() && argType1->isVector())
3455             {
3456                 equalVector();
3457             }
3458             // Even if Arg0 is a vector or matrix, it could also be an array.
3459             if (argType0->isArray() && argType1->isArray())
3460             {
3461                 equalArray();
3462             }
3463             if (argType0->getStruct() && argType1->getStruct() && argType0->isArray() &&
3464                 argType1->isArray())
3465             {
3466                 equalStructArray();
3467             }
3468             if (argType0->isMatrix() && argType1->isMatrix())
3469             {
3470                 equalMatrix();
3471             }
3472             break;
3473 
3474         case TOperator::EOpNotEqual:
3475             if (argType0->isVector() && argType1->isVector())
3476             {
3477                 notEqualVector();
3478             }
3479             else if (argType0->getStruct() && argType1->getStruct())
3480             {
3481                 notEqualStruct();
3482             }
3483             // Same as above.
3484             if (argType0->isArray() && argType1->isArray())
3485             {
3486                 notEqualArray();
3487             }
3488             if (argType0->getStruct() && argType1->getStruct() && argType0->isArray() &&
3489                 argType1->isArray())
3490             {
3491                 notEqualStructArray();
3492             }
3493             if (argType0->isMatrix() && argType1->isMatrix())
3494             {
3495                 notEqualMatrix();
3496             }
3497             break;
3498 
3499         case TOperator::EOpCross:
3500             include_metal_geometric();
3501             break;
3502 
3503         case TOperator::EOpSign:
3504             sign();
3505             break;
3506 
3507         case TOperator::EOpClamp:
3508         case TOperator::EOpMin:
3509         case TOperator::EOpMax:
3510         case TOperator::EOpMix:
3511         case TOperator::EOpStep:
3512         case TOperator::EOpSmoothstep:
3513             include_metal_common();
3514             break;
3515 
3516         case TOperator::EOpAll:
3517         case TOperator::EOpAny:
3518         case TOperator::EOpIsnan:
3519         case TOperator::EOpIsinf:
3520             include_metal_relational();
3521             break;
3522 
3523         case TOperator::EOpDFdx:
3524         case TOperator::EOpDFdy:
3525         case TOperator::EOpFwidth:
3526             include_metal_graphics();
3527             break;
3528 
3529         case TOperator::EOpTranspose:
3530         case TOperator::EOpDeterminant:
3531             include_metal_matrix();
3532             break;
3533 
3534         case TOperator::EOpAdd:
3535             if (argType0->isMatrix() && argType1->isScalar())
3536             {
3537                 addMatrixScalar();
3538             }
3539             break;
3540 
3541         case TOperator::EOpAddAssign:
3542             if (argType0->isMatrix() && argType1->isScalar())
3543             {
3544                 addMatrixScalarAssign();
3545             }
3546             break;
3547 
3548         case TOperator::EOpSub:
3549             if (argType0->isMatrix() && argType1->isScalar())
3550             {
3551                 subMatrixScalar();
3552             }
3553             break;
3554 
3555         case TOperator::EOpSubAssign:
3556             if (argType0->isMatrix() && argType1->isScalar())
3557             {
3558                 subMatrixScalarAssign();
3559             }
3560             break;
3561 
3562         case TOperator::EOpDiv:
3563             if (argType0->isMatrix())
3564             {
3565                 if (argType1->isMatrix())
3566                 {
3567                     componentWiseDivide();
3568                 }
3569                 else if (argType1->isScalar())
3570                 {
3571                     divMatrixScalar();
3572                 }
3573             }
3574             break;
3575 
3576         case TOperator::EOpDivAssign:
3577             if (argType0->isMatrix() && argType1->isMatrix())
3578             {
3579                 componentWiseDivideAssign();
3580             }
3581             break;
3582 
3583         case TOperator::EOpMatrixCompMult:
3584             if (argType0->isMatrix() && argType1->isMatrix())
3585             {
3586                 componentWiseMultiply();
3587             }
3588             break;
3589 
3590         case TOperator::EOpOuterProduct:
3591             outerProduct();
3592             break;
3593 
3594         case TOperator::EOpInverse:
3595             switch (argType0->getCols())
3596             {
3597                 case 2:
3598                     inverse2();
3599                     break;
3600                 case 3:
3601                     inverse3();
3602                     break;
3603                 case 4:
3604                     inverse4();
3605                     break;
3606                 default:
3607                     UNREACHABLE();
3608             }
3609             break;
3610 
3611         case TOperator::EOpMatrixTimesMatrixAssign:
3612             matmulAssign();
3613             break;
3614 
3615         case TOperator::EOpPreIncrement:
3616             if (argType0->isMatrix())
3617             {
3618                 preIncrementMatrix();
3619             }
3620             break;
3621 
3622         case TOperator::EOpPostIncrement:
3623             if (argType0->isMatrix())
3624             {
3625                 postIncrementMatrix();
3626             }
3627             break;
3628 
3629         case TOperator::EOpPreDecrement:
3630             if (argType0->isMatrix())
3631             {
3632                 preDecrementMatrix();
3633             }
3634             break;
3635 
3636         case TOperator::EOpPostDecrement:
3637             if (argType0->isMatrix())
3638             {
3639                 postDecrementMatrix();
3640             }
3641             break;
3642 
3643             break;
3644 
3645         case TOperator::EOpNegative:
3646             if (argType0->isMatrix())
3647             {
3648                 negateMatrix();
3649             }
3650             break;
3651 
3652         case TOperator::EOpComma:
3653         case TOperator::EOpAssign:
3654         case TOperator::EOpInitialize:
3655         case TOperator::EOpMulAssign:
3656         case TOperator::EOpIModAssign:
3657         case TOperator::EOpBitShiftLeftAssign:
3658         case TOperator::EOpBitShiftRightAssign:
3659         case TOperator::EOpBitwiseAndAssign:
3660         case TOperator::EOpBitwiseXorAssign:
3661         case TOperator::EOpBitwiseOrAssign:
3662         case TOperator::EOpMul:
3663         case TOperator::EOpIMod:
3664         case TOperator::EOpBitShiftLeft:
3665         case TOperator::EOpBitShiftRight:
3666         case TOperator::EOpBitwiseAnd:
3667         case TOperator::EOpBitwiseXor:
3668         case TOperator::EOpBitwiseOr:
3669         case TOperator::EOpLessThan:
3670         case TOperator::EOpGreaterThan:
3671         case TOperator::EOpLessThanEqual:
3672         case TOperator::EOpGreaterThanEqual:
3673         case TOperator::EOpLessThanComponentWise:
3674         case TOperator::EOpLessThanEqualComponentWise:
3675         case TOperator::EOpGreaterThanEqualComponentWise:
3676         case TOperator::EOpGreaterThanComponentWise:
3677         case TOperator::EOpLogicalOr:
3678         case TOperator::EOpLogicalXor:
3679         case TOperator::EOpLogicalAnd:
3680         case TOperator::EOpPositive:
3681         case TOperator::EOpLogicalNot:
3682         case TOperator::EOpNotComponentWise:
3683         case TOperator::EOpBitwiseNot:
3684         case TOperator::EOpVectorTimesScalarAssign:
3685         case TOperator::EOpVectorTimesMatrixAssign:
3686         case TOperator::EOpMatrixTimesScalarAssign:
3687         case TOperator::EOpVectorTimesScalar:
3688         case TOperator::EOpVectorTimesMatrix:
3689         case TOperator::EOpMatrixTimesVector:
3690         case TOperator::EOpMatrixTimesScalar:
3691         case TOperator::EOpMatrixTimesMatrix:
3692         case TOperator::EOpReturn:
3693         case TOperator::EOpBreak:
3694         case TOperator::EOpContinue:
3695         case TOperator::EOpEqualComponentWise:
3696         case TOperator::EOpNotEqualComponentWise:
3697         case TOperator::EOpIndexDirect:
3698         case TOperator::EOpIndexIndirect:
3699         case TOperator::EOpIndexDirectStruct:
3700         case TOperator::EOpIndexDirectInterfaceBlock:
3701         case TOperator::EOpFloatBitsToInt:
3702         case TOperator::EOpIntBitsToFloat:
3703         case TOperator::EOpUintBitsToFloat:
3704         case TOperator::EOpFloatBitsToUint:
3705         case TOperator::EOpNull:
3706             // do nothing
3707             break;
3708 
3709         case TOperator::EOpKill:
3710             include_metal_graphics();
3711             break;
3712 
3713         case TOperator::EOpPackUnorm2x16:
3714         case TOperator::EOpPackSnorm2x16:
3715         case TOperator::EOpPackUnorm4x8:
3716         case TOperator::EOpPackSnorm4x8:
3717         case TOperator::EOpUnpackSnorm2x16:
3718         case TOperator::EOpUnpackUnorm2x16:
3719         case TOperator::EOpUnpackUnorm4x8:
3720         case TOperator::EOpUnpackSnorm4x8:
3721             include_metal_pack();
3722             break;
3723 
3724         case TOperator::EOpPackHalf2x16:
3725             pack_half_2x16();
3726             break;
3727         case TOperator::EOpUnpackHalf2x16:
3728             unpack_half_2x16();
3729             break;
3730 
3731         case TOperator::EOpBitfieldExtract:
3732         case TOperator::EOpBitfieldInsert:
3733         case TOperator::EOpBitfieldReverse:
3734         case TOperator::EOpBitCount:
3735         case TOperator::EOpFindLSB:
3736         case TOperator::EOpFindMSB:
3737         case TOperator::EOpUaddCarry:
3738         case TOperator::EOpUsubBorrow:
3739         case TOperator::EOpUmulExtended:
3740         case TOperator::EOpImulExtended:
3741         case TOperator::EOpBarrier:
3742         case TOperator::EOpMemoryBarrier:
3743         case TOperator::EOpMemoryBarrierAtomicCounter:
3744         case TOperator::EOpMemoryBarrierBuffer:
3745         case TOperator::EOpMemoryBarrierImage:
3746         case TOperator::EOpMemoryBarrierShared:
3747         case TOperator::EOpGroupMemoryBarrier:
3748         case TOperator::EOpAtomicAdd:
3749         case TOperator::EOpAtomicMin:
3750         case TOperator::EOpAtomicMax:
3751         case TOperator::EOpAtomicAnd:
3752         case TOperator::EOpAtomicOr:
3753         case TOperator::EOpAtomicXor:
3754         case TOperator::EOpAtomicExchange:
3755         case TOperator::EOpAtomicCompSwap:
3756         case TOperator::EOpEmitVertex:
3757         case TOperator::EOpEndPrimitive:
3758         case TOperator::EOpFtransform:
3759         case TOperator::EOpPackDouble2x32:
3760         case TOperator::EOpUnpackDouble2x32:
3761         case TOperator::EOpArrayLength:
3762             UNIMPLEMENTED();
3763             break;
3764 
3765         case TOperator::EOpConstruct:
3766             ASSERT(!func);
3767             break;
3768 
3769         case TOperator::EOpCallFunctionInAST:
3770         case TOperator::EOpCallInternalRawFunction:
3771         default:
3772             ASSERT(func);
3773             if (mHandled.insert(func).second)
3774             {
3775                 const Name name(*func);
3776                 const auto end = mFuncToEmitter.end();
3777                 auto iter      = mFuncToEmitter.find(name);
3778                 if (iter == end)
3779                 {
3780                     char buffer[32];
3781                     auto mask = MaskTemplateArgs(name, sizeof(buffer), buffer);
3782                     if (mask.second)
3783                     {
3784                         iter = mFuncToEmitter.find(mask.first);
3785                     }
3786                 }
3787                 if (iter != end)
3788                 {
3789                     const auto &emitter = iter->second;
3790                     emitter(*this, *func);
3791                 }
3792             }
3793             break;
3794     }
3795 }
3796 
visitVariable(const Name & name,const TType & type)3797 void ProgramPrelude::visitVariable(const Name &name, const TType &type)
3798 {
3799     if (const TStructure *s = type.getStruct())
3800     {
3801         const Name typeName(*s);
3802         if (typeName.beginsWith(Name("TextureEnv<")))
3803         {
3804             textureEnv();
3805         }
3806     }
3807     else
3808     {
3809         if (name.rawName() == sh::mtl::kRasterizerDiscardEnabledConstName)
3810         {
3811             functionConstants();
3812         }
3813     }
3814 }
3815 
visitVariable(const TVariable & var)3816 void ProgramPrelude::visitVariable(const TVariable &var)
3817 {
3818     if (mHandled.insert(&var).second)
3819     {
3820         visitVariable(Name(var), var.getType());
3821     }
3822 }
3823 
visitStructure(const TStructure & s)3824 void ProgramPrelude::visitStructure(const TStructure &s)
3825 {
3826     if (mHandled.insert(&s).second)
3827     {
3828         for (const TField *field : s.fields())
3829         {
3830             const TType &type = *field->type();
3831             visitVariable(Name(*field), type);
3832         }
3833     }
3834 }
3835 
visitBinary(Visit visit,TIntermBinary * node)3836 bool ProgramPrelude::visitBinary(Visit visit, TIntermBinary *node)
3837 {
3838     const TType &leftType  = node->getLeft()->getType();
3839     const TType &rightType = node->getRight()->getType();
3840     visitOperator(node->getOp(), nullptr, &leftType, &rightType);
3841     return true;
3842 }
3843 
visitUnary(Visit visit,TIntermUnary * node)3844 bool ProgramPrelude::visitUnary(Visit visit, TIntermUnary *node)
3845 {
3846     const TType &argType = node->getOperand()->getType();
3847     visitOperator(node->getOp(), nullptr, &argType);
3848     return true;
3849 }
3850 
visitAggregate(Visit visit,TIntermAggregate * node)3851 bool ProgramPrelude::visitAggregate(Visit visit, TIntermAggregate *node)
3852 {
3853     const size_t argCount = node->getChildCount();
3854 
3855     auto getArgType = [node, argCount](size_t index) -> const TType & {
3856         ASSERT(index < argCount);
3857         TIntermTyped *arg = node->getChildNode(index)->getAsTyped();
3858         ASSERT(arg);
3859         return arg->getType();
3860     };
3861 
3862     const TFunction *func = node->getFunction();
3863 
3864     switch (node->getChildCount())
3865     {
3866         case 0:
3867         {
3868             visitOperator(node->getOp(), func, nullptr);
3869         }
3870         break;
3871 
3872         case 1:
3873         {
3874             const TType &argType0 = getArgType(0);
3875             visitOperator(node->getOp(), func, &argType0);
3876         }
3877         break;
3878 
3879         case 2:
3880         {
3881             const TType &argType0 = getArgType(0);
3882             const TType &argType1 = getArgType(1);
3883             visitOperator(node->getOp(), func, &argType0, &argType1);
3884         }
3885         break;
3886 
3887         default:
3888         {
3889             const TType &argType0 = getArgType(0);
3890             const TType &argType1 = getArgType(1);
3891             visitOperator(node->getOp(), func, &argType0, &argType1);
3892         }
3893         break;
3894     }
3895 
3896     return true;
3897 }
3898 
visitDeclaration(Visit,TIntermDeclaration * node)3899 bool ProgramPrelude::visitDeclaration(Visit, TIntermDeclaration *node)
3900 {
3901     Declaration decl  = ViewDeclaration(*node);
3902     const TType &type = decl.symbol.getType();
3903     if (type.isStructSpecifier())
3904     {
3905         const TStructure *s = type.getStruct();
3906         ASSERT(s);
3907         visitStructure(*s);
3908     }
3909     return true;
3910 }
3911 
visitSymbol(TIntermSymbol * node)3912 void ProgramPrelude::visitSymbol(TIntermSymbol *node)
3913 {
3914     visitVariable(node->variable());
3915 }
3916 
EmitProgramPrelude(TIntermBlock & root,TInfoSinkBase & out,const ProgramPreludeConfig & ppc)3917 bool sh::EmitProgramPrelude(TIntermBlock &root, TInfoSinkBase &out, const ProgramPreludeConfig &ppc)
3918 {
3919     ProgramPrelude programPrelude(out, ppc);
3920     root.traverse(&programPrelude);
3921     return true;
3922 }
3923