kmx git

//
// 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 <map>

#include "common/system_utils.h"
#include "compiler/translator/BaseTypes.h"
#include "compiler/translator/ImmutableStringBuilder.h"
#include "compiler/translator/Name.h"
#include "compiler/translator/OutputTree.h"
#include "compiler/translator/SymbolTable.h"
#include "compiler/translator/msl/AstHelpers.h"
#include "compiler/translator/msl/DebugSink.h"
#include "compiler/translator/msl/EmitMetal.h"
#include "compiler/translator/msl/Layout.h"
#include "compiler/translator/msl/ProgramPrelude.h"
#include "compiler/translator/msl/RewritePipelines.h"
#include "compiler/translator/msl/TranslatorMSL.h"
#include "compiler/translator/tree_util/IntermTraverse.h"

using namespace sh;

////////////////////////////////////////////////////////////////////////////////

#if defined(ANGLE_ENABLE_ASSERTS)
using Sink = DebugSink;
#else
using Sink = TInfoSinkBase;
#endif

////////////////////////////////////////////////////////////////////////////////

namespace
{

struct VarDecl
{
    explicit VarDecl(const TVariable &var) : mVariable(&var), mIsField(false) {}
    explicit VarDecl(const TField &field) : mField(&field), mIsField(true) {}

    ANGLE_INLINE const TVariable &variable() const
    {
        ASSERT(isVariable());
        return *mVariable;
    }

    ANGLE_INLINE const TField &field() const
    {
        ASSERT(isField());
        return *mField;
    }

    ANGLE_INLINE bool isVariable() const { return !mIsField; }

    ANGLE_INLINE bool isField() const { return mIsField; }

    const TType &type() const { return isField() ? *field().type() : variable().getType(); }

    SymbolType symbolType() const
    {
        return isField() ? field().symbolType() : variable().symbolType();
    }

  private:
    union
    {
        const TVariable *mVariable;
        const TField *mField;
    };
    bool mIsField;
};

class GenMetalTraverser : public TIntermTraverser
{
  public:
    ~GenMetalTraverser() override;

    GenMetalTraverser(const TCompiler &compiler,
                      Sink &out,
                      IdGen &idGen,
                      const PipelineStructs &pipelineStructs,
                      SymbolEnv &symbolEnv,
                      const ShCompileOptions &compileOptions);

    void visitSymbol(TIntermSymbol *) override;
    void visitConstantUnion(TIntermConstantUnion *) override;
    bool visitSwizzle(Visit, TIntermSwizzle *) override;
    bool visitBinary(Visit, TIntermBinary *) override;
    bool visitUnary(Visit, TIntermUnary *) override;
    bool visitTernary(Visit, TIntermTernary *) override;
    bool visitIfElse(Visit, TIntermIfElse *) override;
    bool visitSwitch(Visit, TIntermSwitch *) override;
    bool visitCase(Visit, TIntermCase *) override;
    void visitFunctionPrototype(TIntermFunctionPrototype *) override;
    bool visitFunctionDefinition(Visit, TIntermFunctionDefinition *) override;
    bool visitAggregate(Visit, TIntermAggregate *) override;
    bool visitBlock(Visit, TIntermBlock *) override;
    bool visitGlobalQualifierDeclaration(Visit, TIntermGlobalQualifierDeclaration *) override;
    bool visitDeclaration(Visit, TIntermDeclaration *) override;
    bool visitLoop(Visit, TIntermLoop *) override;
    bool visitForLoop(TIntermLoop *);
    bool visitWhileLoop(TIntermLoop *);
    bool visitDoWhileLoop(TIntermLoop *);
    bool visitBranch(Visit, TIntermBranch *) override;

  private:
    using FuncToName = std::map<ImmutableString, Name>;
    static FuncToName BuildFuncToName();

    struct EmitVariableDeclarationConfig
    {
        bool isParameter                = false;
        bool isMainParameter            = false;
        bool emitPostQualifier          = false;
        bool isPacked                   = false;
        bool disableStructSpecifier     = false;
        bool isUBO                      = false;
        const AddressSpace *isPointer   = nullptr;
        const AddressSpace *isReference = nullptr;
    };

    struct EmitTypeConfig
    {
        const EmitVariableDeclarationConfig *evdConfig = nullptr;
    };

    void emitIndentation();
    void emitOpeningPointerParen();
    void emitClosingPointerParen();
    void emitFunctionSignature(const TFunction &func);
    void emitFunctionReturn(const TFunction &func);
    void emitFunctionParameter(const TFunction &func, const TVariable &param);

    void emitNameOf(const TField &object);
    void emitNameOf(const TSymbol &object);
    void emitNameOf(const VarDecl &object);

    void emitBareTypeName(const TType &type, const EmitTypeConfig &etConfig);
    void emitType(const TType &type, const EmitTypeConfig &etConfig);
    void emitPostQualifier(const EmitVariableDeclarationConfig &evdConfig,
                           const VarDecl &decl,
                           const TQualifier qualifier);

    void emitLoopBody(TIntermBlock *bodyNode);

    struct FieldAnnotationIndices
    {
        size_t attribute = 0;
        size_t color     = 0;
    };

    void emitFieldDeclaration(const TField &field,
                              const TStructure &parent,
                              FieldAnnotationIndices &annotationIndices);
    void emitAttributeDeclaration(const TField &field, FieldAnnotationIndices &annotationIndices);
    void emitUniformBufferDeclaration(const TField &field,
                                      FieldAnnotationIndices &annotationIndices);
    void emitStructDeclaration(const TType &type);
    void emitOrdinaryVariableDeclaration(const VarDecl &decl,
                                         const EmitVariableDeclarationConfig &evdConfig);
    void emitVariableDeclaration(const VarDecl &decl,
                                 const EmitVariableDeclarationConfig &evdConfig);

    void emitOpenBrace();
    void emitCloseBrace();

    void groupedTraverse(TIntermNode &node);

    const TField &getDirectField(const TFieldListCollection &fieldsNode,
                                 const TConstantUnion &index);
    const TField &getDirectField(const TIntermTyped &fieldsNode, TIntermTyped &indexNode);

    const TConstantUnion *emitConstantUnionArray(const TConstantUnion *const constUnion,
                                                 const size_t size);

    const TConstantUnion *emitConstantUnion(const TType &type, const TConstantUnion *constUnion);

    void emitSingleConstant(const TConstantUnion *const constUnion);

  private:
    Sink &mOut;
    const TCompiler &mCompiler;
    const PipelineStructs &mPipelineStructs;
    SymbolEnv &mSymbolEnv;
    IdGen &mIdGen;
    int mIndentLevel                  = -1;
    int mLastIndentationPos           = -1;
    int mOpenPointerParenCount        = 0;
    bool mParentIsSwitch              = false;
    bool isTraversingVertexMain       = false;
    bool mTemporarilyDisableSemicolon = false;
    std::unordered_map<const TSymbol *, Name> mRenamedSymbols;
    const FuncToName mFuncToName           = BuildFuncToName();
    size_t mMainTextureIndex               = 0;
    size_t mMainSamplerIndex               = 0;
    size_t mMainUniformBufferIndex         = 0;
    size_t mDriverUniformsBindingIndex     = 0;
    size_t mUBOArgumentBufferBindingIndex  = 0;
    bool mRasterOrderGroupsSupported       = false;
    bool mInjectAsmStatementIntoLoopBodies = false;
};
}  // anonymous namespace

GenMetalTraverser::~GenMetalTraverser()
{
    ASSERT(mIndentLevel == -1);
    ASSERT(!mParentIsSwitch);
    ASSERT(mOpenPointerParenCount == 0);
}

GenMetalTraverser::GenMetalTraverser(const TCompiler &compiler,
                                     Sink &out,
                                     IdGen &idGen,
                                     const PipelineStructs &pipelineStructs,
                                     SymbolEnv &symbolEnv,
                                     const ShCompileOptions &compileOptions)
    : TIntermTraverser(true, false, false),
      mOut(out),
      mCompiler(compiler),
      mPipelineStructs(pipelineStructs),
      mSymbolEnv(symbolEnv),
      mIdGen(idGen),
      mMainUniformBufferIndex(compileOptions.metal.defaultUniformsBindingIndex),
      mDriverUniformsBindingIndex(compileOptions.metal.driverUniformsBindingIndex),
      mUBOArgumentBufferBindingIndex(compileOptions.metal.UBOArgumentBufferBindingIndex),
      mRasterOrderGroupsSupported(compileOptions.pls.fragmentSyncType ==
                                  ShFragmentSynchronizationType::RasterOrderGroups_Metal),
      mInjectAsmStatementIntoLoopBodies(compileOptions.metal.injectAsmStatementIntoLoopBodies)
{}

void GenMetalTraverser::emitIndentation()
{
    ASSERT(mIndentLevel >= 0);

    if (mLastIndentationPos == mOut.size())
    {
        return;  // Line is already indented.
    }

    for (int i = 0; i < mIndentLevel; ++i)
    {
        mOut << "  ";
    }

    mLastIndentationPos = mOut.size();
}

void GenMetalTraverser::emitOpeningPointerParen()
{
    mOut << "(*";
    mOpenPointerParenCount++;
}

void GenMetalTraverser::emitClosingPointerParen()
{
    if (mOpenPointerParenCount > 0)
    {
        mOut << ")";
        mOpenPointerParenCount--;
    }
}

