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