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