static const char *GetOperatorString(TOperator op,
                                     const TType &resultType,
                                     const TType *argType0,
                                     const TType *argType1,
                                     const TType *argType2)
{
    switch (op)
    {
        case TOperator::EOpComma:
            return ",";
        case TOperator::EOpAssign:
            return "=";
        case TOperator::EOpInitialize:
            return "=";
        case TOperator::EOpAddAssign:
            return "+=";
        case TOperator::EOpSubAssign:
            return "-=";
        case TOperator::EOpMulAssign:
            return "*=";
        case TOperator::EOpDivAssign:
            return "/=";
        case TOperator::EOpIModAssign:
            return "%=";
        case TOperator::EOpBitShiftLeftAssign:
            return "<<=";  // TODO: Check logical vs arithmetic shifting.
        case TOperator::EOpBitShiftRightAssign:
            return ">>=";  // TODO: Check logical vs arithmetic shifting.
        case TOperator::EOpBitwiseAndAssign:
            return "&=";
        case TOperator::EOpBitwiseXorAssign:
            return "^=";
        case TOperator::EOpBitwiseOrAssign:
            return "|=";
        case TOperator::EOpAdd:
            return "+";
        case TOperator::EOpSub:
            return "-";
        case TOperator::EOpMul:
            return "*";
        case TOperator::EOpDiv:
            return "/";
        case TOperator::EOpIMod:
            return "%";
        case TOperator::EOpBitShiftLeft:
            return "<<";  // TODO: Check logical vs arithmetic shifting.
        case TOperator::EOpBitShiftRight:
            return ">>";  // TODO: Check logical vs arithmetic shifting.
        case TOperator::EOpBitwiseAnd:
            return "&";
        case TOperator::EOpBitwiseXor:
            return "^";
        case TOperator::EOpBitwiseOr:
            return "|";
        case TOperator::EOpLessThan:
            return "<";
        case TOperator::EOpGreaterThan:
            return ">";
        case TOperator::EOpLessThanEqual:
            return "<=";
        case TOperator::EOpGreaterThanEqual:
            return ">=";
        case TOperator::EOpLessThanComponentWise:
            return "<";
        case TOperator::EOpLessThanEqualComponentWise:
            return "<=";
        case TOperator::EOpGreaterThanEqualComponentWise:
            return ">=";
        case TOperator::EOpGreaterThanComponentWise:
            return ">";
        case TOperator::EOpLogicalOr:
            return "||";
        case TOperator::EOpLogicalXor:
            return "!=/*xor*/";  // XXX: This might need to be handled differently for some obtuse
                                 // use case.
        case TOperator::EOpLogicalAnd:
            return "&&";
        case TOperator::EOpNegative:
            return "-";
        case TOperator::EOpPositive:
            if (argType0->isMatrix())
            {
                return "";
            }
            return "+";
        case TOperator::EOpLogicalNot:
            return "!";
        case TOperator::EOpNotComponentWise:
            return "!";
        case TOperator::EOpBitwiseNot:
            return "~";
        case TOperator::EOpPostIncrement:
            return "++";
        case TOperator::EOpPostDecrement:
            return "--";
        case TOperator::EOpPreIncrement:
            return "++";
        case TOperator::EOpPreDecrement:
            return "--";
        case TOperator::EOpVectorTimesScalarAssign:
            return "*=";
        case TOperator::EOpVectorTimesMatrixAssign:
            return "*=";
        case TOperator::EOpMatrixTimesScalarAssign:
            return "*=";
        case TOperator::EOpMatrixTimesMatrixAssign:
            return "*=";
        case TOperator::EOpVectorTimesScalar:
            return "*";
        case TOperator::EOpVectorTimesMatrix:
            return "*";
        case TOperator::EOpMatrixTimesVector:
            return "*";
        case TOperator::EOpMatrixTimesScalar:
            return "*";
        case TOperator::EOpMatrixTimesMatrix:
            return "*";
        case TOperator::EOpEqualComponentWise:
            return "==";
        case TOperator::EOpNotEqualComponentWise:
            return "!=";

        case TOperator::EOpEqual:
            if ((argType0->getStruct() && argType1->getStruct()) &&
                (argType0->isArray() && argType1->isArray()))
            {
                return "ANGLE_equalStructArray";
            }

            if ((argType0->isVector() && argType1->isVector()) ||
                (argType0->getStruct() && argType1->getStruct()) ||
                (argType0->isArray() && argType1->isArray()) ||
                (argType0->isMatrix() && argType1->isMatrix()))

            {
                return "ANGLE_equal";
            }

            return "==";

        case TOperator::EOpNotEqual:
            if ((argType0->getStruct() && argType1->getStruct()) &&
                (argType0->isArray() && argType1->isArray()))
            {
                return "ANGLE_notEqualStructArray";
            }

            if ((argType0->isVector() && argType1->isVector()) ||
                (argType0->isArray() && argType1->isArray()) ||
                (argType0->isMatrix() && argType1->isMatrix()))
            {
                return "ANGLE_notEqual";
            }
            else if (argType0->getStruct() && argType1->getStruct())
            {
                return "ANGLE_notEqualStruct";
            }
            return "!=";

        case TOperator::EOpKill:
            UNIMPLEMENTED();
            return "kill";
        case TOperator::EOpReturn:
            return "return";
        case TOperator::EOpBreak:
            return "break";
        case TOperator::EOpContinue:
            return "continue";

        case TOperator::EOpRadians:
            return "ANGLE_radians";
        case TOperator::EOpDegrees:
            return "ANGLE_degrees";
        case TOperator::EOpAtan:
            return argType1 == nullptr ? "metal::atan" : "metal::atan2";
        case TOperator::EOpMod:
            return "ANGLE_mod";  // differs from metal::mod
        case TOperator::EOpRefract:
            return argType0->isVector() ? "metal::refract" : "ANGLE_refract_scalar";
        case TOperator::EOpDistance:
            return argType0->isVector() ? "metal::distance" : "ANGLE_distance_scalar";
        case TOperator::EOpLength:
            return argType0->isVector() ? "metal::length" : "metal::abs";
        case TOperator::EOpDot:
            return argType0->isVector() ? "metal::dot" : "*";
        case TOperator::EOpNormalize:
            return argType0->isVector() ? "metal::fast::normalize" : "metal::sign";
        case TOperator::EOpFaceforward:
            return argType0->isVector() ? "metal::faceforward" : "ANGLE_faceforward_scalar";
        case TOperator::EOpReflect:
            return argType0->isVector() ? "metal::reflect" : "ANGLE_reflect_scalar";
        case TOperator::EOpMatrixCompMult:
            return "ANGLE_componentWiseMultiply";
        case TOperator::EOpOuterProduct:
            return "ANGLE_outerProduct";
        case TOperator::EOpSign:
            return argType0->getBasicType() == EbtFloat ? "metal::sign" : "ANGLE_sign_int";

        case TOperator::EOpAbs:
            return "metal::abs";
        case TOperator::EOpAll:
            return "metal::all";
        case TOperator::EOpAny:
            return "metal::any";
        case TOperator::EOpSin:
            return "metal::sin";
        case TOperator::EOpCos:
            return "metal::cos";
        case TOperator::EOpTan:
            return "metal::tan";
        case TOperator::EOpAsin:
            return "metal::asin";
        case TOperator::EOpAcos:
            return "metal::acos";
        case TOperator::EOpSinh:
            return "metal::sinh";
        case TOperator::EOpCosh:
            return "metal::cosh";
        case TOperator::EOpTanh:
            return resultType.getPrecision() == TPrecision::EbpHigh ? "metal::precise::tanh"
                                                                    : "metal::tanh";
        case TOperator::EOpAsinh:
            return "metal::asinh";
        case TOperator::EOpAcosh:
            return "metal::acosh";
        case TOperator::EOpAtanh:
            return "metal::atanh";
        case TOperator::EOpFma:
            return "metal::fma";
        case TOperator::EOpPow:
            return "metal::powr";  // GLSL's pow excludes negative x
        case TOperator::EOpExp:
            return "metal::exp";
        case TOperator::EOpExp2:
            return "metal::exp2";
        case TOperator::EOpLog:
            return "metal::log";
        case TOperator::EOpLog2:
            return "metal::log2";
        case TOperator::EOpSqrt:
            return "metal::sqrt";
        case TOperator::EOpFloor:
            return "metal::floor";
        case TOperator::EOpTrunc:
            return "metal::trunc";
        case TOperator::EOpCeil:
            return "metal::ceil";
        case TOperator::EOpFract:
            return "metal::fract";
        case TOperator::EOpMin:
            return "metal::min";
        case TOperator::EOpMax:
            return "metal::max";
        case TOperator::EOpRound:
            return "metal::round";
        case TOperator::EOpRoundEven:
            return "metal::rint";
        case TOperator::EOpClamp:
            return "metal::clamp";  // TODO fast vs precise namespace
        case TOperator::EOpSaturate:
            return "metal::saturate";  // TODO fast vs precise namespace
        case TOperator::EOpMix:
            if (!argType1->isScalar() && argType2 && argType2->getBasicType() == EbtBool)
            {
                return "ANGLE_mix_bool";
            }
            return "metal::mix";
        case TOperator::EOpStep:
            return "metal::step";
        case TOperator::EOpSmoothstep:
            return "metal::smoothstep";
        case TOperator::EOpModf:
            return "metal::modf";
        case TOperator::EOpIsnan:
            return "metal::isnan";
        case TOperator::EOpIsinf:
            return "metal::isinf";
        case TOperator::EOpLdexp:
            return "metal::ldexp";
        case TOperator::EOpFrexp:
            return "metal::frexp";
        case TOperator::EOpInversesqrt:
            return "metal::rsqrt";
        case TOperator::EOpCross:
            return "metal::cross";
        case TOperator::EOpDFdx:
            return "metal::dfdx";
        case TOperator::EOpDFdy:
            return "metal::dfdy";
        case TOperator::EOpFwidth:
            return "metal::fwidth";
        case TOperator::EOpTranspose:
            return "metal::transpose";
        case TOperator::EOpDeterminant:
            return "metal::determinant";

        case TOperator::EOpInverse:
            return "ANGLE_inverse";

        case TOperator::EOpInterpolateAtCentroid:
            return "ANGLE_interpolateAtCentroid";
        case TOperator::EOpInterpolateAtSample:
            return "ANGLE_interpolateAtSample";
        case TOperator::EOpInterpolateAtOffset:
            return "ANGLE_interpolateAtOffset";
        case TOperator::EOpInterpolateAtCenter:
            return "ANGLE_interpolateAtCenter";

        case TOperator::EOpFloatBitsToInt:
        case TOperator::EOpFloatBitsToUint:
        case TOperator::EOpIntBitsToFloat:
        case TOperator::EOpUintBitsToFloat:
        {
#define RETURN_AS_TYPE_SCALAR()             \
    do                                      \
        switch (resultType.getBasicType())  \
        {                                   \
            case TBasicType::EbtInt:        \
                return "as_type<int>";      \
            case TBasicType::EbtUInt:       \
                return "as_type<uint32_t>"; \
            case TBasicType::EbtFloat:      \
                return "as_type<float>";    \
            default:                        \
                UNIMPLEMENTED();            \
                return "TOperator_TODO";    \
        }                                   \
    while (false)

#define RETURN_AS_TYPE(post)                     \
    do                                           \
        switch (resultType.getBasicType())       \
        {                                        \
            case TBasicType::EbtInt:             \
                return "as_type<int" post ">";   \
            case TBasicType::EbtUInt:            \
                return "as_type<uint" post ">";  \
            case TBasicType::EbtFloat:           \
                return "as_type<float" post ">"; \
            default:                             \
                UNIMPLEMENTED();                 \
                return "TOperator_TODO";         \
        }                                        \
    while (false)

            if (resultType.isScalar())
            {
                RETURN_AS_TYPE_SCALAR();
            }
            else if (resultType.isVector())
            {
                switch (resultType.getNominalSize())
                {
                    case 2:
                        RETURN_AS_TYPE("2");
                    case 3:
                        RETURN_AS_TYPE("3");
                    case 4:
                        RETURN_AS_TYPE("4");
                    default:
                        UNREACHABLE();
                        return nullptr;
                }
            }
            else
            {
                UNIMPLEMENTED();
                return "TOperator_TODO";
            }

#undef RETURN_AS_TYPE
#undef RETURN_AS_TYPE_SCALAR
        }

        case TOperator::EOpPackUnorm2x16:
            return "metal::pack_float_to_unorm2x16";
        case TOperator::EOpPackSnorm2x16:
            return "metal::pack_float_to_snorm2x16";

        case TOperator::EOpPackUnorm4x8:
            return "metal::pack_float_to_unorm4x8";
        case TOperator::EOpPackSnorm4x8:
            return "metal::pack_float_to_snorm4x8";

        case TOperator::EOpUnpackUnorm2x16:
            return "metal::unpack_unorm2x16_to_float";
        case TOperator::EOpUnpackSnorm2x16:
            return "metal::unpack_snorm2x16_to_float";

        case TOperator::EOpUnpackUnorm4x8:
            return "metal::unpack_unorm4x8_to_float";
        case TOperator::EOpUnpackSnorm4x8:
            return "metal::unpack_snorm4x8_to_float";

        case TOperator::EOpPackHalf2x16:
            return "ANGLE_pack_half_2x16";
        case TOperator::EOpUnpackHalf2x16:
            return "ANGLE_unpack_half_2x16";

        case TOperator::EOpNumSamples:
            return "metal::get_num_samples";
        case TOperator::EOpSamplePosition:
            return "metal::get_sample_position";

        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();
            return "TOperator_TODO";

        case TOperator::EOpNull:
        case TOperator::EOpConstruct:
        case TOperator::EOpCallFunctionInAST:
        case TOperator::EOpCallInternalRawFunction:
        case TOperator::EOpIndexDirect:
        case TOperator::EOpIndexIndirect:
        case TOperator::EOpIndexDirectStruct:
        case TOperator::EOpIndexDirectInterfaceBlock:
            UNREACHABLE();
            return nullptr;
        default:
            // Any other built-in function.
            return nullptr;
    }
}

