Google Git
Sign in
android / platform / external / angle / HEAD / . / src / compiler / translator / msl / ProgramPrelude.cpp
blob: ec117053caf8ef376f0f5e27729b46964c6a2d26 [file] [log] [blame]
//
// Copyright 2020 The ANGLE Project Authors. All rights reserved.
// Use of this source code is governed by a BSD-style license that can be
// found in the LICENSE file.
//
#include <cctype>
#include "compiler/translator/InfoSink.h"
#include "compiler/translator/Name.h"
#include "compiler/translator/Symbol.h"
#include "compiler/translator/msl/AstHelpers.h"
#include "compiler/translator/msl/ProgramPrelude.h"
#include "compiler/translator/tree_util/IntermTraverse.h"
#include "compiler/translator/util.h"
using namespace sh;
////////////////////////////////////////////////////////////////////////////////
namespace
{
class ProgramPrelude : public TIntermTraverser
{
using LineTag = unsigned;
using FuncEmitter = void (*)(ProgramPrelude &, const TFunction &);
using FuncToEmitter = std::map<Name, FuncEmitter>;
public:
ProgramPrelude(TInfoSinkBase &out, const ProgramPreludeConfig &ppc)
: TIntermTraverser(true, false, false), mOut(out)
{
mOut << "#include <metal_stdlib>\n\n";
ALWAYS_INLINE();
int_clamp();
switch (ppc.shaderType)
{
case MetalShaderType::None:
ASSERT(0 && "ppc.shaderType should not be ShaderTypeNone");
break;
case MetalShaderType::Vertex:
transform_feedback_guard();
break;
case MetalShaderType::Fragment:
functionConstants();
mOut << "constant bool " << mtl::kSampleMaskWriteEnabledConstName << " = "
<< mtl::kMultisampledRenderingConstName;
if (ppc.usesDerivatives)
{
mOut << " || " << mtl::kWriteHelperSampleMaskConstName;
}
mOut << ";\n";
break;
case MetalShaderType::Compute:
ASSERT(0 && "compute shaders not currently supported");
break;
default:
break;
}
mOut << "#pragma clang diagnostic ignored \"-Wunused-value\"\n";
}
private:
bool emitGuard(LineTag lineTag)
{
if (mEmitted.find(lineTag) != mEmitted.end())
{
return false;
}
mEmitted.insert(lineTag);
return true;
}
static FuncToEmitter BuildFuncToEmitter();
void visitOperator(TOperator op, const TFunction *func, const TType *argType0);
void visitOperator(TOperator op,
const TFunction *func,
const TType *argType0,
const TType *argType1);
void visitOperator(TOperator op,
const TFunction *func,
const TType *argType0,
const TType *argType1,
const TType *argType2);
void visitVariable(const Name &name, const TType &type);
void visitVariable(const TVariable &var);
void visitStructure(const TStructure &s);
bool visitBinary(Visit, TIntermBinary *node) override;
bool visitUnary(Visit, TIntermUnary *node) override;
bool visitAggregate(Visit, TIntermAggregate *node) override;
bool visitDeclaration(Visit, TIntermDeclaration *node) override;
void visitSymbol(TIntermSymbol *node) override;
private:
void ALWAYS_INLINE();
void transform_feedback_guard();
void enable_if();
void addressof();
void distanceScalar();
void faceforwardScalar();
void reflectScalar();
void refractScalar();
void degrees();
void radians();
void mod();
void mixBool();
void postIncrementMatrix();
void preIncrementMatrix();
void postDecrementMatrix();
void preDecrementMatrix();
void negateMatrix();
void matmulAssign();
void int_clamp();
void addMatrixScalarAssign();
void subMatrixScalarAssign();
void addMatrixScalar();
void addScalarMatrix();
void subMatrixScalar();
void subScalarMatrix();
void divScalarMatrix();
void componentWiseDivide();
void componentWiseDivideAssign();
void componentWiseMultiply();
void outerProduct();
void inverse2();
void inverse3();
void inverse4();
void equalScalar();
void equalVector();
void equalMatrix();
void notEqualVector();
void notEqualStruct();
void notEqualStructArray();
void notEqualMatrix();
void equalArray();
void equalStructArray();
void notEqualArray();
void signInt();
void pack_half_2x16();
void unpack_half_2x16();
void vectorElemRef();
void swizzleRef();
void out();
void inout();
void flattenArray();
void castVector();
void castMatrix();
void functionConstants();
void textureEnv();
void texelFetch_2D();
void texelFetch_3D();
void texelFetch_2DArray();
void texelFetch_2DMS();
void texelFetchOffset_2D();
void texelFetchOffset_3D();
void texelFetchOffset_2DArray();
void texture_2D();
void texture_3D();
void texture_Cube();
void texture_2DArray();
void texture_2DShadow();
void texture_CubeShadow();
void texture_2DArrayShadow();
void textureBias_2D();
void textureBias_3D();
void textureBias_Cube();
void textureBias_2DArray();
void textureBias_2DShadow();
void textureBias_CubeShadow();
void textureBias_2DArrayShadow();
void texture2D();
void texture2DBias();
void texture2DGradEXT();
void texture2DLod();
void texture2DLodEXT();
void texture2DProj();
void texture2DProjBias();
void texture2DProjGradEXT();
void texture2DProjLod();
void texture2DProjLodEXT();
void texture3D();
void texture3DBias();
void texture3DLod();
void texture3DProj();
void texture3DProjBias();
void texture3DProjLod();
void textureCube();
void textureCubeBias();
void textureCubeGradEXT();
void textureCubeLod();
void textureCubeLodEXT();
void textureGrad_2D();
void textureGrad_3D();
void textureGrad_Cube();
void textureGrad_2DArray();
void textureGrad_2DShadow();
void textureGrad_CubeShadow();
void textureGrad_2DArrayShadow();
void textureGradOffset_2D();
void textureGradOffset_3D();
void textureGradOffset_2DArray();
void textureGradOffset_2DShadow();
void textureGradOffset_2DArrayShadow();
void textureLod_2D();
void textureLod_3D();
void textureLod_Cube();
void textureLod_2DArray();
void textureLod_2DShadow();
void textureLod_CubeShadow();
void textureLod_2DArrayShadow();
void textureLodOffset_2D();
void textureLodOffset_3D();
void textureLodOffset_2DArray();
void textureLodOffset_2DShadow();
void textureLodOffset_2DArrayShadow();
void textureOffset_2D();
void textureOffset_3D();
void textureOffset_2DArray();
void textureOffset_2DShadow();
void textureOffset_2DArrayShadow();
void textureOffsetBias_2D();
void textureOffsetBias_3D();
void textureOffsetBias_2DArray();
void textureOffsetBias_2DShadow();
void textureOffsetBias_2DArrayShadow();
void textureProj_2D_float3();
void textureProj_2D_float4();
void textureProj_2DShadow();
void textureProj_3D();
void textureProjBias_2D_float3();
void textureProjBias_2D_float4();
void textureProjBias_2DShadow();
void textureProjBias_3D();
void textureProjGrad_2D_float3();
void textureProjGrad_2D_float4();
void textureProjGrad_2DShadow();
void textureProjGrad_3D();
void textureProjGradOffset_2D_float3();
void textureProjGradOffset_2D_float4();
void textureProjGradOffset_2DShadow();
void textureProjGradOffset_3D();
void textureProjLod_2D_float3();
void textureProjLod_2D_float4();
void textureProjLod_2DShadow();
void textureProjLod_3D();
void textureProjLodOffset_2D_float3();
void textureProjLodOffset_2D_float4();
void textureProjLodOffset_2DShadow();
void textureProjLodOffset_3D();
void textureProjOffset_2D_float3();
void textureProjOffset_2D_float4();
void textureProjOffset_2DShadow();
void textureProjOffset_3D();
void textureProjOffsetBias_2D_float3();
void textureProjOffsetBias_2D_float4();
void textureProjOffsetBias_2DShadow();
void textureProjOffsetBias_3D();
void textureSize_2D();
void textureSize_3D();
void textureSize_2DArray();
void textureSize_2DArrayShadow();
void textureSize_2DMS();
void imageLoad();
void imageStore();
void memoryBarrierImage();
void interpolateAtCenter();
void interpolateAtCentroid();
void interpolateAtSample();
void interpolateAtOffset();
private:
TInfoSinkBase &mOut;
std::unordered_set<LineTag> mEmitted;
std::unordered_set<const TSymbol *> mHandled;
const FuncToEmitter mFuncToEmitter = BuildFuncToEmitter();
};
} // anonymous namespace
////////////////////////////////////////////////////////////////////////////////
#define PROGRAM_PRELUDE_DECLARE(name, code, ...) \
void ProgramPrelude::name() \
{ \
ASSERT(code[0] == '\n'); \
if (emitGuard(__LINE__)) \
{ \
__VA_ARGS__; /* dependencies */ \
mOut << (static_cast<const char *>(code "\n") + 1); \
} \
}
////////////////////////////////////////////////////////////////////////////////
PROGRAM_PRELUDE_DECLARE(transform_feedback_guard, R"(
#if TRANSFORM_FEEDBACK_ENABLED
#define __VERTEX_OUT(args) void
#else
#define __VERTEX_OUT(args) args
#endif
)")
PROGRAM_PRELUDE_DECLARE(ALWAYS_INLINE, R"(
#define ANGLE_ALWAYS_INLINE __attribute__((always_inline))
)")
PROGRAM_PRELUDE_DECLARE(enable_if, R"(
template <bool B, typename T = void>
struct ANGLE_enable_if {};
template <typename T>
struct ANGLE_enable_if<true, T>
{
using type = T;
};
template <bool B>
using ANGLE_enable_if_t = typename ANGLE_enable_if<B>::type;
)")
PROGRAM_PRELUDE_DECLARE(addressof,
R"(
template <typename T>
ANGLE_ALWAYS_INLINE thread T * ANGLE_addressof(thread T &ref)
{
return &ref;
}
)")
PROGRAM_PRELUDE_DECLARE(distanceScalar,
R"(
template <typename T>
ANGLE_ALWAYS_INLINE T ANGLE_distance_scalar(T x, T y)
{
return metal::abs(x - y);
}
)")
PROGRAM_PRELUDE_DECLARE(faceforwardScalar,
R"(
template <typename T>
ANGLE_ALWAYS_INLINE T ANGLE_faceforward_scalar(T n, T i, T nref)
{
return nref * i < T(0) ? n : -n;
}
)")
PROGRAM_PRELUDE_DECLARE(reflectScalar,
R"(
template <typename T>
ANGLE_ALWAYS_INLINE T ANGLE_reflect_scalar(T i, T n)
{
return i - T(2) * (n * i) * n;
}
)")
PROGRAM_PRELUDE_DECLARE(refractScalar,
R"(
template <typename T>
ANGLE_ALWAYS_INLINE T ANGLE_refract_scalar(T i, T n, T eta)
{
auto dotNI = n * i;
auto k = T(1) - eta * eta * (T(1) - dotNI * dotNI);
if (k < T(0))
{
return T(0);
}
else
{
return eta * i - (eta * dotNI + metal::sqrt(k)) * n;
}
}
)")
PROGRAM_PRELUDE_DECLARE(signInt,
R"(
ANGLE_ALWAYS_INLINE int ANGLE_sign_int(int x)
{
return (0 < x) - (x < 0);
}
template <int N>
ANGLE_ALWAYS_INLINE metal::vec<int, N> ANGLE_sign_int(metal::vec<int, N> x)
{
metal::vec<int, N> s;
for (int i = 0; i < N; ++i)
{
s[i] = ANGLE_sign_int(x[i]);
}
return s;
}
)")
PROGRAM_PRELUDE_DECLARE(int_clamp,
R"(
ANGLE_ALWAYS_INLINE int ANGLE_int_clamp(int value, int minValue, int maxValue)
{
return ((value < minValue) ? minValue : ((value > maxValue) ? maxValue : value));
};
)")
PROGRAM_PRELUDE_DECLARE(degrees, R"(
template <typename T>
ANGLE_ALWAYS_INLINE T ANGLE_degrees(T x)
{
return static_cast<T>(57.29577951308232) * x;
}
)")
PROGRAM_PRELUDE_DECLARE(radians, R"(
template <typename T>
ANGLE_ALWAYS_INLINE T ANGLE_radians(T x)
{
return static_cast<T>(1.7453292519943295e-2) * x;
}
)")
PROGRAM_PRELUDE_DECLARE(mod,
R"(
template <typename X, typename Y>
ANGLE_ALWAYS_INLINE X ANGLE_mod(X x, Y y)
{
return x - y * metal::floor(x / y);
}
)")
PROGRAM_PRELUDE_DECLARE(mixBool,
R"(
template <typename T, int N>
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)
{
return metal::mix(a, b, static_cast<metal::vec<T,N>>(c));
}
)")
PROGRAM_PRELUDE_DECLARE(pack_half_2x16,
R"(
ANGLE_ALWAYS_INLINE uint32_t ANGLE_pack_half_2x16(float2 v)
{
return as_type<uint32_t>(half2(v));
}
)")
PROGRAM_PRELUDE_DECLARE(unpack_half_2x16,
R"(
ANGLE_ALWAYS_INLINE float2 ANGLE_unpack_half_2x16(uint32_t x)
{
return float2(as_type<half2>(x));
}
)")
PROGRAM_PRELUDE_DECLARE(matmulAssign, R"(
template <typename T, int Cols, int Rows>
ANGLE_ALWAYS_INLINE thread metal::matrix<T, Cols, Rows> &operator*=(thread metal::matrix<T, Cols, Rows> &a, metal::matrix<T, Cols, Cols> b)
{
a = a * b;
return a;
}
)")
PROGRAM_PRELUDE_DECLARE(postIncrementMatrix,
R"(
template <typename T, int Cols, int Rows>
ANGLE_ALWAYS_INLINE metal::matrix<T, Cols, Rows> operator++(thread metal::matrix<T, Cols, Rows> &a, int)
{
auto b = a;
a += T(1);
return b;
}
)",
addMatrixScalarAssign())
PROGRAM_PRELUDE_DECLARE(preIncrementMatrix,
R"(
template <typename T, int Cols, int Rows>
ANGLE_ALWAYS_INLINE thread metal::matrix<T, Cols, Rows> &operator++(thread metal::matrix<T, Cols, Rows> &a)
{
a += T(1);
return a;
}
)",
addMatrixScalarAssign())
PROGRAM_PRELUDE_DECLARE(postDecrementMatrix,
R"(
template <typename T, int Cols, int Rows>
ANGLE_ALWAYS_INLINE metal::matrix<T, Cols, Rows> operator--(thread metal::matrix<T, Cols, Rows> &a, int)
{
auto b = a;
a -= T(1);
return b;
}
)",
subMatrixScalarAssign())
PROGRAM_PRELUDE_DECLARE(preDecrementMatrix,
R"(
template <typename T, int Cols, int Rows>
ANGLE_ALWAYS_INLINE thread metal::matrix<T, Cols, Rows> &operator--(thread metal::matrix<T, Cols, Rows> &a)
{
a -= T(1);
return a;
}
)",
subMatrixScalarAssign())
PROGRAM_PRELUDE_DECLARE(negateMatrix,
R"(
template <typename T, int Cols, int Rows>
ANGLE_ALWAYS_INLINE metal::matrix<T, Cols, Rows> operator-(metal::matrix<T, Cols, Rows> m)
{
for (size_t col = 0; col < Cols; ++col)
{
thread auto &mCol = m[col];
mCol = -mCol;
}
return m;
}
)")
PROGRAM_PRELUDE_DECLARE(addMatrixScalarAssign, R"(
template <typename T, int Cols, int Rows>
ANGLE_ALWAYS_INLINE thread metal::matrix<T, Cols, Rows> &operator+=(thread metal::matrix<T, Cols, Rows> &m, T x)
{
for (size_t col = 0; col < Cols; ++col)
{
m[col] += x;
}
return m;
}
)")
PROGRAM_PRELUDE_DECLARE(addMatrixScalar,
R"(
template <typename T, int Cols, int Rows>
ANGLE_ALWAYS_INLINE metal::matrix<T, Cols, Rows> operator+(metal::matrix<T, Cols, Rows> m, T x)
{
m += x;
return m;
}
)",
addMatrixScalarAssign())
PROGRAM_PRELUDE_DECLARE(addScalarMatrix,
R"(
template <typename T, int Cols, int Rows>
ANGLE_ALWAYS_INLINE metal::matrix<T, Cols, Rows> operator+(T x, metal::matrix<T, Cols, Rows> m)
{
for (size_t col = 0; col < Cols; ++col)
{
m[col] = x + m[col];
}
return m;
}
)")
PROGRAM_PRELUDE_DECLARE(subMatrixScalarAssign,
R"(
template <typename T, int Cols, int Rows>
ANGLE_ALWAYS_INLINE thread metal::matrix<T, Cols, Rows> &operator-=(thread metal::matrix<T, Cols, Rows> &m, T x)
{
for (size_t col = 0; col < Cols; ++col)
{
m[col] -= x;
}
return m;
}
)")
PROGRAM_PRELUDE_DECLARE(subMatrixScalar,
R"(
template <typename T, int Cols, int Rows>
ANGLE_ALWAYS_INLINE metal::matrix<T, Cols, Rows> operator-(metal::matrix<T, Cols, Rows> m, T x)
{
m -= x;
return m;
}
)",
subMatrixScalarAssign())
PROGRAM_PRELUDE_DECLARE(subScalarMatrix,
R"(
template <typename T, int Cols, int Rows>
ANGLE_ALWAYS_INLINE metal::matrix<T, Cols, Rows> operator-(T x, metal::matrix<T, Cols, Rows> m)
{
for (size_t col = 0; col < Cols; ++col)
{
m[col] = x - m[col];
}
return m;
}
)")
PROGRAM_PRELUDE_DECLARE(divScalarMatrix,
R"(
template <typename T, int Cols, int Rows>
ANGLE_ALWAYS_INLINE metal::matrix<T, Cols, Rows> operator/(T x, metal::matrix<T, Cols, Rows> m)
{
for (size_t col = 0; col < Cols; ++col)
{
m[col] = x / m[col];
}
return m;
}
)")
PROGRAM_PRELUDE_DECLARE(componentWiseDivide, R"(
template <typename T, int Cols, int Rows>
ANGLE_ALWAYS_INLINE metal::matrix<T, Cols, Rows> operator/(metal::matrix<T, Cols, Rows> a, metal::matrix<T, Cols, Rows> b)
{
for (size_t col = 0; col < Cols; ++col)
{
a[col] /= b[col];
}
return a;
}
)")
PROGRAM_PRELUDE_DECLARE(componentWiseDivideAssign,
R"(
template <typename T, int Cols, int Rows>
ANGLE_ALWAYS_INLINE thread metal::matrix<T, Cols, Rows> &operator/=(thread metal::matrix<T, Cols, Rows> &a, metal::matrix<T, Cols, Rows> b)
{
a = a / b;
return a;
}
)",
componentWiseDivide())
PROGRAM_PRELUDE_DECLARE(componentWiseMultiply, R"(
template <typename T, int Cols, int Rows>
ANGLE_ALWAYS_INLINE metal::matrix<T, Cols, Rows> ANGLE_componentWiseMultiply(metal::matrix<T, Cols, Rows> a, metal::matrix<T, Cols, Rows> b)
{
for (size_t col = 0; col < Cols; ++col)
{
a[col] *= b[col];
}
return a;
}
)")
PROGRAM_PRELUDE_DECLARE(outerProduct, R"(
template <typename T, int M, int N>
ANGLE_ALWAYS_INLINE metal::matrix<T, N, M> ANGLE_outerProduct(metal::vec<T, M> u, metal::vec<T, N> v)
{
metal::matrix<T, N, M> o;
for (size_t n = 0; n < N; ++n)
{
o[n] = u * v[n];
}
return o;
}
)")
PROGRAM_PRELUDE_DECLARE(inverse2, R"(
template <typename T>
ANGLE_ALWAYS_INLINE metal::matrix<T, 2, 2> ANGLE_inverse(metal::matrix<T, 2, 2> m)
{
metal::matrix<T, 2, 2> adj;
adj[0][0] = m[1][1];
adj[0][1] = -m[0][1];
adj[1][0] = -m[1][0];
adj[1][1] = m[0][0];
T det = (adj[0][0] * m[0][0]) + (adj[0][1] * m[1][0]);
return adj * (T(1) / det);
}
)")
PROGRAM_PRELUDE_DECLARE(inverse3, R"(
template <typename T>
ANGLE_ALWAYS_INLINE metal::matrix<T, 3, 3> ANGLE_inverse(metal::matrix<T, 3, 3> m)
{
T a = m[1][1] * m[2][2] - m[2][1] * m[1][2];
T b = m[1][0] * m[2][2];
T c = m[1][2] * m[2][0];
T d = m[1][0] * m[2][1];
T det = m[0][0] * (a) -
m[0][1] * (b - c) +
m[0][2] * (d - m[1][1] * m[2][0]);
det = T(1) / det;
metal::matrix<T, 3, 3> minv;
minv[0][0] = (a) * det;
minv[0][1] = (m[0][2] * m[2][1] - m[0][1] * m[2][2]) * det;
minv[0][2] = (m[0][1] * m[1][2] - m[0][2] * m[1][1]) * det;
minv[1][0] = (c - b) * det;
minv[1][1] = (m[0][0] * m[2][2] - m[0][2] * m[2][0]) * det;
minv[1][2] = (m[1][0] * m[0][2] - m[0][0] * m[1][2]) * det;
minv[2][0] = (d - m[2][0] * m[1][1]) * det;
minv[2][1] = (m[2][0] * m[0][1] - m[0][0] * m[2][1]) * det;
minv[2][2] = (m[0][0] * m[1][1] - m[1][0] * m[0][1]) * det;
return minv;
}
)")
PROGRAM_PRELUDE_DECLARE(inverse4, R"(
template <typename T>
ANGLE_ALWAYS_INLINE metal::matrix<T, 4, 4> ANGLE_inverse(metal::matrix<T, 4, 4> m)
{
T A2323 = m[2][2] * m[3][3] - m[2][3] * m[3][2];
T A1323 = m[2][1] * m[3][3] - m[2][3] * m[3][1];
T A1223 = m[2][1] * m[3][2] - m[2][2] * m[3][1];
T A0323 = m[2][0] * m[3][3] - m[2][3] * m[3][0];
T A0223 = m[2][0] * m[3][2] - m[2][2] * m[3][0];
T A0123 = m[2][0] * m[3][1] - m[2][1] * m[3][0];
T A2313 = m[1][2] * m[3][3] - m[1][3] * m[3][2];
T A1313 = m[1][1] * m[3][3] - m[1][3] * m[3][1];
T A1213 = m[1][1] * m[3][2] - m[1][2] * m[3][1];
T A2312 = m[1][2] * m[2][3] - m[1][3] * m[2][2];
T A1312 = m[1][1] * m[2][3] - m[1][3] * m[2][1];
T A1212 = m[1][1] * m[2][2] - m[1][2] * m[2][1];
T A0313 = m[1][0] * m[3][3] - m[1][3] * m[3][0];
T A0213 = m[1][0] * m[3][2] - m[1][2] * m[3][0];
T A0312 = m[1][0] * m[2][3] - m[1][3] * m[2][0];
T A0212 = m[1][0] * m[2][2] - m[1][2] * m[2][0];
T A0113 = m[1][0] * m[3][1] - m[1][1] * m[3][0];
T A0112 = m[1][0] * m[2][1] - m[1][1] * m[2][0];
T a = m[1][1] * A2323 - m[1][2] * A1323 + m[1][3] * A1223;
T b = m[1][0] * A2323 - m[1][2] * A0323 + m[1][3] * A0223;
T c = m[1][0] * A1323 - m[1][1] * A0323 + m[1][3] * A0123;
T d = m[1][0] * A1223 - m[1][1] * A0223 + m[1][2] * A0123;
T det = m[0][0] * ( a )
- m[0][1] * ( b )
+ m[0][2] * ( c )
- m[0][3] * ( d );
det = T(1) / det;
metal::matrix<T, 4, 4> im;
im[0][0] = det * ( a );
im[0][1] = det * - ( m[0][1] * A2323 - m[0][2] * A1323 + m[0][3] * A1223 );
im[0][2] = det * ( m[0][1] * A2313 - m[0][2] * A1313 + m[0][3] * A1213 );
im[0][3] = det * - ( m[0][1] * A2312 - m[0][2] * A1312 + m[0][3] * A1212 );
im[1][0] = det * - ( b );
im[1][1] = det * ( m[0][0] * A2323 - m[0][2] * A0323 + m[0][3] * A0223 );
im[1][2] = det * - ( m[0][0] * A2313 - m[0][2] * A0313 + m[0][3] * A0213 );
im[1][3] = det * ( m[0][0] * A2312 - m[0][2] * A0312 + m[0][3] * A0212 );
im[2][0] = det * ( c );
im[2][1] = det * - ( m[0][0] * A1323 - m[0][1] * A0323 + m[0][3] * A0123 );
im[2][2] = det * ( m[0][0] * A1313 - m[0][1] * A0313 + m[0][3] * A0113 );
im[2][3] = det * - ( m[0][0] * A1312 - m[0][1] * A0312 + m[0][3] * A0112 );
im[3][0] = det * - ( d );
im[3][1] = det * ( m[0][0] * A1223 - m[0][1] * A0223 + m[0][2] * A0123 );
im[3][2] = det * - ( m[0][0] * A1213 - m[0][1] * A0213 + m[0][2] * A0113 );
im[3][3] = det * ( m[0][0] * A1212 - m[0][1] * A0212 + m[0][2] * A0112 );
return im;
}
)")
PROGRAM_PRELUDE_DECLARE(equalArray,
R"(
template <typename T, size_t N>
ANGLE_ALWAYS_INLINE bool ANGLE_equal(metal::array<T, N> u, metal::array<T, N> v)
{
for(size_t i = 0; i < N; i++)
if (!ANGLE_equal(u[i], v[i])) return false;
return true;
}
)",
equalScalar(),
equalVector(),
equalMatrix())
PROGRAM_PRELUDE_DECLARE(equalStructArray,
R"(
template <typename T, size_t N>
ANGLE_ALWAYS_INLINE bool ANGLE_equalStructArray(metal::array<T, N> u, metal::array<T, N> v)
{
for(size_t i = 0; i < N; i++)
{
if (!ANGLE_equal(u[i], v[i]))
return false;
}
return true;
}
)")
PROGRAM_PRELUDE_DECLARE(notEqualArray,
R"(
template <typename T, size_t N>
ANGLE_ALWAYS_INLINE bool ANGLE_notEqual(metal::array<T, N> u, metal::array<T, N> v)
{
return !ANGLE_equal(u,v);
}
)",
equalArray())
PROGRAM_PRELUDE_DECLARE(equalScalar,
R"(
template <typename T>
ANGLE_ALWAYS_INLINE bool ANGLE_equal(T u, T v)
{
return u == v;
}
)")
PROGRAM_PRELUDE_DECLARE(equalVector,
R"(
template <typename T, int N>
ANGLE_ALWAYS_INLINE bool ANGLE_equal(metal::vec<T, N> u, metal::vec<T, N> v)
{
return metal::all(u == v);
}
)")
PROGRAM_PRELUDE_DECLARE(equalMatrix,
R"(
template <typename T, int C, int R>
ANGLE_ALWAYS_INLINE bool ANGLE_equal(metal::matrix<T, C, R> a, metal::matrix<T, C, R> b)
{
for (int c = 0; c < C; ++c)
{
if (!ANGLE_equal(a[c], b[c]))
{
return false;
}
}
return true;
}
)",
equalVector())
PROGRAM_PRELUDE_DECLARE(notEqualMatrix,
R"(
template <typename T, int C, int R>
ANGLE_ALWAYS_INLINE bool ANGLE_notEqual(metal::matrix<T, C, R> u, metal::matrix<T, C, R> v)
{
return !ANGLE_equal(u, v);
}
)",
equalMatrix())
PROGRAM_PRELUDE_DECLARE(notEqualVector,
R"(
template <typename T, int N>
ANGLE_ALWAYS_INLINE bool ANGLE_notEqual(metal::vec<T, N> u, metal::vec<T, N> v)
{
return !ANGLE_equal(u, v);
}
)",
equalVector())
PROGRAM_PRELUDE_DECLARE(notEqualStruct,
R"(
template <typename T>
ANGLE_ALWAYS_INLINE bool ANGLE_notEqualStruct(thread const T &a, thread const T &b)
{
return !ANGLE_equal(a, b);
}
template <typename T>
ANGLE_ALWAYS_INLINE bool ANGLE_notEqualStruct(constant const T &a, thread const T &b)
{
return !ANGLE_equal(a, b);
}
template <typename T>
ANGLE_ALWAYS_INLINE bool ANGLE_notEqualStruct(thread const T &a, constant const T &b)
{
return !ANGLE_equal(a, b);
}
template <typename T>
ANGLE_ALWAYS_INLINE bool ANGLE_notEqualStruct(constant const T &a, constant const T &b)
{
return !ANGLE_equal(a, b);
}
)",
equalVector(),
equalMatrix())
PROGRAM_PRELUDE_DECLARE(notEqualStructArray,
R"(
template <typename T, size_t N>
ANGLE_ALWAYS_INLINE bool ANGLE_notEqualStructArray(metal::array<T, N> u, metal::array<T, N> v)
{
for(size_t i = 0; i < N; i++)
{
if (ANGLE_notEqualStruct(u[i], v[i]))
return true;
}
return false;
}
)",
notEqualStruct())
PROGRAM_PRELUDE_DECLARE(vectorElemRef,
R"(
template <typename T, int N>
struct ANGLE_VectorElemRef
{
thread metal::vec<T, N> &mVec;
T mRef;
const int mIndex;
~ANGLE_VectorElemRef() { mVec[mIndex] = mRef; }
ANGLE_VectorElemRef(thread metal::vec<T, N> &vec, int index)
: mVec(vec), mRef(vec[index]), mIndex(index)
{}
operator thread T &() { return mRef; }
};
template <typename T, int N>
ANGLE_ALWAYS_INLINE ANGLE_VectorElemRef<T, N> ANGLE_elem_ref(thread metal::vec<T, N> &vec, int index)
{
return ANGLE_VectorElemRef<T, N>(vec, metal::clamp(index, 0, N - 1));
}
)")
PROGRAM_PRELUDE_DECLARE(swizzleRef,
R"(
template <typename T, int VN, int SN>
struct ANGLE_SwizzleRef
{
thread metal::vec<T, VN> &mVec;
metal::vec<T, SN> mRef;
int mIndices[SN];
~ANGLE_SwizzleRef()
{
for (int i = 0; i < SN; ++i)
{
const int j = mIndices[i];
mVec[j] = mRef[i];
}
}
ANGLE_SwizzleRef(thread metal::vec<T, VN> &vec, thread const int *indices)
: mVec(vec)
{
for (int i = 0; i < SN; ++i)
{
const int j = indices[i];
mIndices[i] = j;
mRef[i] = mVec[j];
}
}
operator thread metal::vec<T, SN> &() { return mRef; }
};
template <typename T, int N>
ANGLE_ALWAYS_INLINE ANGLE_VectorElemRef<T, N> ANGLE_swizzle_ref(thread metal::vec<T, N> &vec, int i0)
{
return ANGLE_VectorElemRef<T, N>(vec, i0);
}
template <typename T, int N>
ANGLE_ALWAYS_INLINE ANGLE_SwizzleRef<T, N, 2> ANGLE_swizzle_ref(thread metal::vec<T, N> &vec, int i0, int i1)
{
const int is[] = { i0, i1 };
return ANGLE_SwizzleRef<T, N, 2>(vec, is);
}
template <typename T, int N>
ANGLE_ALWAYS_INLINE ANGLE_SwizzleRef<T, N, 3> ANGLE_swizzle_ref(thread metal::vec<T, N> &vec, int i0, int i1, int i2)
{
const int is[] = { i0, i1, i2 };
return ANGLE_SwizzleRef<T, N, 3>(vec, is);
}
template <typename T, int N>
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)
{
const int is[] = { i0, i1, i2, i3 };
return ANGLE_SwizzleRef<T, N, 4>(vec, is);
}
)",
vectorElemRef())
PROGRAM_PRELUDE_DECLARE(out, R"(
template <typename T>
struct ANGLE_Out
{
T mTemp;
thread T &mDest;
~ANGLE_Out() { mDest = mTemp; }
ANGLE_Out(thread T &dest)
: mTemp(dest), mDest(dest)
{}
operator thread T &() { return mTemp; }
};
template <typename T>
ANGLE_ALWAYS_INLINE ANGLE_Out<T> ANGLE_out(thread T &dest)
{
return ANGLE_Out<T>(dest);
}
)")
PROGRAM_PRELUDE_DECLARE(inout, R"(
template <typename T>
struct ANGLE_InOut
{
T mTemp;
thread T &mDest;
~ANGLE_InOut() { mDest = mTemp; }
ANGLE_InOut(thread T &dest)
: mTemp(dest), mDest(dest)
{}
operator thread T &() { return mTemp; }
};
template <typename T>
ANGLE_ALWAYS_INLINE ANGLE_InOut<T> ANGLE_inout(thread T &dest)
{
return ANGLE_InOut<T>(dest);
}
)")
PROGRAM_PRELUDE_DECLARE(flattenArray, R"(
template <typename T>
struct ANGLE_flatten_impl
{
static ANGLE_ALWAYS_INLINE thread T *exec(thread T &x)
{
return &x;
}
};
template <typename T, size_t N>
struct ANGLE_flatten_impl<metal::array<T, N>>
{
static ANGLE_ALWAYS_INLINE auto exec(thread metal::array<T, N> &arr) -> T
{
return ANGLE_flatten_impl<T>::exec(arr[0]);
}
};
template <typename T, size_t N>
ANGLE_ALWAYS_INLINE auto ANGLE_flatten(thread metal::array<T, N> &arr) -> T
{
return ANGLE_flatten_impl<T>::exec(arr[0]);
}
)")
PROGRAM_PRELUDE_DECLARE(castVector, R"(
template <typename T, int N1, int N2>
struct ANGLE_castVector {};
template <typename T, int N>
struct ANGLE_castVector<T, N, N>
{
static ANGLE_ALWAYS_INLINE metal::vec<T, N> exec(metal::vec<T, N> const v)
{
return v;
}
};
template <typename T>
struct ANGLE_castVector<T, 2, 3>
{
static ANGLE_ALWAYS_INLINE metal::vec<T, 2> exec(metal::vec<T, 3> const v)
{
return v.xy;
}
};
template <typename T>
struct ANGLE_castVector<T, 2, 4>
{
static ANGLE_ALWAYS_INLINE metal::vec<T, 2> exec(metal::vec<T, 4> const v)
{
return v.xy;
}
};
template <typename T>
struct ANGLE_castVector<T, 3, 4>
{
static ANGLE_ALWAYS_INLINE metal::vec<T, 3> exec(metal::vec<T, 4> const v)
{
return as_type<metal::vec<T, 3>>(v);
}
};
template <int N1, int N2, typename T>
ANGLE_ALWAYS_INLINE metal::vec<T, N1> ANGLE_cast(metal::vec<T, N2> const v)
{
return ANGLE_castVector<T, N1, N2>::exec(v);
}
)")
PROGRAM_PRELUDE_DECLARE(castMatrix,
R"(
template <typename T, int C1, int R1, int C2, int R2, typename Enable = void>
struct ANGLE_castMatrix
{
static ANGLE_ALWAYS_INLINE metal::matrix<T, C1, R1> exec(metal::matrix<T, C2, R2> const m2)
{
metal::matrix<T, C1, R1> m1;
const int MinC = C1 <= C2 ? C1 : C2;
const int MinR = R1 <= R2 ? R1 : R2;
for (int c = 0; c < MinC; ++c)
{
for (int r = 0; r < MinR; ++r)
{
m1[c][r] = m2[c][r];
}
for (int r = R2; r < R1; ++r)
{
m1[c][r] = c == r ? T(1) : T(0);
}
}
for (int c = C2; c < C1; ++c)
{
for (int r = 0; r < R1; ++r)
{
m1[c][r] = c == r ? T(1) : T(0);
}
}
return m1;
}
};
template <typename T, int C1, int R1, int C2, int R2>
struct ANGLE_castMatrix<T, C1, R1, C2, R2, ANGLE_enable_if_t<(C1 <= C2 && R1 <= R2)>>
{
static ANGLE_ALWAYS_INLINE metal::matrix<T, C1, R1> exec(metal::matrix<T, C2, R2> const m2)
{
metal::matrix<T, C1, R1> m1;
for (size_t c = 0; c < C1; ++c)
{
m1[c] = ANGLE_cast<R1>(m2[c]);
}
return m1;
}
};
template <int C1, int R1, int C2, int R2, typename T>
ANGLE_ALWAYS_INLINE metal::matrix<T, C1, R1> ANGLE_cast(metal::matrix<T, C2, R2> const m)
{
return ANGLE_castMatrix<T, C1, R1, C2, R2>::exec(m);
};
)",
enable_if(),
castVector())
PROGRAM_PRELUDE_DECLARE(textureEnv,
R"(
template <typename T>
struct ANGLE_TextureEnv
{
thread T *texture;
thread metal::sampler *sampler;
};
)")
PROGRAM_PRELUDE_DECLARE(functionConstants,
R"(
#define ANGLE_SAMPLE_COMPARE_GRADIENT_INDEX 0
#define ANGLE_RASTERIZATION_DISCARD_INDEX 1
#define ANGLE_MULTISAMPLED_RENDERING_INDEX 2
#define ANGLE_DEPTH_WRITE_ENABLED_INDEX 3
#define ANGLE_EMULATE_ALPHA_TO_COVERAGE_INDEX 4
#define ANGLE_WRITE_HELPER_SAMPLE_MASK_INDEX 5
constant bool ANGLEUseSampleCompareGradient [[function_constant(ANGLE_SAMPLE_COMPARE_GRADIENT_INDEX)]];
constant bool ANGLERasterizerDisabled [[function_constant(ANGLE_RASTERIZATION_DISCARD_INDEX)]];
constant bool ANGLEMultisampledRendering [[function_constant(ANGLE_MULTISAMPLED_RENDERING_INDEX)]];
constant bool ANGLEDepthWriteEnabled [[function_constant(ANGLE_DEPTH_WRITE_ENABLED_INDEX)]];
constant bool ANGLEEmulateAlphaToCoverage [[function_constant(ANGLE_EMULATE_ALPHA_TO_COVERAGE_INDEX)]];
constant bool ANGLEWriteHelperSampleMask [[function_constant(ANGLE_WRITE_HELPER_SAMPLE_MASK_INDEX)]];
#define ANGLE_ALPHA0
)")
PROGRAM_PRELUDE_DECLARE(texelFetch_2D,
R"(
template <typename T>
ANGLE_ALWAYS_INLINE auto ANGLE_texelFetch(
thread ANGLE_TextureEnv<metal::texture2d<T>> &env,
metal::int2 const coord,
int const level)
{
return env.texture->read(uint2(coord), uint32_t(level));
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(texelFetch_3D,
R"(
template <typename T>
ANGLE_ALWAYS_INLINE auto ANGLE_texelFetch(
thread ANGLE_TextureEnv<metal::texture3d<T>> &env,
metal::int3 const coord,
int const level)
{
return env.texture->read(uint3(coord), uint32_t(level));
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(texelFetch_2DArray,
R"(
template <typename T>
ANGLE_ALWAYS_INLINE auto ANGLE_texelFetch(
thread ANGLE_TextureEnv<metal::texture2d_array<T>> &env,
metal::int3 const coord,
int const level)
{
return env.texture->read(uint2(coord.xy), uint32_t(coord.z), uint32_t(level));
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(texelFetch_2DMS,
R"(
template <typename T>
ANGLE_ALWAYS_INLINE auto ANGLE_texelFetch(
thread ANGLE_TextureEnv<metal::texture2d_ms<T>> &env,
metal::int2 const coord,
int const sample)
{
return env.texture->read(uint2(coord), uint32_t(sample));
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(texelFetchOffset_2D,
R"(
template <typename T>
ANGLE_ALWAYS_INLINE auto ANGLE_texelFetchOffset(
thread ANGLE_TextureEnv<metal::texture2d<T>> &env,
metal::int2 const coord,
int const level,
metal::int2 const offset)
{
return env.texture->read(uint2(coord + offset), uint32_t(level));
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(texelFetchOffset_3D,
R"(
template <typename T>
ANGLE_ALWAYS_INLINE auto ANGLE_texelFetchOffset(
thread ANGLE_TextureEnv<metal::texture3d<T>> &env,
metal::int3 const coord,
int const level,
metal::int3 const offset)
{
return env.texture->read(uint3(coord + offset), uint32_t(level));
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(texelFetchOffset_2DArray,
R"(
template <typename T>
ANGLE_ALWAYS_INLINE auto ANGLE_texelFetchOffset(
thread ANGLE_TextureEnv<metal::texture2d_array<T>> &env,
metal::int3 const coord,
int const level,
metal::int2 const offset)
{
return env.texture->read(uint2(coord.xy + offset), uint32_t(coord.z), uint32_t(level));
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(texture_2D,
R"(
template <typename T>
ANGLE_ALWAYS_INLINE auto ANGLE_texture(
thread ANGLE_TextureEnv<metal::texture2d<T>> &env,
metal::float2 const coord)
{
return env.texture->sample(*env.sampler, coord);
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(textureBias_2D,
R"(
template <typename T>
ANGLE_ALWAYS_INLINE auto ANGLE_texture(
thread ANGLE_TextureEnv<metal::texture2d<T>> &env,
metal::float2 const coord,
float const bias)
{
return env.texture->sample(*env.sampler, coord, metal::bias(bias));
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(texture_3D,
R"(
template <typename T>
ANGLE_ALWAYS_INLINE auto ANGLE_texture(
thread ANGLE_TextureEnv<metal::texture3d<T>> &env,
metal::float3 const coord)
{
return env.texture->sample(*env.sampler, coord);
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(textureBias_3D,
R"(
template <typename T>
ANGLE_ALWAYS_INLINE auto ANGLE_texture(
thread ANGLE_TextureEnv<metal::texture3d<T>> &env,
metal::float3 const coord,
float const bias)
{
return env.texture->sample(*env.sampler, coord, metal::bias(bias));
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(texture_Cube,
R"(
template <typename T>
ANGLE_ALWAYS_INLINE auto ANGLE_texture(
thread ANGLE_TextureEnv<metal::texturecube<T>> &env,
metal::float3 const coord)
{
return env.texture->sample(*env.sampler, coord);
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(textureBias_Cube,
R"(
template <typename T>
ANGLE_ALWAYS_INLINE auto ANGLE_texture(
thread ANGLE_TextureEnv<metal::texturecube<T>> &env,
metal::float3 const coord,
float const bias)
{
return env.texture->sample(*env.sampler, coord, metal::bias(bias));
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(texture_2DArray,
R"(
template <typename T>
ANGLE_ALWAYS_INLINE auto ANGLE_texture(
thread ANGLE_TextureEnv<metal::texture2d_array<T>> &env,
metal::float3 const coord)
{
return env.texture->sample(*env.sampler, coord.xy, uint32_t(metal::round(coord.z)));
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(textureBias_2DArray,
R"(
template <typename T>
ANGLE_ALWAYS_INLINE auto ANGLE_texture(
thread ANGLE_TextureEnv<metal::texture2d_array<T>> &env,
metal::float3 const coord,
float const bias)
{
return env.texture->sample(*env.sampler, coord.xy, uint32_t(metal::round(coord.z)), metal::bias(bias));
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(texture_2DShadow,
R"(
ANGLE_ALWAYS_INLINE auto ANGLE_texture(
thread ANGLE_TextureEnv<metal::depth2d<float>> &env,
metal::float3 const coord)
{
return env.texture->sample_compare(*env.sampler, coord.xy, coord.z);
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(textureBias_2DShadow,
R"(
ANGLE_ALWAYS_INLINE auto ANGLE_texture(
thread ANGLE_TextureEnv<metal::depth2d<float>> &env,
metal::float3 const coord,
float const bias)
{
return env.texture->sample_compare(*env.sampler, coord.xy, coord.z, metal::bias(bias));
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(texture_2DArrayShadow,
R"(
ANGLE_ALWAYS_INLINE auto ANGLE_texture(
thread ANGLE_TextureEnv<metal::depth2d_array<float>> &env,
metal::float4 const coord)
{
return env.texture->sample_compare(*env.sampler, coord.xy, uint32_t(metal::round(coord.z)), coord.w);
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(textureBias_2DArrayShadow,
R"(
ANGLE_ALWAYS_INLINE auto ANGLE_texture(
thread ANGLE_TextureEnv<metal::depth2d_array<float>> &env,
metal::float4 const coord,
float const bias)
{
return env.texture->sample_compare(*env.sampler, coord.xy, uint32_t(metal::round(coord.z)), coord.w, metal::bias(bias));
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(texture_CubeShadow,
R"(
ANGLE_ALWAYS_INLINE auto ANGLE_texture(
thread ANGLE_TextureEnv<metal::depthcube<float>> &env,
metal::float4 const coord)
{
return env.texture->sample_compare(*env.sampler, coord.xyz, coord.w);
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(textureBias_CubeShadow,
R"(
ANGLE_ALWAYS_INLINE auto ANGLE_texture(
thread ANGLE_TextureEnv<metal::depthcube<float>> &env,
metal::float4 const coord,
float const bias)
{
return env.texture->sample_compare(*env.sampler, coord.xyz, coord.w, metal::bias(bias));
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(texture2D,
R"(
ANGLE_ALWAYS_INLINE auto ANGLE_texture2D(
thread ANGLE_TextureEnv<metal::texture2d<float>> &env,
metal::float2 const coord)
{
return env.texture->sample(*env.sampler, coord);
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(texture2DBias,
R"(
ANGLE_ALWAYS_INLINE auto ANGLE_texture2D(
thread ANGLE_TextureEnv<metal::texture2d<float>> &env,
metal::float2 const coord,
float const bias)
{
return env.texture->sample(*env.sampler, coord, metal::bias(bias));
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(texture2DGradEXT,
R"(
ANGLE_ALWAYS_INLINE auto ANGLE_texture2DGradEXT(
thread ANGLE_TextureEnv<metal::texture2d<float>> &env,
metal::float2 const coord,
metal::float2 const dPdx,
metal::float2 const dPdy)
{
return env.texture->sample(*env.sampler, coord, metal::gradient2d(dPdx, dPdy));
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(texture2DLod,
R"(
ANGLE_ALWAYS_INLINE auto ANGLE_texture2DLod(
thread ANGLE_TextureEnv<metal::texture2d<float>> &env,
metal::float2 const coord,
float const level)
{
return env.texture->sample(*env.sampler, coord, metal::level(level));
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(texture2DLodEXT,
R"(
#define ANGLE_texture2DLodEXT ANGLE_texture2DLod
)",
texture2DLod())
PROGRAM_PRELUDE_DECLARE(texture2DProj,
R"(
ANGLE_ALWAYS_INLINE auto ANGLE_texture2DProj(
thread ANGLE_TextureEnv<metal::texture2d<float>> &env,
metal::float3 const coord)
{
return env.texture->sample(*env.sampler, coord.xy/coord.z);
}
ANGLE_ALWAYS_INLINE auto ANGLE_texture2DProj(
thread ANGLE_TextureEnv<metal::texture2d<float>> &env,
metal::float4 const coord)
{
return env.texture->sample(*env.sampler, coord.xy/coord.w);
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(texture2DProjBias,
R"(
ANGLE_ALWAYS_INLINE auto ANGLE_texture2DProj(
thread ANGLE_TextureEnv<metal::texture2d<float>> &env,
metal::float3 const coord,
float const bias)
{
return env.texture->sample(*env.sampler, coord.xy/coord.z, metal::bias(bias));
}
ANGLE_ALWAYS_INLINE auto ANGLE_texture2DProj(
thread ANGLE_TextureEnv<metal::texture2d<float>> &env,
metal::float4 const coord,
float const bias)
{
return env.texture->sample(*env.sampler, coord.xy/coord.w, metal::bias(bias));
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(texture2DProjGradEXT,
R"(
ANGLE_ALWAYS_INLINE auto ANGLE_texture2DProjGradEXT(
thread ANGLE_TextureEnv<metal::texture2d<float>> &env,
metal::float3 const coord,
metal::float2 const dPdx,
metal::float2 const dPdy)
{
return env.texture->sample(*env.sampler, coord.xy/coord.z, metal::gradient2d(dPdx, dPdy));
}
ANGLE_ALWAYS_INLINE auto ANGLE_texture2DProjGradEXT(
thread ANGLE_TextureEnv<metal::texture2d<float>> &env,
metal::float4 const coord,
metal::float2 const dPdx,
metal::float2 const dPdy)
{
return env.texture->sample(*env.sampler, coord.xy/coord.w, metal::gradient2d(dPdx, dPdy));
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(texture2DProjLod,
R"(
ANGLE_ALWAYS_INLINE auto ANGLE_texture2DProjLod(
thread ANGLE_TextureEnv<metal::texture2d<float>> &env,
metal::float3 const coord,
float const level)
{
return env.texture->sample(*env.sampler, coord.xy/coord.z, metal::level(level));
}
ANGLE_ALWAYS_INLINE auto ANGLE_texture2DProjLod(
thread ANGLE_TextureEnv<metal::texture2d<float>> &env,
metal::float4 const coord,
float const level)
{
return env.texture->sample(*env.sampler, coord.xy/coord.w, metal::level(level));
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(texture2DProjLodEXT,
R"(
#define ANGLE_texture2DProjLodEXT ANGLE_texture2DProjLod
)",
texture2DProjLod())
PROGRAM_PRELUDE_DECLARE(texture3D,
R"(
ANGLE_ALWAYS_INLINE auto ANGLE_texture3D(
thread ANGLE_TextureEnv<metal::texture3d<float>> &env,
metal::float3 const coord)
{
return env.texture->sample(*env.sampler, coord);
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(texture3DBias,
R"(
ANGLE_ALWAYS_INLINE auto ANGLE_texture3D(
thread ANGLE_TextureEnv<metal::texture3d<float>> &env,
metal::float3 const coord,
float const bias)
{
return env.texture->sample(*env.sampler, coord, metal::bias(bias));
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(texture3DLod,
R"(
ANGLE_ALWAYS_INLINE auto ANGLE_texture3DLod(
thread ANGLE_TextureEnv<metal::texture3d<float>> &env,
metal::float3 const coord,
float const level)
{
return env.texture->sample(*env.sampler, coord, metal::level(level));
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(texture3DProj,
R"(
ANGLE_ALWAYS_INLINE auto ANGLE_texture3DProj(
thread ANGLE_TextureEnv<metal::texture3d<float>> &env,
metal::float4 const coord)
{
return env.texture->sample(*env.sampler, coord.xyz/coord.w);
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(texture3DProjBias,
R"(
ANGLE_ALWAYS_INLINE auto ANGLE_texture3DProj(
thread ANGLE_TextureEnv<metal::texture3d<float>> &env,
metal::float4 const coord,
float const bias)
{
return env.texture->sample(*env.sampler, coord.xyz/coord.w, metal::bias(bias));
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(texture3DProjLod,
R"(
ANGLE_ALWAYS_INLINE auto ANGLE_texture3DProjLod(
thread ANGLE_TextureEnv<metal::texture3d<float>> &env,
metal::float4 const coord,
float const level)
{
return env.texture->sample(*env.sampler, coord.xyz/coord.w, metal::level(level));
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(textureCube,
R"(
ANGLE_ALWAYS_INLINE auto ANGLE_textureCube(
thread ANGLE_TextureEnv<metal::texturecube<float>> &env,
metal::float3 const coord)
{
return env.texture->sample(*env.sampler, coord);
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(textureCubeBias,
R"(
ANGLE_ALWAYS_INLINE auto ANGLE_textureCube(
thread ANGLE_TextureEnv<metal::texturecube<float>> &env,
metal::float3 const coord,
float const bias)
{
return env.texture->sample(*env.sampler, coord, metal::bias(bias));
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(textureCubeGradEXT,
R"(
ANGLE_ALWAYS_INLINE auto ANGLE_textureCubeGradEXT(
thread ANGLE_TextureEnv<metal::texturecube<float>> &env,
metal::float3 const coord,
metal::float3 const dPdx,
metal::float3 const dPdy)
{
return env.texture->sample(*env.sampler, coord, metal::gradientcube(dPdx, dPdy));
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(textureCubeLod,
R"(
ANGLE_ALWAYS_INLINE auto ANGLE_textureCubeLod(
thread ANGLE_TextureEnv<metal::texturecube<float>> &env,
metal::float3 const coord,
float const level)
{
return env.texture->sample(*env.sampler, coord, metal::level(level));
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(textureCubeLodEXT,
R"(
#define ANGLE_textureCubeLodEXT ANGLE_textureCubeLod
)",
textureCubeLod())
PROGRAM_PRELUDE_DECLARE(textureGrad_2D,
R"(
template <typename T>
ANGLE_ALWAYS_INLINE auto ANGLE_textureGrad(
thread ANGLE_TextureEnv<metal::texture2d<T>> &env,
metal::float2 const coord,
metal::float2 const dPdx,
metal::float2 const dPdy)
{
return env.texture->sample(*env.sampler, coord, metal::gradient2d(dPdx, dPdy));
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(textureGrad_3D,
R"(
template <typename T>
ANGLE_ALWAYS_INLINE auto ANGLE_textureGrad(
thread ANGLE_TextureEnv<metal::texture3d<T>> &env,
metal::float3 const coord,
metal::float3 const dPdx,
metal::float3 const dPdy)
{
return env.texture->sample(*env.sampler, coord, metal::gradient3d(dPdx, dPdy));
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(textureGrad_Cube,
R"(
template <typename T>
ANGLE_ALWAYS_INLINE auto ANGLE_textureGrad(
thread ANGLE_TextureEnv<metal::texturecube<T>> &env,
metal::float3 const coord,
metal::float3 const dPdx,
metal::float3 const dPdy)
{
return env.texture->sample(*env.sampler, coord, metal::gradientcube(dPdx, dPdy));
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(textureGrad_2DArray,
R"(
template <typename T>
ANGLE_ALWAYS_INLINE auto ANGLE_textureGrad(
thread ANGLE_TextureEnv<metal::texture2d_array<T>> &env,
metal::float3 const coord,
metal::float2 const dPdx,
metal::float2 const dPdy)
{
return env.texture->sample(*env.sampler, coord.xy, uint32_t(metal::round(coord.z)), metal::gradient2d(dPdx, dPdy));
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(textureGrad_2DShadow,
R"(
ANGLE_ALWAYS_INLINE auto ANGLE_textureGrad(
thread ANGLE_TextureEnv<metal::depth2d<float>> &env,
metal::float3 const coord,
metal::float2 const dPdx,
metal::float2 const dPdy)
{
if (ANGLEUseSampleCompareGradient)
{
return env.texture->sample_compare(*env.sampler, coord.xy, coord.z, metal::gradient2d(dPdx, dPdy));
}
else
{
const float2 dims = float2(env.texture->get_width(0), env.texture->get_height(0));
const float lod = 0.5 * metal::log2(metal::max(metal::length_squared(dPdx * dims), metal::length_squared(dPdy * dims)));
return env.texture->sample_compare(*env.sampler, coord.xy, coord.z, metal::level(lod));
}
}
)",
functionConstants(),
textureEnv())
PROGRAM_PRELUDE_DECLARE(textureGrad_2DArrayShadow,
R"(
ANGLE_ALWAYS_INLINE auto ANGLE_textureGrad(
thread ANGLE_TextureEnv<metal::depth2d_array<float>> &env,
metal::float4 const coord,
metal::float2 const dPdx,
metal::float2 const dPdy)
{
if (ANGLEUseSampleCompareGradient)
{
return env.texture->sample_compare(*env.sampler, coord.xy, uint32_t(metal::round(coord.z)), coord.w, metal::gradient2d(dPdx, dPdy));
}
else
{
const float2 dims = float2(env.texture->get_width(0), env.texture->get_height(0));
const float lod = 0.5 * metal::log2(metal::max(metal::length_squared(dPdx * dims), metal::length_squared(dPdy * dims)));
return env.texture->sample_compare(*env.sampler, coord.xy, uint32_t(metal::round(coord.z)), coord.w, metal::level(lod));
}
}
)",
functionConstants(),
textureEnv())
PROGRAM_PRELUDE_DECLARE(textureGrad_CubeShadow,
R"(
ANGLE_ALWAYS_INLINE auto ANGLE_textureGrad(
thread ANGLE_TextureEnv<metal::depthcube<float>> &env,
metal::float4 const coord,
metal::float3 const dPdx,
metal::float3 const dPdy)
{
if (ANGLEUseSampleCompareGradient)
{
return env.texture->sample_compare(*env.sampler, coord.xyz, coord.w, metal::gradientcube(dPdx, dPdy));
}
else
{
const float3 coord_abs = metal::abs(coord.xyz);
const bool z_major = coord_abs.z >= metal::max(coord_abs.x, coord_abs.y);
const bool y_major = coord_abs.y >= metal::max(coord_abs.x, coord_abs.z);
const float3 Q = z_major ? coord.xyz : (y_major ? coord.xzy : coord.yzx);
const float3 dQdx = z_major ? dPdx : (y_major ? dPdx.xzy : dPdx.yzx);
const float3 dQdy = z_major ? dPdy : (y_major ? dPdy.xzy : dPdy.yzx);
const float4 d = (float4(dQdx.xy, dQdy.xy) - (Q.xy / Q.z).xyxy * float4(dQdx.zz, dQdy.zz)) / Q.z;
const float dim = float(env.texture->get_width(0));
const float lod = -1.0 + 0.5 * metal::log2(dim * dim * metal::max(metal::length_squared(d.xy), metal::length_squared(d.zw)));
return env.texture->sample_compare(*env.sampler, coord.xyz, coord.w, metal::level(lod));
}
}
)",
functionConstants(),
textureEnv())
PROGRAM_PRELUDE_DECLARE(textureGradOffset_2D,
R"(
template <typename T>
ANGLE_ALWAYS_INLINE auto ANGLE_textureGradOffset(
thread ANGLE_TextureEnv<metal::texture2d<T>> &env,
metal::float2 const coord,
metal::float2 const dPdx,
metal::float2 const dPdy,
int2 const offset)
{
return env.texture->sample(*env.sampler, coord, metal::gradient2d(dPdx, dPdy), offset);
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(textureGradOffset_3D,
R"(
template <typename T>
ANGLE_ALWAYS_INLINE auto ANGLE_textureGradOffset(
thread ANGLE_TextureEnv<metal::texture3d<T>> &env,
metal::float3 const coord,
metal::float3 const dPdx,
metal::float3 const dPdy,
int3 const offset)
{
return env.texture->sample(*env.sampler, coord, metal::gradient3d(dPdx, dPdy), offset);
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(textureGradOffset_2DArray,
R"(
template <typename T>
ANGLE_ALWAYS_INLINE auto ANGLE_textureGradOffset(
thread ANGLE_TextureEnv<metal::texture2d_array<T>> &env,
metal::float3 const coord,
metal::float2 const dPdx,
metal::float2 const dPdy,
metal::int2 const offset)
{
return env.texture->sample(*env.sampler, coord.xy, uint32_t(metal::round(coord.z)), metal::gradient2d(dPdx, dPdy), offset);
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(textureGradOffset_2DShadow,
R"(
ANGLE_ALWAYS_INLINE auto ANGLE_textureGradOffset(
thread ANGLE_TextureEnv<metal::depth2d<float>> &env,
metal::float3 const coord,
metal::float2 const dPdx,
metal::float2 const dPdy,
metal::int2 const offset)
{
if (ANGLEUseSampleCompareGradient)
{
return env.texture->sample_compare(*env.sampler, coord.xy, coord.z, metal::gradient2d(dPdx, dPdy), offset);
}
else
{
const float2 dims = float2(env.texture->get_width(0), env.texture->get_height(0));
const float lod = 0.5 * metal::log2(metal::max(metal::length_squared(dPdx * dims), metal::length_squared(dPdy * dims)));
return env.texture->sample_compare(*env.sampler, coord.xy, coord.z, metal::level(lod), offset);
}
}
)",
functionConstants(),
textureEnv())
PROGRAM_PRELUDE_DECLARE(textureGradOffset_2DArrayShadow,
R"(
ANGLE_ALWAYS_INLINE auto ANGLE_textureGradOffset(
thread ANGLE_TextureEnv<metal::depth2d_array<float>> &env,
metal::float4 const coord,
metal::float2 const dPdx,
metal::float2 const dPdy,
metal::int2 const offset)
{
if (ANGLEUseSampleCompareGradient)
{
return env.texture->sample_compare(*env.sampler, coord.xy, uint32_t(metal::round(coord.z)), coord.w, metal::gradient2d(dPdx, dPdy), offset);
}
else
{
const float2 dims = float2(env.texture->get_width(0), env.texture->get_height(0));
const float lod = 0.5 * metal::log2(metal::max(metal::length_squared(dPdx * dims), metal::length_squared(dPdy * dims)));
return env.texture->sample_compare(*env.sampler, coord.xy, uint32_t(metal::round(coord.z)), coord.w, metal::level(lod), offset);
}
}
)",
functionConstants(),
textureEnv())
PROGRAM_PRELUDE_DECLARE(textureLod_2D,
R"(
template <typename T>
ANGLE_ALWAYS_INLINE auto ANGLE_textureLod(
thread ANGLE_TextureEnv<metal::texture2d<T>> &env,
metal::float2 const coord,
float const level)
{
return env.texture->sample(*env.sampler, coord, metal::level(level));
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(textureLod_3D,
R"(
template <typename T>
ANGLE_ALWAYS_INLINE auto ANGLE_textureLod(
thread ANGLE_TextureEnv<metal::texture3d<T>> &env,
metal::float3 const coord,
float const level)
{
return env.texture->sample(*env.sampler, coord, metal::level(level));
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(textureLod_Cube,
R"(
template <typename T>
ANGLE_ALWAYS_INLINE auto ANGLE_textureLod(
thread ANGLE_TextureEnv<metal::texturecube<T>> &env,
metal::float3 const coord,
float const level)
{
return env.texture->sample(*env.sampler, coord, metal::level(level));
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(textureLod_2DShadow,
R"(
ANGLE_ALWAYS_INLINE auto ANGLE_textureLod(
thread ANGLE_TextureEnv<metal::depth2d<float>> &env,
metal::float3 const coord,
float const level)
{
return env.texture->sample_compare(*env.sampler, coord.xy, coord.z, metal::level(level));
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(textureLod_2DArray,
R"(
template <typename T>
ANGLE_ALWAYS_INLINE auto ANGLE_textureLod(
thread ANGLE_TextureEnv<metal::texture2d_array<T>> &env,
metal::float3 const coord,
float const level)
{
return env.texture->sample(*env.sampler, coord.xy, uint32_t(metal::round(coord.z)), metal::level(level));
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(textureLod_CubeShadow,
R"(
ANGLE_ALWAYS_INLINE auto ANGLE_textureLod(
thread ANGLE_TextureEnv<metal::depthcube<float>> &env,
metal::float4 const coord,
float const level)
{
return env.texture->sample_compare(*env.sampler, coord.xyz, coord.w, metal::level(level));
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(textureLod_2DArrayShadow,
R"(
ANGLE_ALWAYS_INLINE auto ANGLE_textureLod(
thread ANGLE_TextureEnv<metal::depth2d_array<float>> &env,
metal::float4 const coord,
float const level)
{
return env.texture->sample_compare(*env.sampler, coord.xy, uint32_t(metal::round(coord.z)), coord.w, metal::level(level));
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(textureLodOffset_2D,
R"(
template <typename T>
ANGLE_ALWAYS_INLINE auto ANGLE_textureLodOffset(
thread ANGLE_TextureEnv<metal::texture2d<T>> &env,
metal::float2 const coord,
float const level,
metal::int2 const offset)
{
return env.texture->sample(*env.sampler, coord, metal::level(level), offset);
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(textureLodOffset_3D,
R"(
template <typename T>
ANGLE_ALWAYS_INLINE auto ANGLE_textureLodOffset(
thread ANGLE_TextureEnv<metal::texture3d<T>> &env,
metal::float3 const coord,
float const level,
metal::int3 const offset)
{
return env.texture->sample(*env.sampler, coord, metal::level(level), offset);
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(textureLodOffset_2DShadow,
R"(
ANGLE_ALWAYS_INLINE auto ANGLE_textureLodOffset(
thread ANGLE_TextureEnv<metal::depth2d<float>> &env,
metal::float3 const coord,
float const level,
int2 const offset)
{
return env.texture->sample_compare(*env.sampler, coord.xy, coord.z, metal::level(level), offset);
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(textureLodOffset_2DArray,
R"(
template <typename T>
ANGLE_ALWAYS_INLINE auto ANGLE_textureLodOffset(
thread ANGLE_TextureEnv<metal::texture2d_array<T>> &env,
metal::float3 const coord,
float const level,
metal::int2 const offset)
{
return env.texture->sample(*env.sampler, coord.xy, uint32_t(metal::round(coord.z)), metal::level(level), offset);
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(textureLodOffset_2DArrayShadow,
R"(
ANGLE_ALWAYS_INLINE auto ANGLE_textureLodOffset(
thread ANGLE_TextureEnv<metal::depth2d_array<float>> &env,
metal::float4 const coord,
float const level,
metal::int2 const offset)
{
return env.texture->sample_compare(*env.sampler, coord.xy, uint32_t(metal::round(coord.z)), coord.w, metal::level(level), offset);
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(textureOffset_2D,
R"(
template <typename T>
ANGLE_ALWAYS_INLINE auto ANGLE_textureOffset(
thread ANGLE_TextureEnv<metal::texture2d<T>> &env,
metal::float2 const coord,
metal::int2 const offset)
{
return env.texture->sample(*env.sampler, coord, offset);
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(textureOffsetBias_2D,
R"(
template <typename T>
ANGLE_ALWAYS_INLINE auto ANGLE_textureOffset(
thread ANGLE_TextureEnv<metal::texture2d<T>> &env,
metal::float2 const coord,
metal::int2 const offset,
float const bias)
{
return env.texture->sample(*env.sampler, coord, metal::bias(bias), offset);
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(textureOffset_2DArray,
R"(
template <typename T>
ANGLE_ALWAYS_INLINE auto ANGLE_textureOffset(
thread ANGLE_TextureEnv<metal::texture2d_array<T>> &env,
metal::float3 const coord,
metal::int2 const offset)
{
return env.texture->sample(*env.sampler, coord.xy, uint32_t(metal::round(coord.z)), offset);
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(textureOffsetBias_2DArray,
R"(
template <typename T>
ANGLE_ALWAYS_INLINE auto ANGLE_textureOffset(
thread ANGLE_TextureEnv<metal::texture2d_array<T>> &env,
metal::float3 const coord,
metal::int2 const offset,
float const bias)
{
return env.texture->sample(*env.sampler, coord.xy, uint32_t(metal::round(coord.z)), metal::bias(bias), offset);
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(textureOffset_3D,
R"(
template <typename T>
ANGLE_ALWAYS_INLINE auto ANGLE_textureOffset(
thread ANGLE_TextureEnv<metal::texture3d<T>> &env,
metal::float3 const coord,
metal::int3 const offset)
{
return env.texture->sample(*env.sampler, coord, offset);
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(textureOffsetBias_3D,
R"(
template <typename T>
ANGLE_ALWAYS_INLINE auto ANGLE_textureOffset(
thread ANGLE_TextureEnv<metal::texture3d<T>> &env,
metal::float3 const coord,
metal::int3 const offset,
float const bias)
{
return env.texture->sample(*env.sampler, coord, metal::bias(bias), offset);
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(textureOffset_2DShadow,
R"(
ANGLE_ALWAYS_INLINE auto ANGLE_textureOffset(
thread ANGLE_TextureEnv<metal::depth2d<float>> &env,
metal::float3 const coord,
metal::int2 const offset)
{
return env.texture->sample_compare(*env.sampler, coord.xy, coord.z, offset);
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(textureOffsetBias_2DShadow,
R"(
ANGLE_ALWAYS_INLINE auto ANGLE_textureOffset(
thread ANGLE_TextureEnv<metal::depth2d<float>> &env,
metal::float3 const coord,
metal::int2 const offset,
float const bias)
{
return env.texture->sample_compare(*env.sampler, coord.xy, coord.z, metal::bias(bias), offset);
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(textureOffset_2DArrayShadow,
R"(
ANGLE_ALWAYS_INLINE auto ANGLE_textureOffset(
thread ANGLE_TextureEnv<metal::depth2d_array<float>> &env,
metal::float4 const coord,
metal::int2 const offset)
{
return env.texture->sample_compare(*env.sampler, coord.xy, uint32_t(metal::round(coord.z)), coord.w, offset);
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(textureOffsetBias_2DArrayShadow,
R"(
ANGLE_ALWAYS_INLINE auto ANGLE_textureOffset(
thread ANGLE_TextureEnv<metal::depth2d_array<float>> &env,
metal::float4 const coord,
metal::int2 const offset,
float const bias)
{
return env.texture->sample_compare(*env.sampler, coord.xy, uint32_t(metal::round(coord.z)), coord.w, metal::bias(bias), offset);
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(textureProj_2D_float3,
R"(
template <typename T>
ANGLE_ALWAYS_INLINE auto ANGLE_textureProj(
thread ANGLE_TextureEnv<metal::texture2d<T>> &env,
metal::float3 const coord)
{
return env.texture->sample(*env.sampler, coord.xy/coord.z);
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(textureProjBias_2D_float3,
R"(
template <typename T>
ANGLE_ALWAYS_INLINE auto ANGLE_textureProj(
thread ANGLE_TextureEnv<metal::texture2d<T>> &env,
metal::float3 const coord,
float const bias)
{
return env.texture->sample(*env.sampler, coord.xy/coord.z, metal::bias(bias));
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(textureProj_2D_float4,
R"(
template <typename T>
ANGLE_ALWAYS_INLINE auto ANGLE_textureProj(
thread ANGLE_TextureEnv<metal::texture2d<T>> &env,
metal::float4 const coord)
{
return env.texture->sample(*env.sampler, coord.xy/coord.w);
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(textureProjBias_2D_float4,
R"(
template <typename T>
ANGLE_ALWAYS_INLINE auto ANGLE_textureProj(
thread ANGLE_TextureEnv<metal::texture2d<T>> &env,
metal::float4 const coord,
float const bias)
{
return env.texture->sample(*env.sampler, coord.xy/coord.w, metal::bias(bias));
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(textureProj_3D,
R"(
template <typename T>
ANGLE_ALWAYS_INLINE auto ANGLE_textureProj(
thread ANGLE_TextureEnv<metal::texture3d<T>> &env,
metal::float4 const coord)
{
return env.texture->sample(*env.sampler, coord.xyz/coord.w);
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(textureProjBias_3D,
R"(
template <typename T>
ANGLE_ALWAYS_INLINE auto ANGLE_textureProj(
thread ANGLE_TextureEnv<metal::texture3d<T>> &env,
metal::float4 const coord,
float const bias)
{
return env.texture->sample(*env.sampler, coord.xyz/coord.w, metal::bias(bias));
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(textureProj_2DShadow,
R"(
ANGLE_ALWAYS_INLINE auto ANGLE_textureProj(
thread ANGLE_TextureEnv<metal::depth2d<float>> &env,
metal::float4 const coord)
{
return env.texture->sample_compare(*env.sampler, coord.xy/coord.w, coord.z/coord.w);
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(textureProjBias_2DShadow,
R"(
ANGLE_ALWAYS_INLINE auto ANGLE_textureProj(
thread ANGLE_TextureEnv<metal::depth2d<float>> &env,
metal::float4 const coord,
float const bias)
{
return env.texture->sample_compare(*env.sampler, coord.xy/coord.w, coord.z/coord.w, metal::bias(bias));
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(textureProjGrad_2D_float3,
R"(
template <typename T>
ANGLE_ALWAYS_INLINE auto ANGLE_textureProjGrad(
thread ANGLE_TextureEnv<metal::texture2d<T>> &env,
metal::float3 const coord,
metal::float2 const dPdx,
metal::float2 const dPdy)
{
return env.texture->sample(*env.sampler, coord.xy/coord.z, metal::gradient2d(dPdx, dPdy));
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(textureProjGrad_2D_float4,
R"(
template <typename T>
ANGLE_ALWAYS_INLINE auto ANGLE_textureProjGrad(
thread ANGLE_TextureEnv<metal::texture2d<T>> &env,
metal::float4 const coord,
metal::float2 const dPdx,
metal::float2 const dPdy)
{
return env.texture->sample(*env.sampler, coord.xy/coord.w, metal::gradient2d(dPdx, dPdy));
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(textureProjGrad_2DShadow,
R"(
ANGLE_ALWAYS_INLINE auto ANGLE_textureProjGrad(
thread ANGLE_TextureEnv<metal::depth2d<float>> &env,
metal::float4 const coord,
metal::float2 const dPdx,
metal::float2 const dPdy)
{
if (ANGLEUseSampleCompareGradient)
{
return env.texture->sample_compare(*env.sampler, coord.xy/coord.w, coord.z/coord.w, metal::gradient2d(dPdx, dPdy));
}
else
{
const float2 dims = float2(env.texture->get_width(0), env.texture->get_height(0));
const float lod = 0.5 * metal::log2(metal::max(metal::length_squared(dPdx * dims), metal::length_squared(dPdy * dims)));
return env.texture->sample_compare(*env.sampler, coord.xy/coord.w, coord.z/coord.w, metal::level(lod));
}
}
)",
functionConstants(),
textureEnv())
PROGRAM_PRELUDE_DECLARE(textureProjGrad_3D,
R"(
template <typename T>
ANGLE_ALWAYS_INLINE auto ANGLE_textureProjGrad(
thread ANGLE_TextureEnv<metal::texture3d<T>> &env,
metal::float4 const coord,
metal::float3 const dPdx,
metal::float3 const dPdy)
{
return env.texture->sample(*env.sampler, coord.xyz/coord.w, metal::gradient3d(dPdx, dPdy));
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(textureProjGradOffset_2D_float3,
R"(
template <typename T>
ANGLE_ALWAYS_INLINE auto ANGLE_textureProjGradOffset(
thread ANGLE_TextureEnv<metal::texture2d<T>> &env,
metal::float3 const coord,
metal::float2 const dPdx,
metal::float2 const dPdy,
int2 const offset)
{
return env.texture->sample(*env.sampler, coord.xy/coord.z, metal::gradient2d(dPdx, dPdy), offset);
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(textureProjGradOffset_2D_float4,
R"(
template <typename T>
ANGLE_ALWAYS_INLINE auto ANGLE_textureProjGradOffset(
thread ANGLE_TextureEnv<metal::texture2d<T>> &env,
metal::float4 const coord,
metal::float2 const dPdx,
metal::float2 const dPdy,
int2 const offset)
{
return env.texture->sample(*env.sampler, coord.xy/coord.w, metal::gradient2d(dPdx, dPdy), offset);
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(textureProjGradOffset_2DShadow,
R"(
ANGLE_ALWAYS_INLINE auto ANGLE_textureProjGradOffset(
thread ANGLE_TextureEnv<metal::depth2d<float>> &env,
metal::float4 const coord,
metal::float2 const dPdx,
metal::float2 const dPdy,
int2 const offset)
{
if (ANGLEUseSampleCompareGradient)
{
return env.texture->sample_compare(*env.sampler, coord.xy/coord.w, coord.z/coord.w, metal::gradient2d(dPdx, dPdy), offset);
}
else
{
const float2 dims = float2(env.texture->get_width(0), env.texture->get_height(0));
const float lod = 0.5 * metal::log2(metal::max(metal::length_squared(dPdx * dims), metal::length_squared(dPdy * dims)));
return env.texture->sample_compare(*env.sampler, coord.xy/coord.w, coord.z/coord.w, metal::level(lod), offset);
}
}
)",
functionConstants(),
textureEnv())
PROGRAM_PRELUDE_DECLARE(textureProjGradOffset_3D,
R"(
template <typename T>
ANGLE_ALWAYS_INLINE auto ANGLE_textureProjGradOffset(
thread ANGLE_TextureEnv<metal::texture3d<T>> &env,
metal::float4 const coord,
metal::float3 const dPdx,
metal::float3 const dPdy,
int3 const offset)
{
return env.texture->sample(*env.sampler, coord.xyz/coord.w, metal::gradient3d(dPdx, dPdy), offset);
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(textureProjLod_2D_float3,
R"(
template <typename T>
ANGLE_ALWAYS_INLINE auto ANGLE_textureProjLod(
thread ANGLE_TextureEnv<metal::texture2d<T>> &env,
metal::float3 const coord,
float const level)
{
return env.texture->sample(*env.sampler, coord.xy/coord.z, metal::level(level));
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(textureProjLod_2D_float4,
R"(
template <typename T>
ANGLE_ALWAYS_INLINE auto ANGLE_textureProjLod(
thread ANGLE_TextureEnv<metal::texture2d<T>> &env,
metal::float4 const coord,
float const level)
{
return env.texture->sample(*env.sampler, coord.xy/coord.w, metal::level(level));
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(textureProjLod_2DShadow,
R"(
ANGLE_ALWAYS_INLINE auto ANGLE_textureProjLod(
thread ANGLE_TextureEnv<metal::depth2d<float>> &env,
metal::float4 const coord,
float const level)
{
return env.texture->sample_compare(*env.sampler, coord.xy/coord.w, coord.z/coord.w, metal::level(level));
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(textureProjLod_3D,
R"(
template <typename T>
ANGLE_ALWAYS_INLINE auto ANGLE_textureProjLod(
thread ANGLE_TextureEnv<metal::texture3d<T>> &env,
metal::float4 const coord,
float const level)
{
return env.texture->sample(*env.sampler, coord.xyz/coord.w, metal::level(level));
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(textureProjLodOffset_2D_float3,
R"(
template <typename T>
ANGLE_ALWAYS_INLINE auto ANGLE_textureProjLodOffset(
thread ANGLE_TextureEnv<metal::texture2d<T>> &env,
metal::float3 const coord,
float const level,
int2 const offset)
{
return env.texture->sample(*env.sampler, coord.xy/coord.z, metal::level(level), offset);
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(textureProjLodOffset_2D_float4,
R"(
template <typename T>
ANGLE_ALWAYS_INLINE auto ANGLE_textureProjLodOffset(
thread ANGLE_TextureEnv<metal::texture2d<T>> &env,
metal::float4 const coord,
float const level,
int2 const offset)
{
return env.texture->sample(*env.sampler, coord.xy/coord.w, metal::level(level), offset);
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(textureProjLodOffset_2DShadow,
R"(
ANGLE_ALWAYS_INLINE auto ANGLE_textureProjLodOffset(
thread ANGLE_TextureEnv<metal::depth2d<float>> &env,
metal::float4 const coord,
float const level,
int2 const offset)
{
return env.texture->sample_compare(*env.sampler, coord.xy/coord.w, coord.z/coord.w, metal::level(level), offset);
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(textureProjLodOffset_3D,
R"(
template <typename T>
ANGLE_ALWAYS_INLINE auto ANGLE_textureProjLodOffset(
thread ANGLE_TextureEnv<metal::texture3d<T>> &env,
metal::float4 const coord,
float const level,
int3 const offset)
{
return env.texture->sample(*env.sampler, coord.xyz/coord.w, metal::level(level), offset);
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(textureProjOffset_2D_float3,
R"(
template <typename T>
ANGLE_ALWAYS_INLINE auto ANGLE_textureProjOffset(
thread ANGLE_TextureEnv<metal::texture2d<T>> &env,
metal::float3 const coord,
int2 const offset)
{
return env.texture->sample(*env.sampler, coord.xy/coord.z, offset);
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(textureProjOffsetBias_2D_float3,
R"(
template <typename T>
ANGLE_ALWAYS_INLINE auto ANGLE_textureProjOffset(
thread ANGLE_TextureEnv<metal::texture2d<T>> &env,
metal::float3 const coord,
int2 const offset,
float const bias)
{
return env.texture->sample(*env.sampler, coord.xy/coord.z, metal::bias(bias), offset);
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(textureProjOffset_2D_float4,
R"(
template <typename T>
ANGLE_ALWAYS_INLINE auto ANGLE_textureProjOffset(
thread ANGLE_TextureEnv<metal::texture2d<T>> &env,
metal::float4 const coord,
int2 const offset)
{
return env.texture->sample(*env.sampler, coord.xy/coord.w, offset);
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(textureProjOffsetBias_2D_float4,
R"(
template <typename T>
ANGLE_ALWAYS_INLINE auto ANGLE_textureProjOffset(
thread ANGLE_TextureEnv<metal::texture2d<T>> &env,
metal::float4 const coord,
int2 const offset,
float const bias)
{
return env.texture->sample(*env.sampler, coord.xy/coord.w, metal::bias(bias), offset);
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(textureProjOffset_3D,
R"(
template <typename T>
ANGLE_ALWAYS_INLINE auto ANGLE_textureProjOffset(
thread ANGLE_TextureEnv<metal::texture3d<T>> &env,
metal::float4 const coord,
int3 const offset)
{
return env.texture->sample(*env.sampler, coord.xyz/coord.w, offset);
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(textureProjOffsetBias_3D,
R"(
template <typename T>
ANGLE_ALWAYS_INLINE auto ANGLE_textureProjOffset(
thread ANGLE_TextureEnv<metal::texture3d<T>> &env,
metal::float4 const coord,
int3 const offset,
float const bias)
{
return env.texture->sample(*env.sampler, coord.xyz/coord.w, metal::bias(bias), offset);
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(textureProjOffset_2DShadow,
R"(
ANGLE_ALWAYS_INLINE auto ANGLE_textureProjOffset(
thread ANGLE_TextureEnv<metal::depth2d<float>> &env,
metal::float4 const coord,
int2 const offset)
{
return env.texture->sample_compare(*env.sampler, coord.xy/coord.w, coord.z/coord.w, offset);
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(textureProjOffsetBias_2DShadow,
R"(
ANGLE_ALWAYS_INLINE auto ANGLE_textureProjOffset(
thread ANGLE_TextureEnv<metal::depth2d<float>> &env,
metal::float4 const coord,
int2 const offset,
float const bias)
{
return env.texture->sample_compare(*env.sampler, coord.xy/coord.w, coord.z/coord.w, metal::bias(bias), offset);
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(textureSize_2D,
R"(
template <typename Texture>
ANGLE_ALWAYS_INLINE auto ANGLE_textureSize(
thread ANGLE_TextureEnv<Texture> &env,
int const level)
{
return int2(env.texture->get_width(uint32_t(level)), env.texture->get_height(uint32_t(level)));
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(textureSize_3D,
R"(
template <typename T>
ANGLE_ALWAYS_INLINE auto ANGLE_textureSize(
thread ANGLE_TextureEnv<metal::texture3d<T>> &env,
int const level)
{
return int3(env.texture->get_width(uint32_t(level)), env.texture->get_height(uint32_t(level)), env.texture->get_depth(uint32_t(level)));
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(textureSize_2DArray,
R"(
template <typename T>
ANGLE_ALWAYS_INLINE auto ANGLE_textureSize(
thread ANGLE_TextureEnv<metal::texture2d_array<T>> &env,
int const level)
{
return int3(env.texture->get_width(uint32_t(level)), env.texture->get_height(uint32_t(level)), env.texture->get_array_size());
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(textureSize_2DArrayShadow,
R"(
ANGLE_ALWAYS_INLINE auto ANGLE_textureSize(
thread ANGLE_TextureEnv<metal::depth2d_array<float>> &env,
int const level)
{
return int3(env.texture->get_width(uint32_t(level)), env.texture->get_height(uint32_t(level)), env.texture->get_array_size());
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(textureSize_2DMS,
R"(
template <typename T>
ANGLE_ALWAYS_INLINE auto ANGLE_textureSize(
thread ANGLE_TextureEnv<metal::texture2d_ms<T>> &env)
{
return int2(env.texture->get_width(), env.texture->get_height());
}
)",
textureEnv())
PROGRAM_PRELUDE_DECLARE(imageLoad, R"(
template <typename T, metal::access Access>
ANGLE_ALWAYS_INLINE auto ANGLE_imageLoad(
thread const metal::texture2d<T, Access> &texture,
metal::int2 coord)
{
return texture.read(uint2(coord));
}
)")
PROGRAM_PRELUDE_DECLARE(imageStore, R"(
template <typename T, metal::access Access>
ANGLE_ALWAYS_INLINE auto ANGLE_imageStore(
thread const metal::texture2d<T, Access> &texture,
metal::int2 coord,
metal::vec<T, 4> value)
{
return texture.write(value, uint2(coord));
}
)")
// TODO(anglebug.com/40096838): When using raster order groups and pixel local storage, which only
// accesses the pixel coordinate, we probably only need an execution barrier (mem_flags::mem_none).
PROGRAM_PRELUDE_DECLARE(memoryBarrierImage, R"(
ANGLE_ALWAYS_INLINE void ANGLE_memoryBarrierImage()
{
simdgroup_barrier(metal::mem_flags::mem_texture);
}
)")
PROGRAM_PRELUDE_DECLARE(interpolateAtCenter,
R"(
template <typename T, typename P>
ANGLE_ALWAYS_INLINE T ANGLE_interpolateAtCenter(
thread metal::interpolant<T, P> &interpolant)
{
return interpolant.interpolate_at_center();
}
)")
PROGRAM_PRELUDE_DECLARE(interpolateAtCentroid,
R"(
template <typename T, typename P>
ANGLE_ALWAYS_INLINE T ANGLE_interpolateAtCentroid(
thread metal::interpolant<T, P> &interpolant)
{
return interpolant.interpolate_at_centroid();
}
template <typename T>
ANGLE_ALWAYS_INLINE T ANGLE_interpolateAtCentroid(T value) { return value; }
)")
PROGRAM_PRELUDE_DECLARE(interpolateAtSample,
R"(
template <typename T, typename P>
ANGLE_ALWAYS_INLINE T ANGLE_interpolateAtSample(
thread metal::interpolant<T, P> &interpolant,
int const sample)
{
if (ANGLEMultisampledRendering)
{
return interpolant.interpolate_at_sample(static_cast<uint32_t>(sample));
}
else
{
return interpolant.interpolate_at_center();
}
}
template <typename T>
ANGLE_ALWAYS_INLINE T ANGLE_interpolateAtSample(T value, int) { return value; }
)")
PROGRAM_PRELUDE_DECLARE(interpolateAtOffset,
R"(
template <typename T, typename P>
ANGLE_ALWAYS_INLINE T ANGLE_interpolateAtOffset(
thread metal::interpolant<T, P> &interpolant,
float2 const offset)
{
return interpolant.interpolate_at_offset(metal::saturate(offset + 0.5f));
}
template <typename T>
ANGLE_ALWAYS_INLINE T ANGLE_interpolateAtOffset(T value, float2) { return value; }
)")
////////////////////////////////////////////////////////////////////////////////
// Returned Name is valid for as long as `buffer` is still alive.
// Returns false if no template args exist.
// Returns false if buffer is not large enough.
//
// Example:
// "foo<1,2>" --> "foo<>"
static std::pair<Name, bool> MaskTemplateArgs(const Name &name, size_t bufferSize, char *buffer)
{
const char *begin = name.rawName().data();
const char *end = strchr(begin, '<');
if (!end)
{
return {{}, false};
}
size_t n = end - begin;
if (n + 3 > bufferSize)
{
return {{}, false};
}
for (size_t i = 0; i < n; ++i)
{
buffer[i] = begin[i];
}
buffer[n + 0] = '<';
buffer[n + 1] = '>';
buffer[n + 2] = '\0';
return {Name(buffer, name.symbolType()), true};
}
ProgramPrelude::FuncToEmitter ProgramPrelude::BuildFuncToEmitter()
{
#define EMIT_METHOD(method) \
[](ProgramPrelude &pp, const TFunction &) -> void { return pp.method(); }
FuncToEmitter map;
auto put = [&](Name name, FuncEmitter emitter) {
FuncEmitter &dest = map[name];
ASSERT(!dest);
dest = emitter;
};
auto putAngle = [&](const char *nameStr, FuncEmitter emitter) {
Name name(nameStr, SymbolType::AngleInternal);
put(name, emitter);
};
auto putBuiltIn = [&](const char *nameStr, FuncEmitter emitter) {
Name name(nameStr, SymbolType::BuiltIn);
put(name, emitter);
};
putAngle("addressof", EMIT_METHOD(addressof));
putAngle("cast<>", EMIT_METHOD(castMatrix));
putAngle("elem_ref", EMIT_METHOD(vectorElemRef));
putAngle("flatten", EMIT_METHOD(flattenArray));
putAngle("inout", EMIT_METHOD(inout));
putAngle("out", EMIT_METHOD(out));
putAngle("swizzle_ref", EMIT_METHOD(swizzleRef));
putBuiltIn("texelFetch", [](ProgramPrelude &pp, const TFunction &func) {
switch (func.getParam(0)->getType().getBasicType())
{
case EbtSampler2D:
case EbtISampler2D:
case EbtUSampler2D:
return pp.texelFetch_2D();
case EbtSampler3D:
case EbtISampler3D:
case EbtUSampler3D:
return pp.texelFetch_3D();
case EbtSampler2DArray:
case EbtISampler2DArray:
case EbtUSampler2DArray:
return pp.texelFetch_2DArray();
case EbtSampler2DMS:
case EbtISampler2DMS:
case EbtUSampler2DMS:
return pp.texelFetch_2DMS();
default:
UNREACHABLE();
}
});
putBuiltIn("texelFetchOffset", [](ProgramPrelude &pp, const TFunction &func) {
switch (func.getParam(0)->getType().getBasicType())
{
case EbtSampler2D:
case EbtISampler2D:
case EbtUSampler2D:
return pp.texelFetchOffset_2D();
case EbtSampler3D:
case EbtISampler3D:
case EbtUSampler3D:
return pp.texelFetchOffset_3D();
case EbtSampler2DArray:
case EbtISampler2DArray:
case EbtUSampler2DArray:
return pp.texelFetchOffset_2DArray();
default:
UNREACHABLE();
}
});
putBuiltIn("texture", [](ProgramPrelude &pp, const TFunction &func) {
const bool bias = func.getParamCount() == 3;
switch (func.getParam(0)->getType().getBasicType())
{
case EbtSampler2D:
case EbtISampler2D:
case EbtUSampler2D:
return bias ? pp.textureBias_2D() : pp.texture_2D();
case EbtSampler3D:
case EbtISampler3D:
case EbtUSampler3D:
return bias ? pp.textureBias_3D() : pp.texture_3D();
case EbtSamplerCube:
case EbtISamplerCube:
case EbtUSamplerCube:
return bias ? pp.textureBias_Cube() : pp.texture_Cube();
case EbtSampler2DArray:
case EbtISampler2DArray:
case EbtUSampler2DArray:
return bias ? pp.textureBias_2DArray() : pp.texture_2DArray();
case EbtSampler2DShadow:
return bias ? pp.textureBias_2DShadow() : pp.texture_2DShadow();
case EbtSamplerCubeShadow:
return bias ? pp.textureBias_CubeShadow() : pp.texture_CubeShadow();
case EbtSampler2DArrayShadow:
return bias ? pp.textureBias_2DArrayShadow() : pp.texture_2DArrayShadow();
default:
UNREACHABLE();
}
});
putBuiltIn("texture2D", [](ProgramPrelude &pp, const TFunction &func) {
switch (func.getParamCount())
{
case 2:
return pp.texture2D();
case 3:
return pp.texture2DBias();
default:
UNREACHABLE();
}
});
putBuiltIn("texture2DGradEXT", EMIT_METHOD(texture2DGradEXT));
putBuiltIn("texture2DLod", EMIT_METHOD(texture2DLod));
putBuiltIn("texture2DLodEXT", EMIT_METHOD(texture2DLodEXT));
putBuiltIn("texture2DProj", [](ProgramPrelude &pp, const TFunction &func) {
switch (func.getParamCount())
{
case 2:
return pp.texture2DProj();
case 3:
return pp.texture2DProjBias();
default:
UNREACHABLE();
}
});
putBuiltIn("texture2DProjGradEXT", EMIT_METHOD(texture2DProjGradEXT));
putBuiltIn("texture2DProjLod", EMIT_METHOD(texture2DProjLod));
putBuiltIn("texture2DProjLodEXT", EMIT_METHOD(texture2DProjLodEXT));
putBuiltIn("texture3D", [](ProgramPrelude &pp, const TFunction &func) {
switch (func.getParamCount())
{
case 2:
return pp.texture3D();
case 3:
return pp.texture3DBias();
default:
UNREACHABLE();
}
});
putBuiltIn("texture3DLod", EMIT_METHOD(texture3DLod));
putBuiltIn("texture3DProj", [](ProgramPrelude &pp, const TFunction &func) {
switch (func.getParamCount())
{
case 2:
return pp.texture3DProj();
case 3:
return pp.texture3DProjBias();
default:
UNREACHABLE();
}
});
putBuiltIn("texture3DProjLod", EMIT_METHOD(texture3DProjLod));
putBuiltIn("textureCube", [](ProgramPrelude &pp, const TFunction &func) {
switch (func.getParamCount())
{
case 2:
return pp.textureCube();
case 3:
return pp.textureCubeBias();
default:
UNREACHABLE();
}
});
putBuiltIn("textureCubeGradEXT", EMIT_METHOD(textureCubeGradEXT));
putBuiltIn("textureCubeLod", EMIT_METHOD(textureCubeLod));
putBuiltIn("textureCubeLodEXT", EMIT_METHOD(textureCubeLodEXT));
putBuiltIn("textureGrad", [](ProgramPrelude &pp, const TFunction &func) {
switch (func.getParam(0)->getType().getBasicType())
{
case EbtSampler2D:
case EbtISampler2D:
case EbtUSampler2D:
return pp.textureGrad_2D();
case EbtSampler3D:
case EbtISampler3D:
case EbtUSampler3D:
return pp.textureGrad_3D();
case EbtSamplerCube:
case EbtISamplerCube:
case EbtUSamplerCube:
return pp.textureGrad_Cube();
case EbtSampler2DArray:
case EbtISampler2DArray:
case EbtUSampler2DArray:
return pp.textureGrad_2DArray();
case EbtSampler2DShadow:
return pp.textureGrad_2DShadow();
case EbtSamplerCubeShadow:
return pp.textureGrad_CubeShadow();
case EbtSampler2DArrayShadow:
return pp.textureGrad_2DArrayShadow();
default:
UNREACHABLE();
}
});
putBuiltIn("textureGradOffset", [](ProgramPrelude &pp, const TFunction &func) {
switch (func.getParam(0)->getType().getBasicType())
{
case EbtSampler2D:
case EbtISampler2D:
case EbtUSampler2D:
return pp.textureGradOffset_2D();
case EbtSampler3D:
case EbtISampler3D:
case EbtUSampler3D:
return pp.textureGradOffset_3D();
case EbtSampler2DArray:
case EbtISampler2DArray:
case EbtUSampler2DArray:
return pp.textureGradOffset_2DArray();
case EbtSampler2DShadow:
return pp.textureGradOffset_2DShadow();
case EbtSampler2DArrayShadow:
return pp.textureGradOffset_2DArrayShadow();
default:
UNREACHABLE();
}
});
putBuiltIn("textureLod", [](ProgramPrelude &pp, const TFunction &func) {
switch (func.getParam(0)->getType().getBasicType())
{
case EbtSampler2D:
case EbtISampler2D:
case EbtUSampler2D:
return pp.textureLod_2D();
case EbtSampler3D:
case EbtISampler3D:
case EbtUSampler3D:
return pp.textureLod_3D();
case EbtSamplerCube:
case EbtISamplerCube:
case EbtUSamplerCube:
return pp.textureLod_Cube();
case EbtSampler2DArray:
case EbtISampler2DArray:
case EbtUSampler2DArray:
return pp.textureLod_2DArray();
case EbtSampler2DShadow:
return pp.textureLod_2DShadow();
case EbtSamplerCubeShadow:
return pp.textureLod_CubeShadow();
case EbtSampler2DArrayShadow:
return pp.textureLod_2DArrayShadow();
default:
UNREACHABLE();
}
});
putBuiltIn("textureLodOffset", [](ProgramPrelude &pp, const TFunction &func) {
switch (func.getParam(0)->getType().getBasicType())
{
case EbtSampler2D:
case EbtISampler2D:
case EbtUSampler2D:
return pp.textureLodOffset_2D();
case EbtSampler3D:
case EbtISampler3D:
case EbtUSampler3D:
return pp.textureLodOffset_3D();
case EbtSampler2DArray:
case EbtISampler2DArray:
case EbtUSampler2DArray:
return pp.textureLodOffset_2DArray();
case EbtSampler2DShadow:
return pp.textureLodOffset_2DShadow();
case EbtSampler2DArrayShadow:
return pp.textureLodOffset_2DArrayShadow();
default:
UNREACHABLE();
}
});
putBuiltIn("textureOffset", [](ProgramPrelude &pp, const TFunction &func) {
const bool bias = func.getParamCount() == 4;
switch (func.getParam(0)->getType().getBasicType())
{
case EbtSampler2D:
case EbtISampler2D:
case EbtUSampler2D:
return bias ? pp.textureOffsetBias_2D() : pp.textureOffset_2D();
case EbtSampler3D:
case EbtISampler3D:
case EbtUSampler3D:
return bias ? pp.textureOffsetBias_3D() : pp.textureOffset_3D();
case EbtSampler2DArray:
case EbtISampler2DArray:
case EbtUSampler2DArray:
return bias ? pp.textureOffsetBias_2DArray() : pp.textureOffset_2DArray();
case EbtSampler2DShadow:
return bias ? pp.textureOffsetBias_2DShadow() : pp.textureOffset_2DShadow();
case EbtSampler2DArrayShadow:
return bias ? pp.textureOffsetBias_2DArrayShadow()
: pp.textureOffset_2DArrayShadow();
default:
UNREACHABLE();
}
});
putBuiltIn("textureProj", [](ProgramPrelude &pp, const TFunction &func) {
const bool bias = func.getParamCount() == 3;
switch (func.getParam(0)->getType().getBasicType())
{
case EbtSampler2D:
case EbtISampler2D:
case EbtUSampler2D:
return func.getParam(1)->getType().getNominalSize() == 4
? (bias ? pp.textureProjBias_2D_float4() : pp.textureProj_2D_float4())
: (bias ? pp.textureProjBias_2D_float3() : pp.textureProj_2D_float3());
case EbtSampler3D:
case EbtISampler3D:
case EbtUSampler3D:
return bias ? pp.textureProjBias_3D() : pp.textureProj_3D();
case EbtSampler2DShadow:
return bias ? pp.textureProjBias_2DShadow() : pp.textureProj_2DShadow();
default:
UNREACHABLE();
}
});
putBuiltIn("textureProjGrad", [](ProgramPrelude &pp, const TFunction &func) {
switch (func.getParam(0)->getType().getBasicType())
{
case EbtSampler2D:
case EbtISampler2D:
case EbtUSampler2D:
return func.getParam(1)->getType().getNominalSize() == 4
? pp.textureProjGrad_2D_float4()
: pp.textureProjGrad_2D_float3();
case EbtSampler3D:
case EbtISampler3D:
case EbtUSampler3D:
return pp.textureProjGrad_3D();
case EbtSampler2DShadow:
return pp.textureProjGrad_2DShadow();
default:
UNREACHABLE();
}
});
putBuiltIn("textureProjGradOffset", [](ProgramPrelude &pp, const TFunction &func) {
switch (func.getParam(0)->getType().getBasicType())
{
case EbtSampler2D:
case EbtISampler2D:
case EbtUSampler2D:
return func.getParam(1)->getType().getNominalSize() == 4
? pp.textureProjGradOffset_2D_float4()
: pp.textureProjGradOffset_2D_float3();
case EbtSampler3D:
case EbtISampler3D:
case EbtUSampler3D:
return pp.textureProjGradOffset_3D();
case EbtSampler2DShadow:
return pp.textureProjGradOffset_2DShadow();
default:
UNREACHABLE();
}
});
putBuiltIn("textureProjLod", [](ProgramPrelude &pp, const TFunction &func) {
switch (func.getParam(0)->getType().getBasicType())
{
case EbtSampler2D:
case EbtISampler2D:
case EbtUSampler2D:
return func.getParam(1)->getType().getNominalSize() == 4
? pp.textureProjLod_2D_float4()
: pp.textureProjLod_2D_float3();
case EbtSampler3D:
case EbtISampler3D:
case EbtUSampler3D:
return pp.textureProjLod_3D();
case EbtSampler2DShadow:
return pp.textureProjLod_2DShadow();
default:
UNREACHABLE();
}
});
putBuiltIn("textureProjLodOffset", [](ProgramPrelude &pp, const TFunction &func) {
switch (func.getParam(0)->getType().getBasicType())
{
case EbtSampler2D:
case EbtISampler2D:
case EbtUSampler2D:
return func.getParam(1)->getType().getNominalSize() == 4
? pp.textureProjLodOffset_2D_float4()
: pp.textureProjLodOffset_2D_float3();
case EbtSampler3D:
case EbtISampler3D:
case EbtUSampler3D:
return pp.textureProjLodOffset_3D();
case EbtSampler2DShadow:
return pp.textureProjLodOffset_2DShadow();
default:
UNREACHABLE();
}
});
putBuiltIn("textureProjOffset", [](ProgramPrelude &pp, const TFunction &func) {
const bool bias = func.getParamCount() == 4;
switch (func.getParam(0)->getType().getBasicType())
{
case EbtSampler2D:
case EbtISampler2D:
case EbtUSampler2D:
return func.getParam(1)->getType().getNominalSize() == 4
? (bias ? pp.textureProjOffsetBias_2D_float4()
: pp.textureProjOffset_2D_float4())
: (bias ? pp.textureProjOffsetBias_2D_float3()
: pp.textureProjOffset_2D_float3());
case EbtSampler3D:
case EbtISampler3D:
case EbtUSampler3D:
return bias ? pp.textureProjOffsetBias_3D() : pp.textureProjOffset_3D();
case EbtSampler2DShadow:
return bias ? pp.textureProjOffsetBias_2DShadow() : pp.textureProjOffset_2DShadow();
default:
UNREACHABLE();
}
});
putBuiltIn("textureSize", [](ProgramPrelude &pp, const TFunction &func) {
switch (func.getParam(0)->getType().getBasicType())
{
case EbtSampler3D:
case EbtISampler3D:
case EbtUSampler3D:
return pp.textureSize_3D();
case EbtSampler2DArray:
case EbtISampler2DArray:
case EbtUSampler2DArray:
return pp.textureSize_2DArray();
case EbtSampler2DArrayShadow:
return pp.textureSize_2DArrayShadow();
case EbtSampler2DMS:
case EbtISampler2DMS:
case EbtUSampler2DMS:
return pp.textureSize_2DMS();
default:
// Same wrapper for 2D, 2D Shadow, Cube, and Cube Shadow
return pp.textureSize_2D();
}
});
putBuiltIn("imageLoad", EMIT_METHOD(imageLoad));
putBuiltIn("imageStore", EMIT_METHOD(imageStore));
putBuiltIn("memoryBarrierImage", EMIT_METHOD(memoryBarrierImage));
putBuiltIn("interpolateAtCenter", EMIT_METHOD(interpolateAtCenter));
putBuiltIn("interpolateAtCentroid", EMIT_METHOD(interpolateAtCentroid));
putBuiltIn("interpolateAtSample", EMIT_METHOD(interpolateAtSample));
putBuiltIn("interpolateAtOffset", EMIT_METHOD(interpolateAtOffset));
return map;
#undef EMIT_METHOD
}
void ProgramPrelude::visitOperator(TOperator op, const TFunction *func, const TType *argType0)
{
visitOperator(op, func, argType0, nullptr, nullptr);
}
void ProgramPrelude::visitOperator(TOperator op,
const TFunction *func,
const TType *argType0,
const TType *argType1)
{
visitOperator(op, func, argType0, argType1, nullptr);
}
void ProgramPrelude::visitOperator(TOperator op,
const TFunction *func,
const TType *argType0,
const TType *argType1,
const TType *argType2)
{
switch (op)
{
case TOperator::EOpRadians:
radians();
break;
case TOperator::EOpDegrees:
degrees();
break;
case TOperator::EOpMod:
mod();
break;
case TOperator::EOpRefract:
if (argType0->isScalar())
{
refractScalar();
}
break;
case TOperator::EOpDistance:
if (argType0->isScalar())
{
distanceScalar();
}
break;
case TOperator::EOpLength:
case TOperator::EOpDot:
case TOperator::EOpNormalize:
break;
case TOperator::EOpFaceforward:
if (argType0->isScalar())
{
faceforwardScalar();
}
break;
case TOperator::EOpReflect:
if (argType0->isScalar())
{
reflectScalar();
}
break;
case TOperator::EOpSin:
case TOperator::EOpCos:
case TOperator::EOpTan:
case TOperator::EOpAsin:
case TOperator::EOpAcos:
case TOperator::EOpAtan:
case TOperator::EOpSinh:
case TOperator::EOpCosh:
case TOperator::EOpTanh:
case TOperator::EOpAsinh:
case TOperator::EOpAcosh:
case TOperator::EOpAtanh:
case TOperator::EOpAbs:
case TOperator::EOpFma:
case TOperator::EOpPow:
case TOperator::EOpExp:
case TOperator::EOpExp2:
case TOperator::EOpLog:
case TOperator::EOpLog2:
case TOperator::EOpSqrt:
case TOperator::EOpFloor:
case TOperator::EOpTrunc:
case TOperator::EOpCeil:
case TOperator::EOpFract:
case TOperator::EOpRound:
case TOperator::EOpRoundEven:
case TOperator::EOpSaturate:
case TOperator::EOpModf:
case TOperator::EOpLdexp:
case TOperator::EOpFrexp:
case TOperator::EOpInversesqrt:
break;
case TOperator::EOpEqual:
if (argType0->isVector() && argType1->isVector())
{
equalVector();
}
// Even if Arg0 is a vector or matrix, it could also be an array.
if (argType0->isArray() && argType1->isArray())
{
equalArray();
}
if (argType0->getStruct() && argType1->getStruct() && argType0->isArray() &&
argType1->isArray())
{
equalStructArray();
}
if (argType0->isMatrix() && argType1->isMatrix())
{
equalMatrix();
}
break;
case TOperator::EOpNotEqual:
if (argType0->isVector() && argType1->isVector())
{
notEqualVector();
}
else if (argType0->getStruct() && argType1->getStruct())
{
notEqualStruct();
}
// Same as above.
if (argType0->isArray() && argType1->isArray())
{
notEqualArray();
}
if (argType0->getStruct() && argType1->getStruct() && argType0->isArray() &&
argType1->isArray())
{
notEqualStructArray();
}
if (argType0->isMatrix() && argType1->isMatrix())
{
notEqualMatrix();
}
break;
case TOperator::EOpCross:
break;
case TOperator::EOpSign:
if (argType0->getBasicType() == TBasicType::EbtInt)
{
signInt();
}
break;
case TOperator::EOpClamp:
case TOperator::EOpMin:
case TOperator::EOpMax:
case TOperator::EOpStep:
case TOperator::EOpSmoothstep:
break;
case TOperator::EOpMix:
if (argType2->getBasicType() == TBasicType::EbtBool)
{
mixBool();
}
break;
case TOperator::EOpAll:
case TOperator::EOpAny:
case TOperator::EOpIsnan:
case TOperator::EOpIsinf:
case TOperator::EOpDFdx:
case TOperator::EOpDFdy:
case TOperator::EOpFwidth:
case TOperator::EOpTranspose:
case TOperator::EOpDeterminant:
break;
case TOperator::EOpAdd:
if (argType0->isMatrix() && argType1->isScalar())
{
addMatrixScalar();
}
if (argType0->isScalar() && argType1->isMatrix())
{
addScalarMatrix();
}
break;
case TOperator::EOpAddAssign:
if (argType0->isMatrix() && argType1->isScalar())
{
addMatrixScalarAssign();
}
break;
case TOperator::EOpSub:
if (argType0->isMatrix() && argType1->isScalar())
{
subMatrixScalar();
}
if (argType0->isScalar() && argType1->isMatrix())
{
subScalarMatrix();
}
break;
case TOperator::EOpSubAssign:
if (argType0->isMatrix() && argType1->isScalar())
{
subMatrixScalarAssign();
}
break;
case TOperator::EOpDiv:
if (argType1->isMatrix())
{
if (argType0->isMatrix())
{
componentWiseDivide();
}
else if (argType0->isScalar())
{
divScalarMatrix();
}
}
break;
case TOperator::EOpDivAssign:
if (argType0->isMatrix() && argType1->isMatrix())
{
componentWiseDivideAssign();
}
break;
case TOperator::EOpMatrixCompMult:
if (argType0->isMatrix() && argType1->isMatrix())
{
componentWiseMultiply();
}
break;
case TOperator::EOpOuterProduct:
outerProduct();
break;
case TOperator::EOpInverse:
switch (argType0->getCols())
{
case 2:
inverse2();
break;
case 3:
inverse3();
break;
case 4:
inverse4();
break;
default:
UNREACHABLE();
}
break;
case TOperator::EOpMatrixTimesMatrixAssign:
matmulAssign();
break;
case TOperator::EOpPreIncrement:
if (argType0->isMatrix())
{
preIncrementMatrix();
}
break;
case TOperator::EOpPostIncrement:
if (argType0->isMatrix())
{
postIncrementMatrix();
}
break;
case TOperator::EOpPreDecrement:
if (argType0->isMatrix())
{
preDecrementMatrix();
}
break;
case TOperator::EOpPostDecrement:
if (argType0->isMatrix())
{
postDecrementMatrix();
}
break;
case TOperator::EOpNegative:
if (argType0->isMatrix())
{
negateMatrix();
}
break;
case TOperator::EOpComma:
case TOperator::EOpAssign:
case TOperator::EOpInitialize:
case TOperator::EOpMulAssign:
case TOperator::EOpIModAssign:
case TOperator::EOpBitShiftLeftAssign:
case TOperator::EOpBitShiftRightAssign:
case TOperator::EOpBitwiseAndAssign:
case TOperator::EOpBitwiseXorAssign:
case TOperator::EOpBitwiseOrAssign:
case TOperator::EOpMul:
case TOperator::EOpIMod:
case TOperator::EOpBitShiftLeft:
case TOperator::EOpBitShiftRight:
case TOperator::EOpBitwiseAnd:
case TOperator::EOpBitwiseXor:
case TOperator::EOpBitwiseOr:
case TOperator::EOpLessThan:
case TOperator::EOpGreaterThan:
case TOperator::EOpLessThanEqual:
case TOperator::EOpGreaterThanEqual:
case TOperator::EOpLessThanComponentWise:
case TOperator::EOpLessThanEqualComponentWise:
case TOperator::EOpGreaterThanEqualComponentWise:
case TOperator::EOpGreaterThanComponentWise:
case TOperator::EOpLogicalOr:
case TOperator::EOpLogicalXor:
case TOperator::EOpLogicalAnd:
case TOperator::EOpPositive:
case TOperator::EOpLogicalNot:
case TOperator::EOpNotComponentWise:
case TOperator::EOpBitwiseNot:
case TOperator::EOpVectorTimesScalarAssign:
case TOperator::EOpVectorTimesMatrixAssign:
case TOperator::EOpMatrixTimesScalarAssign:
case TOperator::EOpVectorTimesScalar:
case TOperator::EOpVectorTimesMatrix:
case TOperator::EOpMatrixTimesVector:
case TOperator::EOpMatrixTimesScalar:
case TOperator::EOpMatrixTimesMatrix:
case TOperator::EOpReturn:
case TOperator::EOpBreak:
case TOperator::EOpContinue:
case TOperator::EOpEqualComponentWise:
case TOperator::EOpNotEqualComponentWise:
case TOperator::EOpIndexDirect:
case TOperator::EOpIndexIndirect:
case TOperator::EOpIndexDirectStruct:
case TOperator::EOpIndexDirectInterfaceBlock:
case TOperator::EOpFloatBitsToInt:
case TOperator::EOpIntBitsToFloat:
case TOperator::EOpUintBitsToFloat:
case TOperator::EOpFloatBitsToUint:
case TOperator::EOpNull:
case TOperator::EOpKill:
case TOperator::EOpPackUnorm2x16:
case TOperator::EOpPackSnorm2x16:
case TOperator::EOpPackUnorm4x8:
case TOperator::EOpPackSnorm4x8:
case TOperator::EOpUnpackSnorm2x16:
case TOperator::EOpUnpackUnorm2x16:
case TOperator::EOpUnpackUnorm4x8:
case TOperator::EOpUnpackSnorm4x8:
break;
case TOperator::EOpPackHalf2x16:
pack_half_2x16();
break;
case TOperator::EOpUnpackHalf2x16:
unpack_half_2x16();
break;
case TOperator::EOpBitfieldExtract:
case TOperator::EOpBitfieldInsert:
case TOperator::EOpBitfieldReverse:
case TOperator::EOpBitCount:
case TOperator::EOpFindLSB:
case TOperator::EOpFindMSB:
case TOperator::EOpUaddCarry:
case TOperator::EOpUsubBorrow:
case TOperator::EOpUmulExtended:
case TOperator::EOpImulExtended:
case TOperator::EOpBarrier:
case TOperator::EOpMemoryBarrier:
case TOperator::EOpMemoryBarrierAtomicCounter:
case TOperator::EOpMemoryBarrierBuffer:
case TOperator::EOpMemoryBarrierShared:
case TOperator::EOpGroupMemoryBarrier:
case TOperator::EOpAtomicAdd:
case TOperator::EOpAtomicMin:
case TOperator::EOpAtomicMax:
case TOperator::EOpAtomicAnd:
case TOperator::EOpAtomicOr:
case TOperator::EOpAtomicXor:
case TOperator::EOpAtomicExchange:
case TOperator::EOpAtomicCompSwap:
case TOperator::EOpEmitVertex:
case TOperator::EOpEndPrimitive:
case TOperator::EOpArrayLength:
UNIMPLEMENTED();
break;
case TOperator::EOpConstruct:
ASSERT(!func);
break;
case TOperator::EOpCallFunctionInAST:
case TOperator::EOpCallInternalRawFunction:
default:
ASSERT(func);
if (mHandled.insert(func).second)
{
const Name name(*func);
const auto end = mFuncToEmitter.end();
auto iter = mFuncToEmitter.find(name);
if (iter == end)
{
char buffer[32];
auto mask = MaskTemplateArgs(name, sizeof(buffer), buffer);
if (mask.second)
{
iter = mFuncToEmitter.find(mask.first);
}
}
if (iter != end)
{
const auto &emitter = iter->second;
emitter(*this, *func);
}
}
break;
}
}
void ProgramPrelude::visitVariable(const Name &name, const TType &type)
{
if (const TStructure *s = type.getStruct())
{
const Name typeName(*s);
if (typeName.beginsWith(Name("TextureEnv<")))
{
textureEnv();
}
}
else
{
if (name.rawName() == sh::mtl::kRasterizerDiscardEnabledConstName ||
name.rawName() == sh::mtl::kDepthWriteEnabledConstName ||
name.rawName() == sh::mtl::kEmulateAlphaToCoverageConstName)
{
functionConstants();
}
}
}
void ProgramPrelude::visitVariable(const TVariable &var)
{
if (mHandled.insert(&var).second)
{
visitVariable(Name(var), var.getType());
}
}
void ProgramPrelude::visitStructure(const TStructure &s)
{
if (mHandled.insert(&s).second)
{
for (const TField *field : s.fields())
{
const TType &type = *field->type();
visitVariable(Name(*field), type);
}
}
}
bool ProgramPrelude::visitBinary(Visit visit, TIntermBinary *node)
{
const TType &leftType = node->getLeft()->getType();
const TType &rightType = node->getRight()->getType();
visitOperator(node->getOp(), nullptr, &leftType, &rightType);
return true;
}
bool ProgramPrelude::visitUnary(Visit visit, TIntermUnary *node)
{
const TType &argType = node->getOperand()->getType();
visitOperator(node->getOp(), nullptr, &argType);
return true;
}
bool ProgramPrelude::visitAggregate(Visit visit, TIntermAggregate *node)
{
const size_t argCount = node->getChildCount();
auto getArgType = [node, argCount](size_t index) -> const TType & {
ASSERT(index < argCount);
TIntermTyped *arg = node->getChildNode(index)->getAsTyped();
ASSERT(arg);
return arg->getType();
};
const TFunction *func = node->getFunction();
switch (node->getChildCount())
{
case 0:
{
visitOperator(node->getOp(), func, nullptr);
}
break;
case 1:
{
const TType &argType0 = getArgType(0);
visitOperator(node->getOp(), func, &argType0);
}
break;
case 2:
{
const TType &argType0 = getArgType(0);
const TType &argType1 = getArgType(1);
visitOperator(node->getOp(), func, &argType0, &argType1);
}
break;
case 3:
{
const TType &argType0 = getArgType(0);
const TType &argType1 = getArgType(1);
const TType &argType2 = getArgType(2);
visitOperator(node->getOp(), func, &argType0, &argType1, &argType2);
}
break;
default:
{
const TType &argType0 = getArgType(0);
const TType &argType1 = getArgType(1);
visitOperator(node->getOp(), func, &argType0, &argType1);
}
break;
}
return true;
}
bool ProgramPrelude::visitDeclaration(Visit, TIntermDeclaration *node)
{
Declaration decl = ViewDeclaration(*node);
const TType &type = decl.symbol.getType();
if (type.isStructSpecifier())
{
const TStructure *s = type.getStruct();
ASSERT(s);
visitStructure(*s);
}
return true;
}
void ProgramPrelude::visitSymbol(TIntermSymbol *node)
{
visitVariable(node->variable());
}
bool sh::EmitProgramPrelude(TIntermBlock &root, TInfoSinkBase &out, const ProgramPreludeConfig &ppc)
{
ProgramPrelude programPrelude(out, ppc);
root.traverse(&programPrelude);
return true;
}