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