static bool IsSymbolicOperator(TOperator op,
                               const TType &resultType,
                               const TType *argType0,
                               const TType *argType1)
{
    const char *operatorString = GetOperatorString(op, resultType, argType0, argType1, nullptr);
    if (operatorString == nullptr)
    {
        return false;
    }
    return !std::isalnum(operatorString[0]);
}

static TIntermBinary *AsSpecificBinaryNode(TIntermNode &node, TOperator op)
{
    TIntermBinary *binaryNode = node.getAsBinaryNode();
    if (binaryNode)
    {
        return binaryNode->getOp() == op ? binaryNode : nullptr;
    }
    return nullptr;
}

static bool Parenthesize(TIntermNode &node)
{
    if (node.getAsSymbolNode())
    {
        return false;
    }
    if (node.getAsConstantUnion())
    {
        return false;
    }
    if (node.getAsAggregate())
    {
        return false;
    }
    if (node.getAsSwizzleNode())
    {
        return false;
    }

    if (TIntermUnary *unaryNode = node.getAsUnaryNode())
    {
        // TODO: Use a precedence and associativity rules instead of this ad-hoc impl.
        const TType &resultType = unaryNode->getType();
        const TType &argType    = unaryNode->getOperand()->getType();
        return IsSymbolicOperator(unaryNode->getOp(), resultType, &argType, nullptr);
    }

    if (TIntermBinary *binaryNode = node.getAsBinaryNode())
    {
        // TODO: Use a precedence and associativity rules instead of this ad-hoc impl.
        const TOperator op = binaryNode->getOp();
        switch (op)
        {
            case TOperator::EOpIndexDirectStruct:
            case TOperator::EOpIndexDirectInterfaceBlock:
            case TOperator::EOpIndexDirect:
            case TOperator::EOpIndexIndirect:
                return Parenthesize(*binaryNode->getLeft());

            case TOperator::EOpAssign:
            case TOperator::EOpInitialize:
                return AsSpecificBinaryNode(*binaryNode->getRight(), TOperator::EOpComma);

            default:
            {
                const TType &resultType = binaryNode->getType();
                const TType &leftType   = binaryNode->getLeft()->getType();
                const TType &rightType  = binaryNode->getRight()->getType();
                return IsSymbolicOperator(binaryNode->getOp(), resultType, &leftType, &rightType);
            }
        }
    }

    return true;
}

void GenMetalTraverser::groupedTraverse(TIntermNode &node)
{
    const bool emitParens = Parenthesize(node);

    if (emitParens)
    {
        mOut << "(";
    }

    node.traverse(this);

    if (emitParens)
    {
        mOut << ")";
    }
}

void GenMetalTraverser::emitPostQualifier(const EmitVariableDeclarationConfig &evdConfig,
                                          const VarDecl &decl,
                                          const TQualifier qualifier)
{
    bool isInvariant = false;
    switch (qualifier)
    {
        case TQualifier::EvqPosition:
            isInvariant = decl.type().isInvariant();
            [[fallthrough]];
        case TQualifier::EvqFragCoord:
            mOut << " [[position]]";
            break;

        case TQualifier::EvqClipDistance:
            mOut << " [[clip_distance]] [" << decl.type().getOutermostArraySize() << "]";
            break;

        case TQualifier::EvqPointSize:
            mOut << " [[point_size]]";
            break;

        case TQualifier::EvqVertexID:
            if (evdConfig.isMainParameter)
            {
                mOut << " [[vertex_id]]";
            }
            break;

        case TQualifier::EvqPointCoord:
            if (evdConfig.isMainParameter)
            {
                mOut << " [[point_coord]]";
            }
            break;

        case TQualifier::EvqFrontFacing:
            if (evdConfig.isMainParameter)
            {
                mOut << " [[front_facing]]";
            }
            break;

        case TQualifier::EvqSampleID:
            if (evdConfig.isMainParameter)
            {
                mOut << " [[sample_id]]";
            }
            break;

        case TQualifier::EvqSampleMaskIn:
            if (evdConfig.isMainParameter)
            {
                mOut << " [[sample_mask]]";
            }
            break;

        default:
            break;
    }

    if (isInvariant)
    {
        mOut << " [[invariant]]";

        TranslatorMetalReflection *reflection = mtl::getTranslatorMetalReflection(&mCompiler);
        reflection->hasInvariance             = true;
    }
}

void GenMetalTraverser::emitLoopBody(TIntermBlock *bodyNode)
{
    if (mInjectAsmStatementIntoLoopBodies)
    {
        emitOpenBrace();

        emitIndentation();
        mOut << "__asm__(\"\");\n";
    }

    bodyNode->traverse(this);

    if (mInjectAsmStatementIntoLoopBodies)
    {
        emitCloseBrace();
    }
}

static void EmitName(Sink &out, const Name &name)
{
#if defined(ANGLE_ENABLE_ASSERTS)
    DebugSink::EscapedSink escapedOut(out.escape());
#else
    TInfoSinkBase &escapedOut = out;
#endif
    name.emit(escapedOut);
}

void GenMetalTraverser::emitNameOf(const TField &object)
{
    EmitName(mOut, Name(object));
}

void GenMetalTraverser::emitNameOf(const TSymbol &object)
{
    auto it = mRenamedSymbols.find(&object);
    if (it == mRenamedSymbols.end())
    {
        EmitName(mOut, Name(object));
    }
    else
    {
        EmitName(mOut, it->second);
    }
}

void GenMetalTraverser::emitNameOf(const VarDecl &object)
{
    if (object.isField())
    {
        emitNameOf(object.field());
    }
    else
    {
        emitNameOf(object.variable());
    }
}

void GenMetalTraverser::emitBareTypeName(const TType &type, const EmitTypeConfig &etConfig)
{
    const TBasicType basicType = type.getBasicType();

    switch (basicType)
    {
        case TBasicType::EbtVoid:
        case TBasicType::EbtBool:
        case TBasicType::EbtFloat:
        case TBasicType::EbtInt:
        {
            mOut << type.getBasicString();
            break;
        }
        case TBasicType::EbtUInt:
        {
            if (type.isScalar())
            {
                mOut << "uint32_t";
            }
            else
            {
                mOut << type.getBasicString();
            }
        }
        break;

        case TBasicType::EbtStruct:
        {
            const TStructure &structure = *type.getStruct();
            emitNameOf(structure);
        }
        break;

        case TBasicType::EbtInterfaceBlock:
        {
            const TInterfaceBlock &interfaceBlock = *type.getInterfaceBlock();
            emitNameOf(interfaceBlock);
        }
        break;

        default:
        {
            if (IsSampler(basicType))
            {
                if (etConfig.evdConfig && etConfig.evdConfig->isMainParameter)
                {
                    EmitName(mOut, GetTextureTypeName(basicType));
                }
                else
                {
                    const TStructure &env = mSymbolEnv.getTextureEnv(basicType);
                    emitNameOf(env);
                }
            }
            else if (IsImage(basicType))
            {
                mOut << "metal::texture2d<";
                switch (type.getBasicType())
                {
                    case EbtImage2D:
                        mOut << "float";
                        break;
                    case EbtIImage2D:
                        mOut << "int";
                        break;
                    case EbtUImage2D:
                        mOut << "uint";
                        break;
                    default:
                        UNIMPLEMENTED();
                        break;
                }
                if (type.getMemoryQualifier().readonly || type.getMemoryQualifier().writeonly)
                {
                    UNIMPLEMENTED();
                }
                mOut << ", metal::access::read_write>";
            }
            else
            {
                UNIMPLEMENTED();
            }
        }
    }
}

