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