void GenMetalTraverser::emitType(const TType &type, const EmitTypeConfig &etConfig)
{
    const bool isUBO = etConfig.evdConfig ? etConfig.evdConfig->isUBO : false;
    if (etConfig.evdConfig)
    {
        const auto &evdConfig = *etConfig.evdConfig;
        if (isUBO)
        {
            if (type.isArray())
            {
                mOut << "metal::array<";
            }
        }
        if (evdConfig.isPointer)
        {
            mOut << toString(*evdConfig.isPointer);
            mOut << " ";
        }
        else if (evdConfig.isReference)
        {
            mOut << toString(*evdConfig.isReference);
            mOut << " ";
        }
    }

    if (!isUBO)
    {
        if (type.isArray())
        {
            mOut << "metal::array<";
        }
    }

    if (type.isInterpolant())
    {
        mOut << "metal::interpolant<";
    }

    if (type.isVector() || type.isMatrix())
    {
        mOut << "metal::";
    }

    if (etConfig.evdConfig && etConfig.evdConfig->isPacked)
    {
        mOut << "packed_";
    }

    emitBareTypeName(type, etConfig);

    if (type.isVector())
    {
        mOut << static_cast<uint32_t>(type.getNominalSize());
    }
    else if (type.isMatrix())
    {
        mOut << static_cast<uint32_t>(type.getCols()) << "x"
             << static_cast<uint32_t>(type.getRows());
    }

    if (type.isInterpolant())
    {
        mOut << ", metal::interpolation::";
        switch (type.getQualifier())
        {
            case EvqNoPerspectiveIn:
            case EvqNoPerspectiveCentroidIn:
            case EvqNoPerspectiveSampleIn:
                mOut << "no_";
                break;
            default:
                break;
        }
        mOut << "perspective>";
    }

    if (!isUBO)
    {
        if (type.isArray())
        {
            for (auto size : type.getArraySizes())
            {
                mOut << ", " << size;
            }
            mOut << ">";
        }
    }

    if (etConfig.evdConfig)
    {
        const auto &evdConfig = *etConfig.evdConfig;
        if (evdConfig.isPointer)
        {
            mOut << " *";
        }
        else if (evdConfig.isReference)
        {
            mOut << " &";
        }
        if (isUBO)
        {
            if (type.isArray())
            {
                for (auto size : type.getArraySizes())
                {
                    mOut << ", " << size;
                }
                mOut << ">";
            }
        }
    }
}

void GenMetalTraverser::emitFieldDeclaration(const TField &field,
                                             const TStructure &parent,
                                             FieldAnnotationIndices &annotationIndices)
{
    const TType &type      = *field.type();
    const TBasicType basic = type.getBasicType();

    EmitVariableDeclarationConfig evdConfig;
    evdConfig.emitPostQualifier      = true;
    evdConfig.disableStructSpecifier = true;
    evdConfig.isPacked               = mSymbolEnv.isPacked(field);
    evdConfig.isUBO                  = mSymbolEnv.isUBO(field);
    evdConfig.isPointer              = mSymbolEnv.isPointer(field);
    evdConfig.isReference            = mSymbolEnv.isReference(field);
    emitVariableDeclaration(VarDecl(field), evdConfig);

    const TQualifier qual = type.getQualifier();
    switch (qual)
    {
        case TQualifier::EvqFlatIn:
            if (mPipelineStructs.fragmentIn.external == &parent)
            {
                mOut << " [[flat]]";
                TranslatorMetalReflection *reflection =
                    mtl::getTranslatorMetalReflection(&mCompiler);
                reflection->hasFlatInput = true;
            }
            break;

        case TQualifier::EvqNoPerspectiveIn:
            if (mPipelineStructs.fragmentIn.external == &parent && !type.isInterpolant())
            {
                mOut << " [[center_no_perspective]]";
            }
            break;

        case TQualifier::EvqCentroidIn:
            if (mPipelineStructs.fragmentIn.external == &parent && !type.isInterpolant())
            {
                mOut << " [[centroid_perspective]]";
            }
            break;

        case TQualifier::EvqSampleIn:
            if (mPipelineStructs.fragmentIn.external == &parent && !type.isInterpolant())
            {
                mOut << " [[sample_perspective]]";
            }
            break;

        case TQualifier::EvqNoPerspectiveCentroidIn:
            if (mPipelineStructs.fragmentIn.external == &parent && !type.isInterpolant())
            {
                mOut << " [[centroid_no_perspective]]";
            }
            break;

        case TQualifier::EvqNoPerspectiveSampleIn:
            if (mPipelineStructs.fragmentIn.external == &parent && !type.isInterpolant())
            {
                mOut << " [[sample_no_perspective]]";
            }
            break;

        case TQualifier::EvqFragColor:
            mOut << " [[color(0)]]";
            break;

        case TQualifier::EvqSecondaryFragColorEXT:
        case TQualifier::EvqSecondaryFragDataEXT:
            mOut << " [[color(0), index(1)]]";
            break;

        case TQualifier::EvqFragmentOut:
        case TQualifier::EvqFragmentInOut:
        case TQualifier::EvqFragData:
            if (mPipelineStructs.fragmentOut.external == &parent ||
                mPipelineStructs.fragmentOut.externalExtra == &parent)
            {
                if ((type.isVector() &&
                     (basic == TBasicType::EbtInt || basic == TBasicType::EbtUInt ||
                      basic == TBasicType::EbtFloat)) ||
                    qual == EvqFragData)
                {
                    // The OpenGL ES 3.0 spec says locations must be specified
                    // unless there is only a single output, in which case the
                    // location is 0. So, when we get to this point the shader
                    // will have been rejected if locations are not specified
                    // and there is more than one output.
                    const TLayoutQualifier &layoutQualifier = type.getLayoutQualifier();
                    if (layoutQualifier.locationsSpecified)
                    {
                        mOut << " [[color(" << layoutQualifier.location << ")";
                        ASSERT(layoutQualifier.index >= -1 && layoutQualifier.index <= 1);
                        if (layoutQualifier.index == 1)
                        {
                            mOut << ", index(1)";
                        }
                    }
                    else if (qual == EvqFragData)
                    {
                        mOut << " [[color(" << annotationIndices.color++ << ")";
                    }
                    else
                    {
                        // Either the only output or EXT_blend_func_extended is used;
                        // actual assignment will happen in UpdateFragmentShaderOutputs.
                        mOut << " [[" << sh::kUnassignedFragmentOutputString;
                    }
                    if (mRasterOrderGroupsSupported && qual == TQualifier::EvqFragmentInOut)
                    {
                        // Put fragment inouts in their own raster order group for better
                        // parallelism.
                        // NOTE: this is not required for the reads to be ordered and coherent.
                        // TODO(anglebug.com/40096838): Consider making raster order groups a PLS
                        // layout qualifier?
                        mOut << ", raster_order_group(0)";
                    }
                    mOut << "]]";
                }
            }
            break;

        case TQualifier::EvqFragDepth:
            mOut << " [[depth(";
            switch (type.getLayoutQualifier().depth)
            {
                case EdGreater:
                    mOut << "greater";
                    break;
                case EdLess:
                    mOut << "less";
                    break;
                default:
                    mOut << "any";
                    break;
            }
            mOut << "), function_constant(" << sh::mtl::kDepthWriteEnabledConstName << ")]]";
            break;

        case TQualifier::EvqSampleMask:
            if (field.symbolType() == SymbolType::AngleInternal)
            {
                mOut << " [[sample_mask, function_constant("
                     << sh::mtl::kSampleMaskWriteEnabledConstName << ")]]";
            }
            break;

        default:
            break;
    }

    if (IsImage(type.getBasicType()))
    {
        if (type.getArraySizeProduct() != 1)
        {
            UNIMPLEMENTED();
        }
        mOut << " [[";
        if (type.getLayoutQualifier().rasterOrdered)
        {
            mOut << "raster_order_group(0), ";
        }
        mOut << "texture(" << mMainTextureIndex << ")]]";
        TranslatorMetalReflection *reflection = mtl::getTranslatorMetalReflection(&mCompiler);
        reflection->addRWTextureBinding(type.getLayoutQualifier().binding,
                                        static_cast<int>(mMainTextureIndex));
        ++mMainTextureIndex;
    }
}

static std::map<Name, size_t> BuildExternalAttributeIndexMap(
    const TCompiler &compiler,
    const PipelineScoped<TStructure> &structure)
{
    ASSERT(structure.isTotallyFull());

    const auto &shaderVars     = compiler.getAttributes();
    const size_t shaderVarSize = shaderVars.size();
    size_t shaderVarIndex      = 0;

    const auto &externalFields = structure.external->fields();
    const size_t externalSize  = externalFields.size();
    size_t externalIndex       = 0;

    const auto &internalFields = structure.internal->fields();
    const size_t internalSize  = internalFields.size();
    size_t internalIndex       = 0;

    // Internal fields are never split. External fields are sometimes split.
    ASSERT(externalSize >= internalSize);

    // Structures do not contain any inactive fields.
    ASSERT(shaderVarSize >= internalSize);

    std::map<Name, size_t> externalNameToAttributeIndex;
    size_t attributeIndex = 0;

    while (internalIndex < internalSize)
    {
        const TField &internalField = *internalFields[internalIndex];
        const Name internalName     = Name(internalField);
        const TType &internalType   = *internalField.type();
        while (internalName.rawName() != shaderVars[shaderVarIndex].name &&
               internalName.rawName() != shaderVars[shaderVarIndex].mappedName)
        {
            // This case represents an inactive field.

            ++shaderVarIndex;
            ASSERT(shaderVarIndex < shaderVarSize);

            ++attributeIndex;  // TODO: Might need to increment more if shader var type is a matrix.
        }

        const size_t cols =
            (internalType.isMatrix() && !externalFields[externalIndex]->type()->isMatrix())
                ? internalType.getCols()
                : 1;

        for (size_t c = 0; c < cols; ++c)
        {
            const TField &externalField = *externalFields[externalIndex];
            const Name externalName     = Name(externalField);

            externalNameToAttributeIndex[externalName] = attributeIndex;

            ++externalIndex;
            ++attributeIndex;
        }

        ++shaderVarIndex;
        ++internalIndex;
    }

    ASSERT(shaderVarIndex <= shaderVarSize);
    ASSERT(externalIndex <= externalSize);  // less than if padding was introduced
    ASSERT(internalIndex == internalSize);

    return externalNameToAttributeIndex;
}

void GenMetalTraverser::emitAttributeDeclaration(const TField &field,
                                                 FieldAnnotationIndices &annotationIndices)
{
    EmitVariableDeclarationConfig evdConfig;
    evdConfig.disableStructSpecifier = true;
    emitVariableDeclaration(VarDecl(field), evdConfig);
    mOut << sh::kUnassignedAttributeString;
}

void GenMetalTraverser::emitUniformBufferDeclaration(const TField &field,
                                                     FieldAnnotationIndices &annotationIndices)
{
    EmitVariableDeclarationConfig evdConfig;
    evdConfig.disableStructSpecifier = true;
    evdConfig.isUBO                  = mSymbolEnv.isUBO(field);
    evdConfig.isPointer              = mSymbolEnv.isPointer(field);
    emitVariableDeclaration(VarDecl(field), evdConfig);
    mOut << "[[id(" << annotationIndices.attribute << ")]]";

    const TType &type   = *field.type();
    const int arraySize = type.isArray() ? type.getArraySizeProduct() : 1;

    TranslatorMetalReflection *reflection = mtl::getTranslatorMetalReflection(&mCompiler);
    ASSERT(type.getBasicType() == TBasicType::EbtStruct);
    const TStructure *structure    = type.getStruct();
    const std::string originalName = reflection->getOriginalName(structure->uniqueId().get());
    reflection->addUniformBufferBinding(
        originalName,
        {.bindIndex = annotationIndices.attribute, .arraySize = static_cast<size_t>(arraySize)});

    annotationIndices.attribute += arraySize;
}

void GenMetalTraverser::emitStructDeclaration(const TType &type)
{
    ASSERT(type.getBasicType() == TBasicType::EbtStruct);
    ASSERT(type.isStructSpecifier());

    mOut << "struct ";
    emitBareTypeName(type, {});

    mOut << "\n";
    emitOpenBrace();

    const TStructure &structure = *type.getStruct();
    std::map<Name, size_t> fieldToAttributeIndex;
    const bool hasAttributeIndices      = mPipelineStructs.vertexIn.external == &structure;
    const bool hasUniformBufferIndicies = mPipelineStructs.uniformBuffers.external == &structure;
    const bool reclaimUnusedAttributeIndices = mCompiler.getShaderVersion() < 300;

    if (hasAttributeIndices)
    {
        // When attribute aliasing is supported, external attribute struct is filled post-link.
        if (mCompiler.supportsAttributeAliasing())
        {
            mtl::getTranslatorMetalReflection(&mCompiler)->hasAttributeAliasing = true;
            mOut << "@@Attrib-Bindings@@\n";
            emitCloseBrace();
            return;
        }

        fieldToAttributeIndex =
            BuildExternalAttributeIndexMap(mCompiler, mPipelineStructs.vertexIn);
    }

    FieldAnnotationIndices annotationIndices;

    for (const TField *field : structure.fields())
    {
        emitIndentation();
        if (hasAttributeIndices)
        {
            const auto it = fieldToAttributeIndex.find(Name(*field));
            if (it == fieldToAttributeIndex.end())
            {
                ASSERT(field->symbolType() == SymbolType::AngleInternal);
                ASSERT(field->name().beginsWith("_"));
                ASSERT(angle::EndsWith(field->name().data(), "_pad"));
                emitFieldDeclaration(*field, structure, annotationIndices);
            }
            else
            {
                ASSERT(field->symbolType() != SymbolType::AngleInternal ||
                       !field->name().beginsWith("_") ||
                       !angle::EndsWith(field->name().data(), "_pad"));
                if (!reclaimUnusedAttributeIndices)
                {
                    annotationIndices.attribute = it->second;
                }
                emitAttributeDeclaration(*field, annotationIndices);
            }
        }
        else if (hasUniformBufferIndicies)
        {
            emitUniformBufferDeclaration(*field, annotationIndices);
        }
        else
        {
            emitFieldDeclaration(*field, structure, annotationIndices);
        }
        mOut << ";\n";
    }

    emitCloseBrace();
}

void GenMetalTraverser::emitOrdinaryVariableDeclaration(
    const VarDecl &decl,
    const EmitVariableDeclarationConfig &evdConfig)
{
    EmitTypeConfig etConfig;
    etConfig.evdConfig = &evdConfig;

    const TType &type = decl.type();
    if (type.getQualifier() == TQualifier::EvqClipDistance)
    {
        // Clip distance output uses float[n] type instead of metal::array.
        // The element count is emitted after the post qualifier.
        ASSERT(type.getBasicType() == TBasicType::EbtFloat);
        mOut << "float";
    }
    else if (type.getQualifier() == TQualifier::EvqSampleID && evdConfig.isMainParameter)
    {
        // Metal's [[sample_id]] must be unsigned
        ASSERT(type.getBasicType() == TBasicType::EbtInt);
        mOut << "uint32_t";
    }
    else
    {
        emitType(type, etConfig);
    }
    if (decl.symbolType() != SymbolType::Empty)
    {
        mOut << " ";
        emitNameOf(decl);
    }
}

void GenMetalTraverser::emitVariableDeclaration(const VarDecl &decl,
                                                const EmitVariableDeclarationConfig &evdConfig)
{
    const SymbolType symbolType = decl.symbolType();
    const TType &type           = decl.type();
    const TBasicType basicType  = type.getBasicType();

    switch (basicType)
    {
        case TBasicType::EbtStruct:
        {
            if (type.isStructSpecifier() && !evdConfig.disableStructSpecifier)
            {
                // It's invalid to declare a struct inside a function argument. When emitting a
                // function parameter, the callsite should set evdConfig.disableStructSpecifier.
                ASSERT(!evdConfig.isParameter);
                emitStructDeclaration(type);
                if (symbolType != SymbolType::Empty)
                {
                    mOut << " ";
                    emitNameOf(decl);
                }
            }
            else
            {
                emitOrdinaryVariableDeclaration(decl, evdConfig);
            }
        }
        break;

        default:
        {
            ASSERT(symbolType != SymbolType::Empty || evdConfig.isParameter);
            emitOrdinaryVariableDeclaration(decl, evdConfig);
        }
    }

    if (evdConfig.emitPostQualifier)
    {
        emitPostQualifier(evdConfig, decl, type.getQualifier());
    }
}

void GenMetalTraverser::visitSymbol(TIntermSymbol *symbolNode)
{
    const TVariable &var = symbolNode->variable();
    const TType &type    = var.getType();
    ASSERT(var.symbolType() != SymbolType::Empty);

    if (type.getBasicType() == TBasicType::EbtVoid)
    {
        mOut << "/*";
        emitNameOf(var);
        mOut << "*/";
    }
    else
    {
        emitNameOf(var);
    }
}

void GenMetalTraverser::emitSingleConstant(const TConstantUnion *const constUnion)
{
    switch (constUnion->getType())
    {
        case TBasicType::EbtBool:
        {
            mOut << (constUnion->getBConst() ? "true" : "false");
        }
        break;

        case TBasicType::EbtFloat:
        {
            float value = constUnion->getFConst();
            if (std::isnan(value))
            {
                mOut << "NAN";
            }
            else if (std::isinf(value))
            {
                if (value < 0)
                {
                    mOut << "-";
                }
                mOut << "INFINITY";
            }
            else
            {
                mOut << value << "f";
            }
        }
        break;

        case TBasicType::EbtInt:
        {
            mOut << constUnion->getIConst();
        }
        break;

        case TBasicType::EbtUInt:
        {
            mOut << constUnion->getUConst() << "u";
        }
        break;

        default:
        {
            UNIMPLEMENTED();
        }
    }
}

const TConstantUnion *GenMetalTraverser::emitConstantUnionArray(
    const TConstantUnion *const constUnion,
    const size_t size)
{
    const TConstantUnion *constUnionIterated = constUnion;
    for (size_t i = 0; i < size; i++, constUnionIterated++)
    {
        emitSingleConstant(constUnionIterated);

        if (i != size - 1)
        {
            mOut << ", ";
        }
    }
    return constUnionIterated;
}

const TConstantUnion *GenMetalTraverser::emitConstantUnion(const TType &type,
                                                           const TConstantUnion *constUnionBegin)
{
    const TConstantUnion *constUnionCurr = constUnionBegin;
    const TStructure *structure          = type.getStruct();
    if (structure)
    {
        EmitTypeConfig config = EmitTypeConfig{nullptr};
        emitType(type, config);
        mOut << "{";
        const TFieldList &fields = structure->fields();
        for (size_t i = 0; i < fields.size(); ++i)
        {
            const TType *fieldType = fields[i]->type();
            constUnionCurr         = emitConstantUnion(*fieldType, constUnionCurr);
            if (i != fields.size() - 1)
            {
                mOut << ", ";
            }
        }
        mOut << "}";
    }
    else
    {
        size_t size    = type.getObjectSize();
        bool writeType = size > 1;
        if (writeType)
        {
            EmitTypeConfig config = EmitTypeConfig{nullptr};
            emitType(type, config);
            mOut << "(";
        }
        constUnionCurr = emitConstantUnionArray(constUnionCurr, size);
        if (writeType)
        {
            mOut << ")";
        }
    }
    return constUnionCurr;
}

void GenMetalTraverser::visitConstantUnion(TIntermConstantUnion *constValueNode)
{
    emitConstantUnion(constValueNode->getType(), constValueNode->getConstantValue());
}

bool GenMetalTraverser::visitSwizzle(Visit, TIntermSwizzle *swizzleNode)
{
    groupedTraverse(*swizzleNode->getOperand());
    mOut << ".";

    {
#if defined(ANGLE_ENABLE_ASSERTS)
        DebugSink::EscapedSink escapedOut(mOut.escape());
        TInfoSinkBase &out = escapedOut.get();
#else
        TInfoSinkBase &out = mOut;
#endif
        swizzleNode->writeOffsetsAsXYZW(&out);
    }

    return false;
}

const TField &GenMetalTraverser::getDirectField(const TFieldListCollection &fieldListCollection,
                                                const TConstantUnion &index)
{
    ASSERT(index.getType() == TBasicType::EbtInt);

    const TFieldList &fieldList = fieldListCollection.fields();
    const int indexVal          = index.getIConst();
    const TField &field         = *fieldList[indexVal];

    return field;
}

const TField &GenMetalTraverser::getDirectField(const TIntermTyped &fieldsNode,
                                                TIntermTyped &indexNode)
{
    const TType &fieldsType = fieldsNode.getType();

    const TFieldListCollection *fieldListCollection = fieldsType.getStruct();
    if (fieldListCollection == nullptr)
    {
        fieldListCollection = fieldsType.getInterfaceBlock();
    }
    ASSERT(fieldListCollection);

    const TIntermConstantUnion *indexNode_ = indexNode.getAsConstantUnion();
    ASSERT(indexNode_);
    const TConstantUnion &index = *indexNode_->getConstantValue();

    return getDirectField(*fieldListCollection, index);
}

bool GenMetalTraverser::visitBinary(Visit, TIntermBinary *binaryNode)
{
    const TOperator op      = binaryNode->getOp();
    TIntermTyped &leftNode  = *binaryNode->getLeft();
    TIntermTyped &rightNode = *binaryNode->getRight();

    switch (op)
    {
        case TOperator::EOpIndexDirectStruct:
        case TOperator::EOpIndexDirectInterfaceBlock:
        {
            const TField &field = getDirectField(leftNode, rightNode);
            if (mSymbolEnv.isPointer(field) && mSymbolEnv.isUBO(field))
            {
                emitOpeningPointerParen();
            }
            groupedTraverse(leftNode);
            if (!mSymbolEnv.isPointer(field))
            {
                emitClosingPointerParen();
            }
            mOut << ".";
            emitNameOf(field);
        }
        break;

        case TOperator::EOpIndexDirect:
        case TOperator::EOpIndexIndirect:
        {
            TType leftType = leftNode.getType();
            groupedTraverse(leftNode);
            mOut << "[";
            const TConstantUnion *constIndex = rightNode.getConstantValue();
            // TODO(anglebug.com/42266914): Convert type and bound checks to
            // assertions after AST validation is enabled for MSL translation.
            if (!leftType.isUnsizedArray() && constIndex != nullptr &&
                constIndex->getType() == EbtInt && constIndex->getIConst() >= 0 &&
                constIndex->getIConst() < static_cast<int>(leftType.isArray()
                                                               ? leftType.getOutermostArraySize()
                                                               : leftType.getNominalSize()))
            {
                emitSingleConstant(constIndex);
            }
            else
            {
                mOut << "ANGLE_int_clamp(";
                groupedTraverse(rightNode);
                mOut << ", 0, ";
                if (leftType.isUnsizedArray())
                {
                    groupedTraverse(leftNode);
                    mOut << ".size()";
                }
                else
                {
                    uint32_t maxSize;
                    if (leftType.isArray())
                    {
                        maxSize = leftType.getOutermostArraySize() - 1;
                    }
                    else
                    {
                        maxSize = leftType.getNominalSize() - 1;
                    }
                    mOut << maxSize;
                }
                mOut << ")";
            }
            mOut << "]";
        }
        break;

        default:
        {
            const TType &resultType = binaryNode->getType();
            const TType &leftType   = leftNode.getType();
            const TType &rightType  = rightNode.getType();

            if (IsSymbolicOperator(op, resultType, &leftType, &rightType))
            {
                groupedTraverse(leftNode);
                if (op != TOperator::EOpComma)
                {
                    mOut << " ";
                }
                else
                {
                    emitClosingPointerParen();
                }
                mOut << GetOperatorString(op, resultType, &leftType, &rightType, nullptr) << " ";
                groupedTraverse(rightNode);
            }
            else
            {
                emitClosingPointerParen();
                mOut << GetOperatorString(op, resultType, &leftType, &rightType, nullptr) << "(";
                leftNode.traverse(this);
                mOut << ", ";
                rightNode.traverse(this);
                mOut << ")";
            }
        }
    }

    return false;
}

static bool IsPostfix(TOperator op)
{
    switch (op)
    {
        case TOperator::EOpPostIncrement:
        case TOperator::EOpPostDecrement:
            return true;

        default:
            return false;
    }
}

bool GenMetalTraverser::visitUnary(Visit, TIntermUnary *unaryNode)
{
    const TOperator op      = unaryNode->getOp();
    const TType &resultType = unaryNode->getType();

    TIntermTyped &arg    = *unaryNode->getOperand();
    const TType &argType = arg.getType();

    if (op == TOperator::EOpIsnan || op == TOperator::EOpIsinf)
    {
        mtl::getTranslatorMetalReflection(&mCompiler)->hasIsnanOrIsinf = true;
    }

    const char *name = GetOperatorString(op, resultType, &argType, nullptr, nullptr);

    if (IsSymbolicOperator(op, resultType, &argType, nullptr))
    {
        const bool postfix = IsPostfix(op);
        if (!postfix)
        {
            mOut << name;
        }
        groupedTraverse(arg);
        if (postfix)
        {
            mOut << name;
        }
    }
    else
    {
        mOut << name << "(";
        arg.traverse(this);
        mOut << ")";
    }

    return false;
}

bool GenMetalTraverser::visitTernary(Visit, TIntermTernary *conditionalNode)
{
    groupedTraverse(*conditionalNode->getCondition());
    mOut << " ? ";
    groupedTraverse(*conditionalNode->getTrueExpression());
    mOut << " : ";
    groupedTraverse(*conditionalNode->getFalseExpression());

    return false;
}

bool GenMetalTraverser::visitIfElse(Visit, TIntermIfElse *ifThenElseNode)
{
    TIntermTyped &condNode = *ifThenElseNode->getCondition();
    TIntermBlock *thenNode = ifThenElseNode->getTrueBlock();
    TIntermBlock *elseNode = ifThenElseNode->getFalseBlock();

    emitIndentation();
    mOut << "if (";
    condNode.traverse(this);
    mOut << ")";

    if (thenNode)
    {
        mOut << "\n";
        thenNode->traverse(this);
    }
    else
    {
        mOut << " {}";
    }

    if (elseNode)
    {
        mOut << "\n";
        emitIndentation();
        mOut << "else\n";
        elseNode->traverse(this);
    }
    else
    {
        // Always emit "else" even when empty block to avoid nested if-stmt issues.
        mOut << " else {}";
    }

    return false;
}

bool GenMetalTraverser::visitSwitch(Visit, TIntermSwitch *switchNode)
{
    emitIndentation();
    mOut << "switch (";
    switchNode->getInit()->traverse(this);
    mOut << ")\n";

    ASSERT(!mParentIsSwitch);
    mParentIsSwitch = true;
    switchNode->getStatementList()->traverse(this);
    mParentIsSwitch = false;

    return false;
}

bool GenMetalTraverser::visitCase(Visit, TIntermCase *caseNode)
{
    emitIndentation();

    if (caseNode->hasCondition())
    {
        TIntermTyped *condExpr = caseNode->getCondition();
        mOut << "case ";
        condExpr->traverse(this);
        mOut << ":";
    }
    else
    {
        mOut << "default:\n";
    }

    return false;
}

void GenMetalTraverser::emitFunctionSignature(const TFunction &func)
{
    const bool isMain = func.isMain();

    emitFunctionReturn(func);

    mOut << " ";
    emitNameOf(func);
    if (isMain)
    {
        mOut << "0";
    }
    mOut << "(";

    bool emitComma          = false;
    const size_t paramCount = func.getParamCount();
    for (size_t i = 0; i < paramCount; ++i)
    {
        if (emitComma)
        {
            mOut << ", ";
        }
        emitComma = true;

        const TVariable &param = *func.getParam(i);
        emitFunctionParameter(func, param);
    }

    if (isTraversingVertexMain)
    {
        mOut << " @@XFB-Bindings@@ ";
    }

    mOut << ")";
}

void GenMetalTraverser::emitFunctionReturn(const TFunction &func)
{
    const bool isMain       = func.isMain();
    bool isVertexMain       = false;
    const TType &returnType = func.getReturnType();
    if (isMain)
    {
        const TStructure *structure = returnType.getStruct();
        if (mPipelineStructs.fragmentOut.matches(*structure))
        {
            if (mCompiler.isEarlyFragmentTestsSpecified())
            {
                mOut << "[[early_fragment_tests]]\n";
            }
            mOut << "fragment ";
        }
        else if (mPipelineStructs.vertexOut.matches(*structure))
        {
            ASSERT(structure != nullptr);
            mOut << "vertex __VERTEX_OUT(";
            isVertexMain = true;
        }
        else
        {
            UNIMPLEMENTED();
        }
    }
    emitType(returnType, EmitTypeConfig());
    if (isVertexMain)
        mOut << ") ";
}

void GenMetalTraverser::emitFunctionParameter(const TFunction &func, const TVariable &param)
{
    const bool isMain = func.isMain();

    const TType &type           = param.getType();
    const TStructure *structure = type.getStruct();

    EmitVariableDeclarationConfig evdConfig;
    evdConfig.isParameter            = true;
    evdConfig.disableStructSpecifier = true;  // It's invalid to declare a struct in a function arg.
    evdConfig.isMainParameter        = isMain;
    evdConfig.emitPostQualifier      = isMain;
    evdConfig.isUBO                  = mSymbolEnv.isUBO(param);
    evdConfig.isPointer              = mSymbolEnv.isPointer(param);
    evdConfig.isReference            = mSymbolEnv.isReference(param);
    emitVariableDeclaration(VarDecl(param), evdConfig);

    if (isMain)
    {
        TranslatorMetalReflection *reflection = mtl::getTranslatorMetalReflection(&mCompiler);
        if (structure)
        {
            if (mPipelineStructs.fragmentIn.matches(*structure) ||
                mPipelineStructs.vertexIn.matches(*structure))
            {
                mOut << " [[stage_in]]";
            }
            else if (mPipelineStructs.angleUniforms.matches(*structure))
            {
                mOut << " [[buffer(" << mDriverUniformsBindingIndex << ")]]";
            }
            else if (mPipelineStructs.uniformBuffers.matches(*structure))
            {
                mOut << " [[buffer(" << mUBOArgumentBufferBindingIndex << ")]]";
                reflection->hasUBOs = true;
            }
            else if (mPipelineStructs.userUniforms.matches(*structure))
            {
                mOut << " [[buffer(" << mMainUniformBufferIndex << ")]]";
                reflection->addUserUniformBufferBinding(param.name().data(),
                                                        mMainUniformBufferIndex);
                mMainUniformBufferIndex += type.getArraySizeProduct();
            }
            else if (structure->name() == "metal::sampler")
            {
                mOut << " [[sampler(" << (mMainSamplerIndex) << ")]]";
                const std::string originalName =
                    reflection->getOriginalName(param.uniqueId().get());
                reflection->addSamplerBinding(originalName, mMainSamplerIndex);
                mMainSamplerIndex += type.getArraySizeProduct();
            }
        }
        else if (IsSampler(type.getBasicType()))
        {
            mOut << " [[texture(" << (mMainTextureIndex) << ")]]";
            const std::string originalName = reflection->getOriginalName(param.uniqueId().get());
            reflection->addTextureBinding(originalName, mMainTextureIndex);
            mMainTextureIndex += type.getArraySizeProduct();
        }
        else if (Name(param) == Pipeline{Pipeline::Type::InstanceId, nullptr}.getStructInstanceName(
                                    Pipeline::Variant::Modified))
        {
            mOut << " [[instance_id]]";
        }
        else if (Name(param) == kBaseInstanceName)
        {
            mOut << " [[base_instance]]";
        }
    }
}

void GenMetalTraverser::visitFunctionPrototype(TIntermFunctionPrototype *funcProtoNode)
{
    const TFunction &func = *funcProtoNode->getFunction();

    emitIndentation();
    emitFunctionSignature(func);
}

bool GenMetalTraverser::visitFunctionDefinition(Visit, TIntermFunctionDefinition *funcDefNode)
{
    const TFunction &func = *funcDefNode->getFunction();
    TIntermBlock &body    = *funcDefNode->getBody();
    if (func.isMain())
    {
        const TType &returnType     = func.getReturnType();
        const TStructure *structure = returnType.getStruct();
        isTraversingVertexMain      = (mPipelineStructs.vertexOut.matches(*structure)) &&
                                 mCompiler.getShaderType() == GL_VERTEX_SHADER;
    }
    emitIndentation();
    emitFunctionSignature(func);
    mOut << "\n";
    body.traverse(this);
    if (isTraversingVertexMain)
    {
        isTraversingVertexMain = false;
    }
    return false;
}

GenMetalTraverser::FuncToName GenMetalTraverser::BuildFuncToName()
{
    FuncToName map;

    auto putAngle = [&](const char *nameStr) {
        const ImmutableString name(nameStr);
        ASSERT(map.find(name) == map.end());
        map[name] = Name(nameStr, SymbolType::AngleInternal);
    };

    putAngle("texelFetch");
    putAngle("texelFetchOffset");
    putAngle("texture");
    putAngle("texture2D");
    putAngle("texture2DGradEXT");
    putAngle("texture2DLod");
    putAngle("texture2DLodEXT");
    putAngle("texture2DProj");
    putAngle("texture2DProjGradEXT");
    putAngle("texture2DProjLod");
    putAngle("texture2DProjLodEXT");
    putAngle("texture3D");
    putAngle("texture3DLod");
    putAngle("texture3DProj");
    putAngle("texture3DProjLod");
    putAngle("textureCube");
    putAngle("textureCubeGradEXT");
    putAngle("textureCubeLod");
    putAngle("textureCubeLodEXT");
    putAngle("textureGrad");
    putAngle("textureGradOffset");
    putAngle("textureLod");
    putAngle("textureLodOffset");
    putAngle("textureOffset");
    putAngle("textureProj");
    putAngle("textureProjGrad");
    putAngle("textureProjGradOffset");
    putAngle("textureProjLod");
    putAngle("textureProjLodOffset");
    putAngle("textureProjOffset");
    putAngle("textureSize");
    putAngle("imageLoad");
    putAngle("imageStore");
    putAngle("memoryBarrierImage");

    return map;
}

bool GenMetalTraverser::visitAggregate(Visit, TIntermAggregate *aggregateNode)
{
    const TIntermSequence &args = *aggregateNode->getSequence();

    auto emitArgList = [&](const char *open, const char *close) {
        mOut << open;

        bool emitComma = false;
        for (TIntermNode *arg : args)
        {
            if (emitComma)
            {
                emitClosingPointerParen();
                mOut << ", ";
            }
            emitComma = true;
            arg->traverse(this);
        }

        mOut << close;
    };

    const TType &retType = aggregateNode->getType();

    if (aggregateNode->isConstructor())
    {
        const bool isStandalone = getParentNode()->getAsBlock();
        if (isStandalone)
        {
            // Prevent constructor from being interpreted as a declaration by wrapping in parens.
            // This can happen if given something like:
            //      int(symbol); // <- This will be treated like `int symbol;`... don't want that.
            // So instead emit:
            //      (int(symbol));
            mOut << "(";
        }

        const EmitTypeConfig etConfig;

        if (retType.isArray())
        {
            emitType(retType, etConfig);
            emitArgList("{", "}");
        }
        else if (retType.getStruct())
        {
            emitType(retType, etConfig);
            emitArgList("{", "}");
        }
        else
        {
            emitType(retType, etConfig);
            emitArgList("(", ")");
        }

        if (isStandalone)
        {
            mOut << ")";
        }

        return false;
    }
    else
    {
        const TOperator op = aggregateNode->getOp();
        switch (op)
        {
            case TOperator::EOpCallFunctionInAST:
            case TOperator::EOpCallInternalRawFunction:
            {
                const TFunction &func = *aggregateNode->getFunction();
                emitNameOf(func);
                //'@' symbol in name specifices a macro substitution marker.
                if (!func.name().contains("@"))
                {
                    emitArgList("(", ")");
                }
                else
                {
                    mTemporarilyDisableSemicolon =
                        true;  // Disable semicolon for macro substitution.
                }
                return false;
            }

            default:
            {
                auto getArgType = [&](size_t index) -> const TType * {
                    if (index < args.size())
                    {
                        TIntermTyped *arg = args[index]->getAsTyped();
                        ASSERT(arg);
                        return &arg->getType();
                    }
                    return nullptr;
                };

                const TType *argType0 = getArgType(0);
                const TType *argType1 = getArgType(1);
                const TType *argType2 = getArgType(2);

                const char *opName = GetOperatorString(op, retType, argType0, argType1, argType2);

                if (IsSymbolicOperator(op, retType, argType0, argType1))
                {
                    switch (args.size())
                    {
                        case 1:
                        {
                            TIntermNode &operandNode = *aggregateNode->getChildNode(0);
                            if (IsPostfix(op))
                            {
                                mOut << opName;
                                groupedTraverse(operandNode);
                            }
                            else
                            {
                                groupedTraverse(operandNode);
                                mOut << opName;
                            }
                            return false;
                        }

                        case 2:
                        {
                            TIntermNode &leftNode  = *aggregateNode->getChildNode(0);
                            TIntermNode &rightNode = *aggregateNode->getChildNode(1);
                            groupedTraverse(leftNode);
                            mOut << " " << opName << " ";
                            groupedTraverse(rightNode);
                            return false;
                        }

                        default:
                            UNREACHABLE();
                            return false;
                    }
                }
                else if (opName == nullptr)
                {
                    const TFunction &func = *aggregateNode->getFunction();
                    auto it               = mFuncToName.find(func.name());
                    ASSERT(it != mFuncToName.end());
                    EmitName(mOut, it->second);
                    emitArgList("(", ")");
                    return false;
                }
                else
                {
                    mOut << opName;
                    emitArgList("(", ")");
                    return false;
                }
            }
        }
    }
}

void GenMetalTraverser::emitOpenBrace()
{
    ASSERT(mIndentLevel >= 0);

    emitIndentation();
    mOut << "{\n";
    ++mIndentLevel;
}

void GenMetalTraverser::emitCloseBrace()
{
    ASSERT(mIndentLevel >= 1);

    --mIndentLevel;
    emitIndentation();
    mOut << "}";
}

static bool RequiresSemicolonTerminator(TIntermNode &node)
{
    if (node.getAsBlock())
    {
        return false;
    }
    if (node.getAsLoopNode())
    {
        return false;
    }
    if (node.getAsSwitchNode())
    {
        return false;
    }
    if (node.getAsIfElseNode())
    {
        return false;
    }
    if (node.getAsFunctionDefinition())
    {
        return false;
    }
    if (node.getAsCaseNode())
    {
        return false;
    }

    return true;
}

static bool NewlinePad(TIntermNode &node)
{
    if (node.getAsFunctionDefinition())
    {
        return true;
    }
    if (TIntermDeclaration *declNode = node.getAsDeclarationNode())
    {
        ASSERT(declNode->getChildCount() == 1);
        TIntermNode &childNode = *declNode->getChildNode(0);
        if (TIntermSymbol *symbolNode = childNode.getAsSymbolNode())
        {
            const TVariable &var = symbolNode->variable();
            return var.getType().isStructSpecifier();
        }
        return false;
    }
    return false;
}

bool GenMetalTraverser::visitBlock(Visit, TIntermBlock *blockNode)
{
    ASSERT(mIndentLevel >= -1);
    const bool isGlobalScope  = mIndentLevel == -1;
    const bool parentIsSwitch = mParentIsSwitch;
    mParentIsSwitch           = false;

    if (isGlobalScope)
    {
        ++mIndentLevel;
    }
    else
    {
        emitOpenBrace();
        if (parentIsSwitch)
        {
            ++mIndentLevel;
        }
    }

    TIntermNode *prevStmtNode = nullptr;

    const size_t stmtCount = blockNode->getChildCount();
    for (size_t i = 0; i < stmtCount; ++i)
    {
        TIntermNode &stmtNode = *blockNode->getChildNode(i);

        if (isGlobalScope && prevStmtNode && (NewlinePad(*prevStmtNode) || NewlinePad(stmtNode)))
        {
            mOut << "\n";
        }
        const bool isCase = stmtNode.getAsCaseNode();
        mIndentLevel -= isCase;
        emitIndentation();
        mIndentLevel += isCase;
        stmtNode.traverse(this);
        if (RequiresSemicolonTerminator(stmtNode) && !mTemporarilyDisableSemicolon)
        {
            mOut << ";";
        }
        mTemporarilyDisableSemicolon = false;
        mOut << "\n";

        prevStmtNode = &stmtNode;
    }

    if (isGlobalScope)
    {
        ASSERT(mIndentLevel == 0);
        --mIndentLevel;
    }
    else
    {
        if (parentIsSwitch)
        {
            ASSERT(mIndentLevel >= 1);
            --mIndentLevel;
        }
        emitCloseBrace();
        mParentIsSwitch = parentIsSwitch;
    }

    return false;
}

bool GenMetalTraverser::visitGlobalQualifierDeclaration(Visit, TIntermGlobalQualifierDeclaration *)
{
    return false;
}

bool GenMetalTraverser::visitDeclaration(Visit, TIntermDeclaration *declNode)
{
    ASSERT(declNode->getChildCount() == 1);
    TIntermNode &node = *declNode->getChildNode(0);

    EmitVariableDeclarationConfig evdConfig;

    if (TIntermSymbol *symbolNode = node.getAsSymbolNode())
    {
        const TVariable &var = symbolNode->variable();
        emitVariableDeclaration(VarDecl(var), evdConfig);
        if (var.getType().isArray() && var.getType().getQualifier() == EvqTemporary)
        {
            // The translator frontend injects a loop-based init for user arrays when the source
            // shader is using ESSL 1.00. Some Metal drivers may fail to access elements of such
            // arrays at runtime depending on the array size. An empty literal initializer added
            // to the generated MSL bypasses the issue. The frontend may be further optimized to
            // skip the loop-based init when targeting MSL.
            mOut << "{}";
        }
    }
    else if (TIntermBinary *initNode = node.getAsBinaryNode())
    {
        ASSERT(initNode->getOp() == TOperator::EOpInitialize);
        TIntermSymbol *leftSymbolNode = initNode->getLeft()->getAsSymbolNode();
        TIntermTyped *valueNode       = initNode->getRight()->getAsTyped();
        ASSERT(leftSymbolNode && valueNode);

        if (getRootNode() == getParentBlock())
        {
            // DeferGlobalInitializers should have turned non-const global initializers into
            // deferred initializers. Note that variables marked as EvqGlobal can be treated as
            // EvqConst in some ANGLE code but not actually have their qualifier actually changed to
            // EvqConst. Thus just assume all EvqGlobal are actually EvqConst for all code run after
            // DeferGlobalInitializers.
            mOut << "constant ";
        }

        const TVariable &var = leftSymbolNode->variable();
        const Name varName(var);

        if (ExpressionContainsName(varName, *valueNode))
        {
            mRenamedSymbols[&var] = mIdGen.createNewName(varName);
        }

        emitVariableDeclaration(VarDecl(var), evdConfig);
        mOut << " = ";
        groupedTraverse(*valueNode);
    }
    else
    {
        UNREACHABLE();
    }

    return false;
}

bool GenMetalTraverser::visitLoop(Visit, TIntermLoop *loopNode)
{
    const TLoopType loopType = loopNode->getType();

    switch (loopType)
    {
        case TLoopType::ELoopFor:
            return visitForLoop(loopNode);
        case TLoopType::ELoopWhile:
            return visitWhileLoop(loopNode);
        case TLoopType::ELoopDoWhile:
            return visitDoWhileLoop(loopNode);
    }
}

bool GenMetalTraverser::visitForLoop(TIntermLoop *loopNode)
{
    ASSERT(loopNode->getType() == TLoopType::ELoopFor);

    TIntermNode *initNode  = loopNode->getInit();
    TIntermTyped *condNode = loopNode->getCondition();
    TIntermTyped *exprNode = loopNode->getExpression();

    mOut << "for (";

    if (initNode)
    {
        initNode->traverse(this);
    }
    else
    {
        mOut << " ";
    }

    mOut << "; ";

    if (condNode)
    {
        condNode->traverse(this);
    }

    mOut << "; ";

    if (exprNode)
    {
        exprNode->traverse(this);
    }

    mOut << ")\n";

    emitLoopBody(loopNode->getBody());

    return false;
}

bool GenMetalTraverser::visitWhileLoop(TIntermLoop *loopNode)
{
    ASSERT(loopNode->getType() == TLoopType::ELoopWhile);

    TIntermNode *initNode  = loopNode->getInit();
    TIntermTyped *condNode = loopNode->getCondition();
    TIntermTyped *exprNode = loopNode->getExpression();
    ASSERT(condNode);
    ASSERT(!initNode && !exprNode);

    emitIndentation();
    mOut << "while (";
    condNode->traverse(this);
    mOut << ")\n";
    emitLoopBody(loopNode->getBody());

    return false;
}

bool GenMetalTraverser::visitDoWhileLoop(TIntermLoop *loopNode)
{
    ASSERT(loopNode->getType() == TLoopType::ELoopDoWhile);

    TIntermNode *initNode  = loopNode->getInit();
    TIntermTyped *condNode = loopNode->getCondition();
    TIntermTyped *exprNode = loopNode->getExpression();
    ASSERT(condNode);
    ASSERT(!initNode && !exprNode);

    emitIndentation();
    mOut << "do\n";
    emitLoopBody(loopNode->getBody());
    mOut << "\n";
    emitIndentation();
    mOut << "while (";
    condNode->traverse(this);
    mOut << ");";

    return false;
}

bool GenMetalTraverser::visitBranch(Visit, TIntermBranch *branchNode)
{
    const TOperator flowOp = branchNode->getFlowOp();
    TIntermTyped *exprNode = branchNode->getExpression();

    emitIndentation();

    switch (flowOp)
    {
        case TOperator::EOpKill:
        {
            ASSERT(exprNode == nullptr);
            mOut << "metal::discard_fragment()";
        }
        break;

        case TOperator::EOpReturn:
        {
            if (isTraversingVertexMain)
            {
                mOut << "#if TRANSFORM_FEEDBACK_ENABLED\n";
                emitIndentation();
                mOut << "return;\n";
                emitIndentation();
                mOut << "#else\n";
                emitIndentation();
            }
            mOut << "return";
            if (exprNode)
            {
                mOut << " ";
                exprNode->traverse(this);
                mOut << ";";
            }
            if (isTraversingVertexMain)
            {
                mOut << "\n";
                emitIndentation();
                mOut << "#endif\n";
                mTemporarilyDisableSemicolon = true;
            }
        }
        break;

        case TOperator::EOpBreak:
        {
            ASSERT(exprNode == nullptr);
            mOut << "break";
        }
        break;

        case TOperator::EOpContinue:
        {
            ASSERT(exprNode == nullptr);
            mOut << "continue";
        }
        break;

        default:
        {
            UNREACHABLE();
        }
    }

    return false;
}

static size_t emitMetalCallCount = 0;

bool sh::EmitMetal(TCompiler &compiler,
                   TIntermBlock &root,
                   IdGen &idGen,
                   const PipelineStructs &pipelineStructs,
                   SymbolEnv &symbolEnv,
                   const ProgramPreludeConfig &ppc,
                   const ShCompileOptions &compileOptions)
{
    TInfoSinkBase &out = compiler.getInfoSink().obj;

    {
        ++emitMetalCallCount;
        std::string filenameProto = angle::GetEnvironmentVar("GMD_FIXED_EMIT");
        if (!filenameProto.empty())
        {
            if (filenameProto != "/dev/null")
            {
                auto tryOpen = [&](char const *ext) {
                    auto filename = filenameProto;
                    filename += std::to_string(emitMetalCallCount);
                    filename += ".";
                    filename += ext;
                    return fopen(filename.c_str(), "rb");
                };
                FILE *file = tryOpen("metal");
                if (!file)
                {
                    file = tryOpen("cpp");
                }
                ASSERT(file);

                fseek(file, 0, SEEK_END);
                size_t fileSize = ftell(file);
                fseek(file, 0, SEEK_SET);

                std::vector<char> buff;
                buff.resize(fileSize + 1);
                fread(buff.data(), fileSize, 1, file);
                buff.back() = '\0';

                fclose(file);

                out << buff.data();
            }

            return true;
        }
    }

    out << "\n\n";

    if (!EmitProgramPrelude(root, out, ppc))
    {
        return false;
    }

    {
#if defined(ANGLE_ENABLE_ASSERTS)
        DebugSink outWrapper(out, angle::GetBoolEnvironmentVar("GMD_STDOUT"));
        outWrapper.watch(angle::GetEnvironmentVar("GMD_WATCH_STRING"));
#else
        TInfoSinkBase &outWrapper = out;
#endif
        GenMetalTraverser gen(compiler, outWrapper, idGen, pipelineStructs, symbolEnv,
                              compileOptions);
        root.traverse(&gen);
    }

    out << "\n";

    return true;
}
kc3-lang/angle/src/compiler/translator/msl/EmitMetal.cpp

Commit

src/compiler/translator/msl/EmitMetal.cpp
