// Copyright 2020 yuzu Emulator Project
// Licensed under GPLv2 or any later version
// Refer to the license.txt file included.

#include <algorithm>
#include <array>
#include <cstddef>
#include <string>
#include <string_view>
#include <utility>
#include <variant>

#include <fmt/format.h>

#include "common/alignment.h"
#include "common/assert.h"
#include "common/common_types.h"
#include "video_core/renderer_opengl/gl_arb_decompiler.h"
#include "video_core/renderer_opengl/gl_device.h"
#include "video_core/shader/registry.h"
#include "video_core/shader/shader_ir.h"

// Predicates in the decompiled code follow the convention that -1 means true and 0 means false.
// GLASM lacks booleans, so they have to be implemented as integers.
// Using -1 for true is useful because both CMP.S and NOT.U can negate it, and CMP.S can be used to
// select between two values, because -1 will be evaluated as true and 0 as false.

namespace OpenGL {

namespace {

using Tegra::Engines::ShaderType;
using Tegra::Shader::Attribute;
using Tegra::Shader::PixelImap;
using Tegra::Shader::Register;
using namespace VideoCommon::Shader;
using Operation = const OperationNode&;

constexpr std::array INTERNAL_FLAG_NAMES = {"ZERO", "SIGN", "CARRY", "OVERFLOW"};

char Swizzle(std::size_t component) {
    static constexpr std::string_view SWIZZLE{"xyzw"};
    return SWIZZLE.at(component);
}

constexpr bool IsGenericAttribute(Attribute::Index index) {
    return index >= Attribute::Index::Attribute_0 && index <= Attribute::Index::Attribute_31;
}

u32 GetGenericAttributeIndex(Attribute::Index index) {
    ASSERT(IsGenericAttribute(index));
    return static_cast<u32>(index) - static_cast<u32>(Attribute::Index::Attribute_0);
}

std::string_view Modifiers(Operation operation) {
    const auto meta = std::get_if<MetaArithmetic>(&operation.GetMeta());
    if (meta && meta->precise) {
        return ".PREC";
    }
    return "";
}

std::string_view GetInputFlags(PixelImap attribute) {
    switch (attribute) {
    case PixelImap::Perspective:
        return "";
    case PixelImap::Constant:
        return "FLAT ";
    case PixelImap::ScreenLinear:
        return "NOPERSPECTIVE ";
    case PixelImap::Unused:
        break;
    }
    UNIMPLEMENTED_MSG("Unknown attribute usage index={}", attribute);
    return {};
}

std::string_view ImageType(Tegra::Shader::ImageType image_type) {
    switch (image_type) {
    case Tegra::Shader::ImageType::Texture1D:
        return "1D";
    case Tegra::Shader::ImageType::TextureBuffer:
        return "BUFFER";
    case Tegra::Shader::ImageType::Texture1DArray:
        return "ARRAY1D";
    case Tegra::Shader::ImageType::Texture2D:
        return "2D";
    case Tegra::Shader::ImageType::Texture2DArray:
        return "ARRAY2D";
    case Tegra::Shader::ImageType::Texture3D:
        return "3D";
    }
    UNREACHABLE();
    return {};
}

std::string_view StackName(MetaStackClass stack) {
    switch (stack) {
    case MetaStackClass::Ssy:
        return "SSY";
    case MetaStackClass::Pbk:
        return "PBK";
    }
    UNREACHABLE();
    return "";
};

std::string_view PrimitiveDescription(Tegra::Engines::Maxwell3D::Regs::PrimitiveTopology topology) {
    switch (topology) {
    case Tegra::Engines::Maxwell3D::Regs::PrimitiveTopology::Points:
        return "POINTS";
    case Tegra::Engines::Maxwell3D::Regs::PrimitiveTopology::Lines:
    case Tegra::Engines::Maxwell3D::Regs::PrimitiveTopology::LineStrip:
        return "LINES";
    case Tegra::Engines::Maxwell3D::Regs::PrimitiveTopology::LinesAdjacency:
    case Tegra::Engines::Maxwell3D::Regs::PrimitiveTopology::LineStripAdjacency:
        return "LINES_ADJACENCY";
    case Tegra::Engines::Maxwell3D::Regs::PrimitiveTopology::Triangles:
    case Tegra::Engines::Maxwell3D::Regs::PrimitiveTopology::TriangleStrip:
    case Tegra::Engines::Maxwell3D::Regs::PrimitiveTopology::TriangleFan:
        return "TRIANGLES";
    case Tegra::Engines::Maxwell3D::Regs::PrimitiveTopology::TrianglesAdjacency:
    case Tegra::Engines::Maxwell3D::Regs::PrimitiveTopology::TriangleStripAdjacency:
        return "TRIANGLES_ADJACENCY";
    default:
        UNIMPLEMENTED_MSG("topology={}", topology);
        return "POINTS";
    }
}

std::string_view TopologyName(Tegra::Shader::OutputTopology topology) {
    switch (topology) {
    case Tegra::Shader::OutputTopology::PointList:
        return "POINTS";
    case Tegra::Shader::OutputTopology::LineStrip:
        return "LINE_STRIP";
    case Tegra::Shader::OutputTopology::TriangleStrip:
        return "TRIANGLE_STRIP";
    default:
        UNIMPLEMENTED_MSG("Unknown output topology: {}", topology);
        return "points";
    }
}

std::string_view StageInputName(ShaderType stage) {
    switch (stage) {
    case ShaderType::Vertex:
    case ShaderType::Geometry:
        return "vertex";
    case ShaderType::Fragment:
        return "fragment";
    case ShaderType::Compute:
        return "invocation";
    default:
        UNREACHABLE();
        return "";
    }
}

std::string TextureType(const MetaTexture& meta) {
    if (meta.sampler.is_buffer) {
        return "BUFFER";
    }
    std::string type;
    if (meta.sampler.is_shadow) {
        type += "SHADOW";
    }
    if (meta.sampler.is_array) {
        type += "ARRAY";
    }
    type += [&meta] {
        switch (meta.sampler.type) {
        case Tegra::Shader::TextureType::Texture1D:
            return "1D";
        case Tegra::Shader::TextureType::Texture2D:
            return "2D";
        case Tegra::Shader::TextureType::Texture3D:
            return "3D";
        case Tegra::Shader::TextureType::TextureCube:
            return "CUBE";
        }
        UNREACHABLE();
        return "2D";
    }();
    return type;
}

class ARBDecompiler final {
public:
    explicit ARBDecompiler(const Device& device_, const ShaderIR& ir_, const Registry& registry_,
                           ShaderType stage_, std::string_view identifier);

    std::string Code() const {
        return shader_source;
    }

private:
    void DefineGlobalMemory();

    void DeclareHeader();
    void DeclareVertex();
    void DeclareGeometry();
    void DeclareFragment();
    void DeclareCompute();
    void DeclareInputAttributes();
    void DeclareOutputAttributes();
    void DeclareLocalMemory();
    void DeclareGlobalMemory();
    void DeclareConstantBuffers();
    void DeclareRegisters();
    void DeclareTemporaries();
    void DeclarePredicates();
    void DeclareInternalFlags();

    void InitializeVariables();

    void DecompileAST();
    void DecompileBranchMode();

    void VisitAST(const ASTNode& node);
    std::string VisitExpression(const Expr& node);

    void VisitBlock(const NodeBlock& bb);

    std::string Visit(const Node& node);

    std::tuple<std::string, std::string, std::size_t> BuildCoords(Operation);
    std::string BuildAoffi(Operation);
    std::string GlobalMemoryPointer(const GmemNode& gmem);
    void Exit();

    std::string Assign(Operation);
    std::string Select(Operation);
    std::string FClamp(Operation);
    std::string FCastHalf0(Operation);
    std::string FCastHalf1(Operation);
    std::string FSqrt(Operation);
    std::string FSwizzleAdd(Operation);
    std::string HAdd2(Operation);
    std::string HMul2(Operation);
    std::string HFma2(Operation);
    std::string HAbsolute(Operation);
    std::string HNegate(Operation);
    std::string HClamp(Operation);
    std::string HCastFloat(Operation);
    std::string HUnpack(Operation);
    std::string HMergeF32(Operation);
    std::string HMergeH0(Operation);
    std::string HMergeH1(Operation);
    std::string HPack2(Operation);
    std::string LogicalAssign(Operation);
    std::string LogicalPick2(Operation);
    std::string LogicalAnd2(Operation);
    std::string FloatOrdered(Operation);
    std::string FloatUnordered(Operation);
    std::string LogicalAddCarry(Operation);
    std::string Texture(Operation);
    std::string TextureGather(Operation);
    std::string TextureQueryDimensions(Operation);
    std::string TextureQueryLod(Operation);
    std::string TexelFetch(Operation);
    std::string TextureGradient(Operation);
    std::string ImageLoad(Operation);
    std::string ImageStore(Operation);
    std::string Branch(Operation);
    std::string BranchIndirect(Operation);
    std::string PushFlowStack(Operation);
    std::string PopFlowStack(Operation);
    std::string Exit(Operation);
    std::string Discard(Operation);
    std::string EmitVertex(Operation);
    std::string EndPrimitive(Operation);
    std::string InvocationId(Operation);
    std::string YNegate(Operation);
    std::string ThreadId(Operation);
    std::string ShuffleIndexed(Operation);
    std::string Barrier(Operation);
    std::string MemoryBarrierGroup(Operation);
    std::string MemoryBarrierGlobal(Operation);

    template <const std::string_view& op>
    std::string Unary(Operation operation) {
        std::string temporary = AllocTemporary();
        AddLine("{}{} {}, {};", op, Modifiers(operation), temporary, Visit(operation[0]));
        return temporary;
    }

    template <const std::string_view& op>
    std::string Binary(Operation operation) {
        std::string temporary = AllocTemporary();
        AddLine("{}{} {}, {}, {};", op, Modifiers(operation), temporary, Visit(operation[0]),
                Visit(operation[1]));
        return temporary;
    }

    template <const std::string_view& op>
    std::string Trinary(Operation operation) {
        std::string temporary = AllocTemporary();
        AddLine("{}{} {}, {}, {}, {};", op, Modifiers(operation), temporary, Visit(operation[0]),
                Visit(operation[1]), Visit(operation[2]));
        return temporary;
    }

    template <const std::string_view& op, bool unordered>
    std::string FloatComparison(Operation operation) {
        std::string temporary = AllocTemporary();
        AddLine("TRUNC.U.CC RC.x, {};", Binary<op>(operation));
        AddLine("MOV.S {}, 0;", temporary);
        AddLine("MOV.S {} (NE.x), -1;", temporary);

        const std::string op_a = Visit(operation[0]);
        const std::string op_b = Visit(operation[1]);
        if constexpr (unordered) {
            AddLine("SNE.F RC.x, {}, {};", op_a, op_a);
            AddLine("TRUNC.U.CC RC.x, RC.x;");
            AddLine("MOV.S {} (NE.x), -1;", temporary);
            AddLine("SNE.F RC.x, {}, {};", op_b, op_b);
            AddLine("TRUNC.U.CC RC.x, RC.x;");
            AddLine("MOV.S {} (NE.x), -1;", temporary);
        } else if (op == SNE_F) {
            AddLine("SNE.F RC.x, {}, {};", op_a, op_a);
            AddLine("TRUNC.U.CC RC.x, RC.x;");
            AddLine("MOV.S {} (NE.x), 0;", temporary);
            AddLine("SNE.F RC.x, {}, {};", op_b, op_b);
            AddLine("TRUNC.U.CC RC.x, RC.x;");
            AddLine("MOV.S {} (NE.x), 0;", temporary);
        }
        return temporary;
    }

    template <const std::string_view& op, bool is_nan>
    std::string HalfComparison(Operation operation) {
        std::string tmp1 = AllocVectorTemporary();
        const std::string tmp2 = AllocVectorTemporary();
        const std::string op_a = Visit(operation[0]);
        const std::string op_b = Visit(operation[1]);
        AddLine("UP2H.F {}, {};", tmp1, op_a);
        AddLine("UP2H.F {}, {};", tmp2, op_b);
        AddLine("{} {}, {}, {};", op, tmp1, tmp1, tmp2);
        AddLine("TRUNC.U.CC RC.xy, {};", tmp1);
        AddLine("MOV.S {}.xy, {{0, 0, 0, 0}};", tmp1);
        AddLine("MOV.S {}.x (NE.x), -1;", tmp1);
        AddLine("MOV.S {}.y (NE.y), -1;", tmp1);
        if constexpr (is_nan) {
            AddLine("MOVC.F RC.x, {};", op_a);
            AddLine("MOV.S {}.x (NAN.x), -1;", tmp1);
            AddLine("MOVC.F RC.x, {};", op_b);
            AddLine("MOV.S {}.y (NAN.x), -1;", tmp1);
        }
        return tmp1;
    }

    template <const std::string_view& op, const std::string_view& type>
    std::string AtomicImage(Operation operation) {
        const auto& meta = std::get<MetaImage>(operation.GetMeta());
        const u32 image_id = device.GetBaseBindings(stage).image + meta.image.index;
        const std::size_t num_coords = operation.GetOperandsCount();
        const std::size_t num_values = meta.values.size();

        const std::string coord = AllocVectorTemporary();
        const std::string value = AllocVectorTemporary();
        for (std::size_t i = 0; i < num_coords; ++i) {
            AddLine("MOV.S {}.{}, {};", coord, Swizzle(i), Visit(operation[i]));
        }
        for (std::size_t i = 0; i < num_values; ++i) {
            AddLine("MOV.F {}.{}, {};", value, Swizzle(i), Visit(meta.values[i]));
        }

        AddLine("ATOMIM.{}.{} {}.x, {}, {}, image[{}], {};", op, type, coord, value, coord,
                image_id, ImageType(meta.image.type));
        return fmt::format("{}.x", coord);
    }

    template <const std::string_view& op, const std::string_view& type>
    std::string Atomic(Operation operation) {
        std::string temporary = AllocTemporary();
        std::string address;
        std::string_view opname;
        bool robust = false;
        if (const auto gmem = std::get_if<GmemNode>(&*operation[0])) {
            address = GlobalMemoryPointer(*gmem);
            opname = "ATOM";
            robust = true;
        } else if (const auto smem = std::get_if<SmemNode>(&*operation[0])) {
            address = fmt::format("shared_mem[{}]", Visit(smem->GetAddress()));
            opname = "ATOMS";
        } else {
            UNREACHABLE();
            return "{0, 0, 0, 0}";
        }
        if (robust) {
            AddLine("IF NE.x;");
        }
        AddLine("{}.{}.{} {}, {}, {};", opname, op, type, temporary, Visit(operation[1]), address);
        if (robust) {
            AddLine("ELSE;");
            AddLine("MOV.S {}, 0;", temporary);
            AddLine("ENDIF;");
        }
        return temporary;
    }

    template <char type>
    std::string Negate(Operation operation) {
        std::string temporary = AllocTemporary();
        if constexpr (type == 'F') {
            AddLine("MOV.F32 {}, -{};", temporary, Visit(operation[0]));
        } else {
            AddLine("MOV.{} {}, -{};", type, temporary, Visit(operation[0]));
        }
        return temporary;
    }

    template <char type>
    std::string Absolute(Operation operation) {
        std::string temporary = AllocTemporary();
        AddLine("MOV.{} {}, |{}|;", type, temporary, Visit(operation[0]));
        return temporary;
    }

    template <char type>
    std::string BitfieldInsert(Operation operation) {
        const std::string temporary = AllocVectorTemporary();
        AddLine("MOV.{} {}.x, {};", type, temporary, Visit(operation[3]));
        AddLine("MOV.{} {}.y, {};", type, temporary, Visit(operation[2]));
        AddLine("BFI.{} {}.x, {}, {}, {};", type, temporary, temporary, Visit(operation[1]),
                Visit(operation[0]));
        return fmt::format("{}.x", temporary);
    }

    template <char type>
    std::string BitfieldExtract(Operation operation) {
        const std::string temporary = AllocVectorTemporary();
        AddLine("MOV.{} {}.x, {};", type, temporary, Visit(operation[2]));
        AddLine("MOV.{} {}.y, {};", type, temporary, Visit(operation[1]));
        AddLine("BFE.{} {}.x, {}, {};", type, temporary, temporary, Visit(operation[0]));
        return fmt::format("{}.x", temporary);
    }

    template <char swizzle>
    std::string LocalInvocationId(Operation) {
        return fmt::format("invocation.localid.{}", swizzle);
    }

    template <char swizzle>
    std::string WorkGroupId(Operation) {
        return fmt::format("invocation.groupid.{}", swizzle);
    }

    template <char c1, char c2>
    std::string ThreadMask(Operation) {
        return fmt::format("{}.thread{}{}mask", StageInputName(stage), c1, c2);
    }

    template <typename... Args>
    void AddExpression(std::string_view text, Args&&... args) {
        shader_source += fmt::format(text, std::forward<Args>(args)...);
    }

    template <typename... Args>
    void AddLine(std::string_view text, Args&&... args) {
        AddExpression(text, std::forward<Args>(args)...);
        shader_source += '\n';
    }

    std::string AllocLongVectorTemporary() {
        max_long_temporaries = std::max(max_long_temporaries, num_long_temporaries + 1);
        return fmt::format("L{}", num_long_temporaries++);
    }

    std::string AllocLongTemporary() {
        return fmt::format("{}.x", AllocLongVectorTemporary());
    }

    std::string AllocVectorTemporary() {
        max_temporaries = std::max(max_temporaries, num_temporaries + 1);
        return fmt::format("T{}", num_temporaries++);
    }

    std::string AllocTemporary() {
        return fmt::format("{}.x", AllocVectorTemporary());
    }

    void ResetTemporaries() noexcept {
        num_temporaries = 0;
        num_long_temporaries = 0;
    }

    const Device& device;
    const ShaderIR& ir;
    const Registry& registry;
    const ShaderType stage;

    std::size_t num_temporaries = 0;
    std::size_t max_temporaries = 0;

    std::size_t num_long_temporaries = 0;
    std::size_t max_long_temporaries = 0;

    std::map<GlobalMemoryBase, u32> global_memory_names;

    std::string shader_source;

    static constexpr std::string_view ADD_F32 = "ADD.F32";
    static constexpr std::string_view ADD_S = "ADD.S";
    static constexpr std::string_view ADD_U = "ADD.U";
    static constexpr std::string_view MUL_F32 = "MUL.F32";
    static constexpr std::string_view MUL_S = "MUL.S";
    static constexpr std::string_view MUL_U = "MUL.U";
    static constexpr std::string_view DIV_F32 = "DIV.F32";
    static constexpr std::string_view DIV_S = "DIV.S";
    static constexpr std::string_view DIV_U = "DIV.U";
    static constexpr std::string_view MAD_F32 = "MAD.F32";
    static constexpr std::string_view RSQ_F32 = "RSQ.F32";
    static constexpr std::string_view COS_F32 = "COS.F32";
    static constexpr std::string_view SIN_F32 = "SIN.F32";
    static constexpr std::string_view EX2_F32 = "EX2.F32";
    static constexpr std::string_view LG2_F32 = "LG2.F32";
    static constexpr std::string_view SLT_F = "SLT.F32";
    static constexpr std::string_view SLT_S = "SLT.S";
    static constexpr std::string_view SLT_U = "SLT.U";
    static constexpr std::string_view SEQ_F = "SEQ.F32";
    static constexpr std::string_view SEQ_S = "SEQ.S";
    static constexpr std::string_view SEQ_U = "SEQ.U";
    static constexpr std::string_view SLE_F = "SLE.F32";
    static constexpr std::string_view SLE_S = "SLE.S";
    static constexpr std::string_view SLE_U = "SLE.U";
    static constexpr std::string_view SGT_F = "SGT.F32";
    static constexpr std::string_view SGT_S = "SGT.S";
    static constexpr std::string_view SGT_U = "SGT.U";
    static constexpr std::string_view SNE_F = "SNE.F32";
    static constexpr std::string_view SNE_S = "SNE.S";
    static constexpr std::string_view SNE_U = "SNE.U";
    static constexpr std::string_view SGE_F = "SGE.F32";
    static constexpr std::string_view SGE_S = "SGE.S";
    static constexpr std::string_view SGE_U = "SGE.U";
    static constexpr std::string_view AND_S = "AND.S";
    static constexpr std::string_view AND_U = "AND.U";
    static constexpr std::string_view TRUNC_F = "TRUNC.F";
    static constexpr std::string_view TRUNC_S = "TRUNC.S";
    static constexpr std::string_view TRUNC_U = "TRUNC.U";
    static constexpr std::string_view SHL_S = "SHL.S";
    static constexpr std::string_view SHL_U = "SHL.U";
    static constexpr std::string_view SHR_S = "SHR.S";
    static constexpr std::string_view SHR_U = "SHR.U";
    static constexpr std::string_view OR_S = "OR.S";
    static constexpr std::string_view OR_U = "OR.U";
    static constexpr std::string_view XOR_S = "XOR.S";
    static constexpr std::string_view XOR_U = "XOR.U";
    static constexpr std::string_view NOT_S = "NOT.S";
    static constexpr std::string_view NOT_U = "NOT.U";
    static constexpr std::string_view BTC_S = "BTC.S";
    static constexpr std::string_view BTC_U = "BTC.U";
    static constexpr std::string_view BTFM_S = "BTFM.S";
    static constexpr std::string_view BTFM_U = "BTFM.U";
    static constexpr std::string_view ROUND_F = "ROUND.F";
    static constexpr std::string_view CEIL_F = "CEIL.F";
    static constexpr std::string_view FLR_F = "FLR.F";
    static constexpr std::string_view I2F_S = "I2F.S";
    static constexpr std::string_view I2F_U = "I2F.U";
    static constexpr std::string_view MIN_F = "MIN.F";
    static constexpr std::string_view MIN_S = "MIN.S";
    static constexpr std::string_view MIN_U = "MIN.U";
    static constexpr std::string_view MAX_F = "MAX.F";
    static constexpr std::string_view MAX_S = "MAX.S";
    static constexpr std::string_view MAX_U = "MAX.U";
    static constexpr std::string_view MOV_U = "MOV.U";
    static constexpr std::string_view TGBALLOT_U = "TGBALLOT.U";
    static constexpr std::string_view TGALL_U = "TGALL.U";
    static constexpr std::string_view TGANY_U = "TGANY.U";
    static constexpr std::string_view TGEQ_U = "TGEQ.U";
    static constexpr std::string_view EXCH = "EXCH";
    static constexpr std::string_view ADD = "ADD";
    static constexpr std::string_view MIN = "MIN";
    static constexpr std::string_view MAX = "MAX";
    static constexpr std::string_view AND = "AND";
    static constexpr std::string_view OR = "OR";
    static constexpr std::string_view XOR = "XOR";
    static constexpr std::string_view U32 = "U32";
    static constexpr std::string_view S32 = "S32";

    static constexpr std::size_t NUM_ENTRIES = static_cast<std::size_t>(OperationCode::Amount);
    using DecompilerType = std::string (ARBDecompiler::*)(Operation);
    static constexpr std::array<DecompilerType, NUM_ENTRIES> OPERATION_DECOMPILERS = {
        &ARBDecompiler::Assign,

        &ARBDecompiler::Select,

        &ARBDecompiler::Binary<ADD_F32>,
        &ARBDecompiler::Binary<MUL_F32>,
        &ARBDecompiler::Binary<DIV_F32>,
        &ARBDecompiler::Trinary<MAD_F32>,
        &ARBDecompiler::Negate<'F'>,
        &ARBDecompiler::Absolute<'F'>,
        &ARBDecompiler::FClamp,
        &ARBDecompiler::FCastHalf0,
        &ARBDecompiler::FCastHalf1,
        &ARBDecompiler::Binary<MIN_F>,
        &ARBDecompiler::Binary<MAX_F>,
        &ARBDecompiler::Unary<COS_F32>,
        &ARBDecompiler::Unary<SIN_F32>,
        &ARBDecompiler::Unary<EX2_F32>,
        &ARBDecompiler::Unary<LG2_F32>,
        &ARBDecompiler::Unary<RSQ_F32>,
        &ARBDecompiler::FSqrt,
        &ARBDecompiler::Unary<ROUND_F>,
        &ARBDecompiler::Unary<FLR_F>,
        &ARBDecompiler::Unary<CEIL_F>,
        &ARBDecompiler::Unary<TRUNC_F>,
        &ARBDecompiler::Unary<I2F_S>,
        &ARBDecompiler::Unary<I2F_U>,
        &ARBDecompiler::FSwizzleAdd,

        &ARBDecompiler::Binary<ADD_S>,
        &ARBDecompiler::Binary<MUL_S>,
        &ARBDecompiler::Binary<DIV_S>,
        &ARBDecompiler::Negate<'S'>,
        &ARBDecompiler::Absolute<'S'>,
        &ARBDecompiler::Binary<MIN_S>,
        &ARBDecompiler::Binary<MAX_S>,

        &ARBDecompiler::Unary<TRUNC_S>,
        &ARBDecompiler::Unary<MOV_U>,
        &ARBDecompiler::Binary<SHL_S>,
        &ARBDecompiler::Binary<SHR_U>,
        &ARBDecompiler::Binary<SHR_S>,
        &ARBDecompiler::Binary<AND_S>,
        &ARBDecompiler::Binary<OR_S>,
        &ARBDecompiler::Binary<XOR_S>,
        &ARBDecompiler::Unary<NOT_S>,
        &ARBDecompiler::BitfieldInsert<'S'>,
        &ARBDecompiler::BitfieldExtract<'S'>,
        &ARBDecompiler::Unary<BTC_S>,
        &ARBDecompiler::Unary<BTFM_S>,

        &ARBDecompiler::Binary<ADD_U>,
        &ARBDecompiler::Binary<MUL_U>,
        &ARBDecompiler::Binary<DIV_U>,
        &ARBDecompiler::Binary<MIN_U>,
        &ARBDecompiler::Binary<MAX_U>,
        &ARBDecompiler::Unary<TRUNC_U>,
        &ARBDecompiler::Unary<MOV_U>,
        &ARBDecompiler::Binary<SHL_U>,
        &ARBDecompiler::Binary<SHR_U>,
        &ARBDecompiler::Binary<SHR_U>,
        &ARBDecompiler::Binary<AND_U>,
        &ARBDecompiler::Binary<OR_U>,
        &ARBDecompiler::Binary<XOR_U>,
        &ARBDecompiler::Unary<NOT_U>,
        &ARBDecompiler::BitfieldInsert<'U'>,
        &ARBDecompiler::BitfieldExtract<'U'>,
        &ARBDecompiler::Unary<BTC_U>,
        &ARBDecompiler::Unary<BTFM_U>,

        &ARBDecompiler::HAdd2,
        &ARBDecompiler::HMul2,
        &ARBDecompiler::HFma2,
        &ARBDecompiler::HAbsolute,
        &ARBDecompiler::HNegate,
        &ARBDecompiler::HClamp,
        &ARBDecompiler::HCastFloat,
        &ARBDecompiler::HUnpack,
        &ARBDecompiler::HMergeF32,
        &ARBDecompiler::HMergeH0,
        &ARBDecompiler::HMergeH1,
        &ARBDecompiler::HPack2,

        &ARBDecompiler::LogicalAssign,
        &ARBDecompiler::Binary<AND_U>,
        &ARBDecompiler::Binary<OR_U>,
        &ARBDecompiler::Binary<XOR_U>,
        &ARBDecompiler::Unary<NOT_U>,
        &ARBDecompiler::LogicalPick2,
        &ARBDecompiler::LogicalAnd2,

        &ARBDecompiler::FloatComparison<SLT_F, false>,
        &ARBDecompiler::FloatComparison<SEQ_F, false>,
        &ARBDecompiler::FloatComparison<SLE_F, false>,
        &ARBDecompiler::FloatComparison<SGT_F, false>,
        &ARBDecompiler::FloatComparison<SNE_F, false>,
        &ARBDecompiler::FloatComparison<SGE_F, false>,
        &ARBDecompiler::FloatOrdered,
        &ARBDecompiler::FloatUnordered,
        &ARBDecompiler::FloatComparison<SLT_F, true>,
        &ARBDecompiler::FloatComparison<SEQ_F, true>,
        &ARBDecompiler::FloatComparison<SLE_F, true>,
        &ARBDecompiler::FloatComparison<SGT_F, true>,
        &ARBDecompiler::FloatComparison<SNE_F, true>,
        &ARBDecompiler::FloatComparison<SGE_F, true>,

        &ARBDecompiler::Binary<SLT_S>,
        &ARBDecompiler::Binary<SEQ_S>,
        &ARBDecompiler::Binary<SLE_S>,
        &ARBDecompiler::Binary<SGT_S>,
        &ARBDecompiler::Binary<SNE_S>,
        &ARBDecompiler::Binary<SGE_S>,

        &ARBDecompiler::Binary<SLT_U>,
        &ARBDecompiler::Binary<SEQ_U>,
        &ARBDecompiler::Binary<SLE_U>,
        &ARBDecompiler::Binary<SGT_U>,
        &ARBDecompiler::Binary<SNE_U>,
        &ARBDecompiler::Binary<SGE_U>,

        &ARBDecompiler::LogicalAddCarry,

        &ARBDecompiler::HalfComparison<SLT_F, false>,
        &ARBDecompiler::HalfComparison<SEQ_F, false>,
        &ARBDecompiler::HalfComparison<SLE_F, false>,
        &ARBDecompiler::HalfComparison<SGT_F, false>,
        &ARBDecompiler::HalfComparison<SNE_F, false>,
        &ARBDecompiler::HalfComparison<SGE_F, false>,
        &ARBDecompiler::HalfComparison<SLT_F, true>,
        &ARBDecompiler::HalfComparison<SEQ_F, true>,
        &ARBDecompiler::HalfComparison<SLE_F, true>,
        &ARBDecompiler::HalfComparison<SGT_F, true>,
        &ARBDecompiler::HalfComparison<SNE_F, true>,
        &ARBDecompiler::HalfComparison<SGE_F, true>,

        &ARBDecompiler::Texture,
        &ARBDecompiler::Texture,
        &ARBDecompiler::TextureGather,
        &ARBDecompiler::TextureQueryDimensions,
        &ARBDecompiler::TextureQueryLod,
        &ARBDecompiler::TexelFetch,
        &ARBDecompiler::TextureGradient,

        &ARBDecompiler::ImageLoad,
        &ARBDecompiler::ImageStore,

        &ARBDecompiler::AtomicImage<ADD, U32>,
        &ARBDecompiler::AtomicImage<AND, U32>,
        &ARBDecompiler::AtomicImage<OR, U32>,
        &ARBDecompiler::AtomicImage<XOR, U32>,
        &ARBDecompiler::AtomicImage<EXCH, U32>,

        &ARBDecompiler::Atomic<EXCH, U32>,
        &ARBDecompiler::Atomic<ADD, U32>,
        &ARBDecompiler::Atomic<MIN, U32>,
        &ARBDecompiler::Atomic<MAX, U32>,
        &ARBDecompiler::Atomic<AND, U32>,
        &ARBDecompiler::Atomic<OR, U32>,
        &ARBDecompiler::Atomic<XOR, U32>,

        &ARBDecompiler::Atomic<EXCH, S32>,
        &ARBDecompiler::Atomic<ADD, S32>,
        &ARBDecompiler::Atomic<MIN, S32>,
        &ARBDecompiler::Atomic<MAX, S32>,
        &ARBDecompiler::Atomic<AND, S32>,
        &ARBDecompiler::Atomic<OR, S32>,
        &ARBDecompiler::Atomic<XOR, S32>,

        &ARBDecompiler::Atomic<ADD, U32>,
        &ARBDecompiler::Atomic<MIN, U32>,
        &ARBDecompiler::Atomic<MAX, U32>,
        &ARBDecompiler::Atomic<AND, U32>,
        &ARBDecompiler::Atomic<OR, U32>,
        &ARBDecompiler::Atomic<XOR, U32>,

        &ARBDecompiler::Atomic<ADD, S32>,
        &ARBDecompiler::Atomic<MIN, S32>,
        &ARBDecompiler::Atomic<MAX, S32>,
        &ARBDecompiler::Atomic<AND, S32>,
        &ARBDecompiler::Atomic<OR, S32>,
        &ARBDecompiler::Atomic<XOR, S32>,

        &ARBDecompiler::Branch,
        &ARBDecompiler::BranchIndirect,
        &ARBDecompiler::PushFlowStack,
        &ARBDecompiler::PopFlowStack,
        &ARBDecompiler::Exit,
        &ARBDecompiler::Discard,

        &ARBDecompiler::EmitVertex,
        &ARBDecompiler::EndPrimitive,

        &ARBDecompiler::InvocationId,
        &ARBDecompiler::YNegate,
        &ARBDecompiler::LocalInvocationId<'x'>,
        &ARBDecompiler::LocalInvocationId<'y'>,
        &ARBDecompiler::LocalInvocationId<'z'>,
        &ARBDecompiler::WorkGroupId<'x'>,
        &ARBDecompiler::WorkGroupId<'y'>,
        &ARBDecompiler::WorkGroupId<'z'>,

        &ARBDecompiler::Unary<TGBALLOT_U>,
        &ARBDecompiler::Unary<TGALL_U>,
        &ARBDecompiler::Unary<TGANY_U>,
        &ARBDecompiler::Unary<TGEQ_U>,

        &ARBDecompiler::ThreadId,
        &ARBDecompiler::ThreadMask<'e', 'q'>,
        &ARBDecompiler::ThreadMask<'g', 'e'>,
        &ARBDecompiler::ThreadMask<'g', 't'>,
        &ARBDecompiler::ThreadMask<'l', 'e'>,
        &ARBDecompiler::ThreadMask<'l', 't'>,
        &ARBDecompiler::ShuffleIndexed,

        &ARBDecompiler::Barrier,
        &ARBDecompiler::MemoryBarrierGroup,
        &ARBDecompiler::MemoryBarrierGlobal,
    };
};

ARBDecompiler::ARBDecompiler(const Device& device_, const ShaderIR& ir_, const Registry& registry_,
                             ShaderType stage_, std::string_view identifier)
    : device{device_}, ir{ir_}, registry{registry_}, stage{stage_} {
    DefineGlobalMemory();

    AddLine("TEMP RC;");
    AddLine("TEMP FSWZA[4];");
    AddLine("TEMP FSWZB[4];");
    if (ir.IsDecompiled()) {
        DecompileAST();
    } else {
        DecompileBranchMode();
    }
    AddLine("END");

    const std::string code = std::move(shader_source);
    DeclareHeader();
    DeclareVertex();
    DeclareGeometry();
    DeclareFragment();
    DeclareCompute();
    DeclareInputAttributes();
    DeclareOutputAttributes();
    DeclareLocalMemory();
    DeclareGlobalMemory();
    DeclareConstantBuffers();
    DeclareRegisters();
    DeclareTemporaries();
    DeclarePredicates();
    DeclareInternalFlags();

    shader_source += code;
}

std::string_view HeaderStageName(ShaderType stage) {
    switch (stage) {
    case ShaderType::Vertex:
        return "vp";
    case ShaderType::Geometry:
        return "gp";
    case ShaderType::Fragment:
        return "fp";
    case ShaderType::Compute:
        return "cp";
    default:
        UNREACHABLE();
        return "";
    }
}

void ARBDecompiler::DefineGlobalMemory() {
    u32 binding = 0;
    for (const auto& pair : ir.GetGlobalMemory()) {
        const GlobalMemoryBase base = pair.first;
        global_memory_names.emplace(base, binding);
        ++binding;
    }
}

void ARBDecompiler::DeclareHeader() {
    AddLine("!!NV{}5.0", HeaderStageName(stage));
    // Enabling this allows us to cheat on some instructions like TXL with SHADOWARRAY2D
    AddLine("OPTION NV_internal;");
    AddLine("OPTION NV_gpu_program_fp64;");
    AddLine("OPTION NV_shader_thread_group;");
    if (ir.UsesWarps() && device.HasWarpIntrinsics()) {
        AddLine("OPTION NV_shader_thread_shuffle;");
    }
    if (stage == ShaderType::Vertex) {
        if (device.HasNvViewportArray2()) {
            AddLine("OPTION NV_viewport_array2;");
        }
    }
    if (stage == ShaderType::Fragment) {
        AddLine("OPTION ARB_draw_buffers;");
    }
    if (device.HasImageLoadFormatted()) {
        AddLine("OPTION EXT_shader_image_load_formatted;");
    }
}

void ARBDecompiler::DeclareVertex() {
    if (stage != ShaderType::Vertex) {
        return;
    }
    AddLine("OUTPUT result_clip[] = {{ result.clip[0..7] }};");
}

void ARBDecompiler::DeclareGeometry() {
    if (stage != ShaderType::Geometry) {
        return;
    }
    const auto& info = registry.GetGraphicsInfo();
    const auto& header = ir.GetHeader();
    AddLine("PRIMITIVE_IN {};", PrimitiveDescription(info.primitive_topology));
    AddLine("PRIMITIVE_OUT {};", TopologyName(header.common3.output_topology));
    AddLine("VERTICES_OUT {};", header.common4.max_output_vertices.Value());
    AddLine("ATTRIB vertex_position = vertex.position;");
}

void ARBDecompiler::DeclareFragment() {
    if (stage != ShaderType::Fragment) {
        return;
    }
    AddLine("OUTPUT result_color7 = result.color[7];");
    AddLine("OUTPUT result_color6 = result.color[6];");
    AddLine("OUTPUT result_color5 = result.color[5];");
    AddLine("OUTPUT result_color4 = result.color[4];");
    AddLine("OUTPUT result_color3 = result.color[3];");
    AddLine("OUTPUT result_color2 = result.color[2];");
    AddLine("OUTPUT result_color1 = result.color[1];");
    AddLine("OUTPUT result_color0 = result.color;");
}

void ARBDecompiler::DeclareCompute() {
    if (stage != ShaderType::Compute) {
        return;
    }
    const ComputeInfo& info = registry.GetComputeInfo();
    AddLine("GROUP_SIZE {} {} {};", info.workgroup_size[0], info.workgroup_size[1],
            info.workgroup_size[2]);
    if (info.shared_memory_size_in_words == 0) {
        return;
    }
    const u32 limit = device.GetMaxComputeSharedMemorySize();
    u32 size_in_bytes = info.shared_memory_size_in_words * 4;
    if (size_in_bytes > limit) {
        LOG_ERROR(Render_OpenGL, "Shared memory size {} is clamped to host's limit {}",
                  size_in_bytes, limit);
        size_in_bytes = limit;
    }

    AddLine("SHARED_MEMORY {};", size_in_bytes);
    AddLine("SHARED shared_mem[] = {{program.sharedmem}};");
}

void ARBDecompiler::DeclareInputAttributes() {
    if (stage == ShaderType::Compute) {
        return;
    }
    const std::string_view stage_name = StageInputName(stage);
    for (const auto attribute : ir.GetInputAttributes()) {
        if (!IsGenericAttribute(attribute)) {
            continue;
        }
        const u32 index = GetGenericAttributeIndex(attribute);

        std::string_view suffix;
        if (stage == ShaderType::Fragment) {
            const auto input_mode{ir.GetHeader().ps.GetPixelImap(index)};
            if (input_mode == PixelImap::Unused) {
                return;
            }
            suffix = GetInputFlags(input_mode);
        }
        AddLine("{}ATTRIB in_attr{}[] = {{ {}.attrib[{}..{}] }};", suffix, index, stage_name, index,
                index);
    }
}

void ARBDecompiler::DeclareOutputAttributes() {
    if (stage == ShaderType::Compute) {
        return;
    }
    for (const auto attribute : ir.GetOutputAttributes()) {
        if (!IsGenericAttribute(attribute)) {
            continue;
        }
        const u32 index = GetGenericAttributeIndex(attribute);
        AddLine("OUTPUT out_attr{}[] = {{ result.attrib[{}..{}] }};", index, index, index);
    }
}

void ARBDecompiler::DeclareLocalMemory() {
    u64 size = 0;
    if (stage == ShaderType::Compute) {
        size = registry.GetComputeInfo().local_memory_size_in_words * 4ULL;
    } else {
        size = ir.GetHeader().GetLocalMemorySize();
    }
    if (size == 0) {
        return;
    }
    const u64 element_count = Common::AlignUp(size, 4) / 4;
    AddLine("TEMP lmem[{}];", element_count);
}

void ARBDecompiler::DeclareGlobalMemory() {
    const size_t num_entries = ir.GetGlobalMemory().size();
    if (num_entries > 0) {
        AddLine("PARAM c[{}] = {{ program.local[0..{}] }};", num_entries, num_entries - 1);
    }
}

void ARBDecompiler::DeclareConstantBuffers() {
    u32 binding = 0;
    for (const auto& cbuf : ir.GetConstantBuffers()) {
        AddLine("CBUFFER cbuf{}[] = {{ program.buffer[{}] }};", cbuf.first, binding);
        ++binding;
    }
}

void ARBDecompiler::DeclareRegisters() {
    for (const u32 gpr : ir.GetRegisters()) {
        AddLine("TEMP R{};", gpr);
    }
}

void ARBDecompiler::DeclareTemporaries() {
    for (std::size_t i = 0; i < max_temporaries; ++i) {
        AddLine("TEMP T{};", i);
    }
    for (std::size_t i = 0; i < max_long_temporaries; ++i) {
        AddLine("LONG TEMP L{};", i);
    }
}

void ARBDecompiler::DeclarePredicates() {
    for (const Tegra::Shader::Pred pred : ir.GetPredicates()) {
        AddLine("TEMP P{};", static_cast<u64>(pred));
    }
}

void ARBDecompiler::DeclareInternalFlags() {
    for (const char* name : INTERNAL_FLAG_NAMES) {
        AddLine("TEMP {};", name);
    }
}

void ARBDecompiler::InitializeVariables() {
    AddLine("MOV.F32 FSWZA[0], -1;");
    AddLine("MOV.F32 FSWZA[1], 1;");
    AddLine("MOV.F32 FSWZA[2], -1;");
    AddLine("MOV.F32 FSWZA[3], 0;");
    AddLine("MOV.F32 FSWZB[0], -1;");
    AddLine("MOV.F32 FSWZB[1], -1;");
    AddLine("MOV.F32 FSWZB[2], 1;");
    AddLine("MOV.F32 FSWZB[3], -1;");

    if (stage == ShaderType::Vertex || stage == ShaderType::Geometry) {
        AddLine("MOV.F result.position, {{0, 0, 0, 1}};");
    }
    for (const auto attribute : ir.GetOutputAttributes()) {
        if (!IsGenericAttribute(attribute)) {
            continue;
        }
        const u32 index = GetGenericAttributeIndex(attribute);
        AddLine("MOV.F result.attrib[{}], {{0, 0, 0, 1}};", index);
    }
    for (const u32 gpr : ir.GetRegisters()) {
        AddLine("MOV.F R{}, {{0, 0, 0, 0}};", gpr);
    }
    for (const Tegra::Shader::Pred pred : ir.GetPredicates()) {
        AddLine("MOV.U P{}, {{0, 0, 0, 0}};", static_cast<u64>(pred));
    }
}

void ARBDecompiler::DecompileAST() {
    const u32 num_flow_variables = ir.GetASTNumVariables();
    for (u32 i = 0; i < num_flow_variables; ++i) {
        AddLine("TEMP F{};", i);
    }
    for (u32 i = 0; i < num_flow_variables; ++i) {
        AddLine("MOV.U F{}, {{0, 0, 0, 0}};", i);
    }

    InitializeVariables();

    VisitAST(ir.GetASTProgram());
}

void ARBDecompiler::DecompileBranchMode() {
    static constexpr u32 FLOW_STACK_SIZE = 20;
    if (!ir.IsFlowStackDisabled()) {
        AddLine("TEMP SSY[{}];", FLOW_STACK_SIZE);
        AddLine("TEMP PBK[{}];", FLOW_STACK_SIZE);
        AddLine("TEMP SSY_TOP;");
        AddLine("TEMP PBK_TOP;");
    }

    AddLine("TEMP PC;");

    if (!ir.IsFlowStackDisabled()) {
        AddLine("MOV.U SSY_TOP.x, 0;");
        AddLine("MOV.U PBK_TOP.x, 0;");
    }

    InitializeVariables();

    const auto basic_block_end = ir.GetBasicBlocks().end();
    auto basic_block_it = ir.GetBasicBlocks().begin();
    const u32 first_address = basic_block_it->first;
    AddLine("MOV.U PC.x, {};", first_address);

    AddLine("REP;");

    std::size_t num_blocks = 0;
    while (basic_block_it != basic_block_end) {
        const auto& [address, bb] = *basic_block_it;
        ++num_blocks;

        AddLine("SEQ.S.CC RC.x, PC.x, {};", address);
        AddLine("IF NE.x;");

        VisitBlock(bb);

        ++basic_block_it;

        if (basic_block_it != basic_block_end) {
            const auto op = std::get_if<OperationNode>(&*bb[bb.size() - 1]);
            if (!op || op->GetCode() != OperationCode::Branch) {
                const u32 next_address = basic_block_it->first;
                AddLine("MOV.U PC.x, {};", next_address);
                AddLine("CONT;");
            }
        }

        AddLine("ELSE;");
    }
    AddLine("RET;");
    while (num_blocks--) {
        AddLine("ENDIF;");
    }

    AddLine("ENDREP;");
}

void ARBDecompiler::VisitAST(const ASTNode& node) {
    if (const auto ast = std::get_if<ASTProgram>(&*node->GetInnerData())) {
        for (ASTNode current = ast->nodes.GetFirst(); current; current = current->GetNext()) {
            VisitAST(current);
        }
    } else if (const auto if_then = std::get_if<ASTIfThen>(&*node->GetInnerData())) {
        const std::string condition = VisitExpression(if_then->condition);
        ResetTemporaries();

        AddLine("MOVC.U RC.x, {};", condition);
        AddLine("IF NE.x;");
        for (ASTNode current = if_then->nodes.GetFirst(); current; current = current->GetNext()) {
            VisitAST(current);
        }
        AddLine("ENDIF;");
    } else if (const auto if_else = std::get_if<ASTIfElse>(&*node->GetInnerData())) {
        AddLine("ELSE;");
        for (ASTNode current = if_else->nodes.GetFirst(); current; current = current->GetNext()) {
            VisitAST(current);
        }
    } else if (const auto decoded = std::get_if<ASTBlockDecoded>(&*node->GetInnerData())) {
        VisitBlock(decoded->nodes);
    } else if (const auto var_set = std::get_if<ASTVarSet>(&*node->GetInnerData())) {
        AddLine("MOV.U F{}, {};", var_set->index, VisitExpression(var_set->condition));
        ResetTemporaries();
    } else if (const auto do_while = std::get_if<ASTDoWhile>(&*node->GetInnerData())) {
        const std::string condition = VisitExpression(do_while->condition);
        ResetTemporaries();
        AddLine("REP;");
        for (ASTNode current = do_while->nodes.GetFirst(); current; current = current->GetNext()) {
            VisitAST(current);
        }
        AddLine("MOVC.U RC.x, {};", condition);
        AddLine("BRK (NE.x);");
        AddLine("ENDREP;");
    } else if (const auto ast_return = std::get_if<ASTReturn>(&*node->GetInnerData())) {
        const bool is_true = ExprIsTrue(ast_return->condition);
        if (!is_true) {
            AddLine("MOVC.U RC.x, {};", VisitExpression(ast_return->condition));
            AddLine("IF NE.x;");
            ResetTemporaries();
        }
        if (ast_return->kills) {
            AddLine("KIL TR;");
        } else {
            Exit();
        }
        if (!is_true) {
            AddLine("ENDIF;");
        }
    } else if (const auto ast_break = std::get_if<ASTBreak>(&*node->GetInnerData())) {
        if (ExprIsTrue(ast_break->condition)) {
            AddLine("BRK;");
        } else {
            AddLine("MOVC.U RC.x, {};", VisitExpression(ast_break->condition));
            AddLine("BRK (NE.x);");
            ResetTemporaries();
        }
    } else if (std::holds_alternative<ASTLabel>(*node->GetInnerData())) {
        // Nothing to do
    } else {
        UNREACHABLE();
    }
}

std::string ARBDecompiler::VisitExpression(const Expr& node) {
    if (const auto expr = std::get_if<ExprAnd>(&*node)) {
        std::string result = AllocTemporary();
        AddLine("AND.U {}, {}, {};", result, VisitExpression(expr->operand1),
                VisitExpression(expr->operand2));
        return result;
    }
    if (const auto expr = std::get_if<ExprOr>(&*node)) {
        std::string result = AllocTemporary();
        AddLine("OR.U {}, {}, {};", result, VisitExpression(expr->operand1),
                VisitExpression(expr->operand2));
        return result;
    }
    if (const auto expr = std::get_if<ExprNot>(&*node)) {
        std::string result = AllocTemporary();
        AddLine("CMP.S {}, {}, 0, -1;", result, VisitExpression(expr->operand1));
        return result;
    }
    if (const auto expr = std::get_if<ExprPredicate>(&*node)) {
        return fmt::format("P{}.x", static_cast<u64>(expr->predicate));
    }
    if (const auto expr = std::get_if<ExprCondCode>(&*node)) {
        return Visit(ir.GetConditionCode(expr->cc));
    }
    if (const auto expr = std::get_if<ExprVar>(&*node)) {
        return fmt::format("F{}.x", expr->var_index);
    }
    if (const auto expr = std::get_if<ExprBoolean>(&*node)) {
        return expr->value ? "0xffffffff" : "0";
    }
    if (const auto expr = std::get_if<ExprGprEqual>(&*node)) {
        std::string result = AllocTemporary();
        AddLine("SEQ.U {}, R{}.x, {};", result, expr->gpr, expr->value);
        return result;
    }
    UNREACHABLE();
    return "0";
}

void ARBDecompiler::VisitBlock(const NodeBlock& bb) {
    for (const auto& node : bb) {
        Visit(node);
    }
}

std::string ARBDecompiler::Visit(const Node& node) {
    if (const auto operation = std::get_if<OperationNode>(&*node)) {
        if (const auto amend_index = operation->GetAmendIndex()) {
            Visit(ir.GetAmendNode(*amend_index));
        }
        const std::size_t index = static_cast<std::size_t>(operation->GetCode());
        if (index >= OPERATION_DECOMPILERS.size()) {
            UNREACHABLE_MSG("Out of bounds operation: {}", index);
            return {};
        }
        const auto decompiler = OPERATION_DECOMPILERS[index];
        if (decompiler == nullptr) {
            UNREACHABLE_MSG("Undefined operation: {}", index);
            return {};
        }
        return (this->*decompiler)(*operation);
    }

    if (const auto gpr = std::get_if<GprNode>(&*node)) {
        const u32 index = gpr->GetIndex();
        if (index == Register::ZeroIndex) {
            return "{0, 0, 0, 0}.x";
        }
        return fmt::format("R{}.x", index);
    }

    if (const auto cv = std::get_if<CustomVarNode>(&*node)) {
        return fmt::format("CV{}.x", cv->GetIndex());
    }

    if (const auto immediate = std::get_if<ImmediateNode>(&*node)) {
        std::string temporary = AllocTemporary();
        AddLine("MOV.U {}, {};", temporary, immediate->GetValue());
        return temporary;
    }

    if (const auto predicate = std::get_if<PredicateNode>(&*node)) {
        std::string temporary = AllocTemporary();
        switch (const auto index = predicate->GetIndex(); index) {
        case Tegra::Shader::Pred::UnusedIndex:
            AddLine("MOV.S {}, -1;", temporary);
            break;
        case Tegra::Shader::Pred::NeverExecute:
            AddLine("MOV.S {}, 0;", temporary);
            break;
        default:
            AddLine("MOV.S {}, P{}.x;", temporary, static_cast<u64>(index));
            break;
        }
        if (predicate->IsNegated()) {
            AddLine("CMP.S {}, {}, 0, -1;", temporary, temporary);
        }
        return temporary;
    }

    if (const auto abuf = std::get_if<AbufNode>(&*node)) {
        if (abuf->IsPhysicalBuffer()) {
            UNIMPLEMENTED_MSG("Physical buffers are not implemented");
            return "{0, 0, 0, 0}.x";
        }

        const Attribute::Index index = abuf->GetIndex();
        const u32 element = abuf->GetElement();
        const char swizzle = Swizzle(element);
        switch (index) {
        case Attribute::Index::Position: {
            if (stage == ShaderType::Geometry) {
                return fmt::format("{}_position[{}].{}", StageInputName(stage),
                                   Visit(abuf->GetBuffer()), swizzle);
            } else {
                return fmt::format("{}.position.{}", StageInputName(stage), swizzle);
            }
        }
        case Attribute::Index::TessCoordInstanceIDVertexID:
            ASSERT(stage == ShaderType::Vertex);
            switch (element) {
            case 2:
                return "vertex.instance";
            case 3:
                return "vertex.id";
            }
            UNIMPLEMENTED_MSG("Unmanaged TessCoordInstanceIDVertexID element={}", element);
            break;
        case Attribute::Index::PointCoord:
            switch (element) {
            case 0:
                return "fragment.pointcoord.x";
            case 1:
                return "fragment.pointcoord.y";
            }
            UNIMPLEMENTED();
            break;
        case Attribute::Index::FrontFacing: {
            ASSERT(stage == ShaderType::Fragment);
            ASSERT(element == 3);
            const std::string temporary = AllocVectorTemporary();
            AddLine("SGT.S RC.x, fragment.facing, {{0, 0, 0, 0}};");
            AddLine("MOV.U.CC RC.x, -RC;");
            AddLine("MOV.S {}.x, 0;", temporary);
            AddLine("MOV.S {}.x (NE.x), -1;", temporary);
            return fmt::format("{}.x", temporary);
        }
        default:
            if (IsGenericAttribute(index)) {
                if (stage == ShaderType::Geometry) {
                    return fmt::format("in_attr{}[{}][0].{}", GetGenericAttributeIndex(index),
                                       Visit(abuf->GetBuffer()), swizzle);
                } else {
                    return fmt::format("{}.attrib[{}].{}", StageInputName(stage),
                                       GetGenericAttributeIndex(index), swizzle);
                }
            }
            UNIMPLEMENTED_MSG("Unimplemented input attribute={}", index);
            break;
        }
        return "{0, 0, 0, 0}.x";
    }

    if (const auto cbuf = std::get_if<CbufNode>(&*node)) {
        std::string offset_string;
        const auto& offset = cbuf->GetOffset();
        if (const auto imm = std::get_if<ImmediateNode>(&*offset)) {
            offset_string = std::to_string(imm->GetValue());
        } else {
            offset_string = Visit(offset);
        }
        std::string temporary = AllocTemporary();
        AddLine("LDC.F32 {}, cbuf{}[{}];", temporary, cbuf->GetIndex(), offset_string);
        return temporary;
    }

    if (const auto gmem = std::get_if<GmemNode>(&*node)) {
        std::string temporary = AllocTemporary();
        AddLine("MOV {}, 0;", temporary);
        AddLine("LOAD.U32 {} (NE.x), {};", temporary, GlobalMemoryPointer(*gmem));
        return temporary;
    }

    if (const auto lmem = std::get_if<LmemNode>(&*node)) {
        std::string temporary = Visit(lmem->GetAddress());
        AddLine("SHR.U {}, {}, 2;", temporary, temporary);
        AddLine("MOV.U {}, lmem[{}].x;", temporary, temporary);
        return temporary;
    }

    if (const auto smem = std::get_if<SmemNode>(&*node)) {
        std::string temporary = Visit(smem->GetAddress());
        AddLine("LDS.U32 {}, shared_mem[{}];", temporary, temporary);
        return temporary;
    }

    if (const auto internal_flag = std::get_if<InternalFlagNode>(&*node)) {
        const std::size_t index = static_cast<std::size_t>(internal_flag->GetFlag());
        return fmt::format("{}.x", INTERNAL_FLAG_NAMES[index]);
    }

    if (const auto conditional = std::get_if<ConditionalNode>(&*node)) {
        if (const auto amend_index = conditional->GetAmendIndex()) {
            Visit(ir.GetAmendNode(*amend_index));
        }
        AddLine("MOVC.U RC.x, {};", Visit(conditional->GetCondition()));
        AddLine("IF NE.x;");
        VisitBlock(conditional->GetCode());
        AddLine("ENDIF;");
        return {};
    }

    if ([[maybe_unused]] const auto cmt = std::get_if<CommentNode>(&*node)) {
        // Uncommenting this will generate invalid code. GLASM lacks comments.
        // AddLine("// {}", cmt->GetText());
        return {};
    }

    UNIMPLEMENTED();
    return {};
}

std::tuple<std::string, std::string, std::size_t> ARBDecompiler::BuildCoords(Operation operation) {
    const auto& meta = std::get<MetaTexture>(operation.GetMeta());
    UNIMPLEMENTED_IF(meta.sampler.is_indexed);

    const bool is_extended = meta.sampler.is_shadow && meta.sampler.is_array &&
                             meta.sampler.type == Tegra::Shader::TextureType::TextureCube;
    const std::size_t count = operation.GetOperandsCount();
    std::string temporary = AllocVectorTemporary();
    std::size_t i = 0;
    for (; i < count; ++i) {
        AddLine("MOV.F {}.{}, {};", temporary, Swizzle(i), Visit(operation[i]));
    }
    if (meta.sampler.is_array) {
        AddLine("I2F.S {}.{}, {};", temporary, Swizzle(i), Visit(meta.array));
        ++i;
    }
    if (meta.sampler.is_shadow) {
        std::string compare = Visit(meta.depth_compare);
        if (is_extended) {
            ASSERT(i == 4);
            std::string extra_coord = AllocVectorTemporary();
            AddLine("MOV.F {}.x, {};", extra_coord, compare);
            return {fmt::format("{}, {}", temporary, extra_coord), extra_coord, 0};
        }
        AddLine("MOV.F {}.{}, {};", temporary, Swizzle(i), compare);
        ++i;
    }
    return {temporary, temporary, i};
}

std::string ARBDecompiler::BuildAoffi(Operation operation) {
    const auto& meta = std::get<MetaTexture>(operation.GetMeta());
    if (meta.aoffi.empty()) {
        return {};
    }
    const std::string temporary = AllocVectorTemporary();
    std::size_t i = 0;
    for (auto& node : meta.aoffi) {
        AddLine("MOV.S {}.{}, {};", temporary, Swizzle(i++), Visit(node));
    }
    return fmt::format(", offset({})", temporary);
}

std::string ARBDecompiler::GlobalMemoryPointer(const GmemNode& gmem) {
    // Read a bindless SSBO, return its address and set CC accordingly
    // address = c[binding].xy
    // length  = c[binding].z
    const u32 binding = global_memory_names.at(gmem.GetDescriptor());

    const std::string pointer = AllocLongVectorTemporary();
    std::string temporary = AllocTemporary();

    AddLine("PK64.U {}, c[{}];", pointer, binding);
    AddLine("SUB.U {}, {}, {};", temporary, Visit(gmem.GetRealAddress()),
            Visit(gmem.GetBaseAddress()));
    AddLine("CVT.U64.U32 {}.z, {};", pointer, temporary);
    AddLine("ADD.U64 {}.x, {}.x, {}.z;", pointer, pointer, pointer);
    // Compare offset to length and set CC
    AddLine("SLT.U.CC RC.x, {}, c[{}].z;", temporary, binding);
    return fmt::format("{}.x", pointer);
}

void ARBDecompiler::Exit() {
    if (stage != ShaderType::Fragment) {
        AddLine("RET;");
        return;
    }

    const auto safe_get_register = [this](u32 reg) -> std::string {
        if (ir.GetRegisters().contains(reg)) {
            return fmt::format("R{}.x", reg);
        }
        return "{0, 0, 0, 0}.x";
    };

    const auto& header = ir.GetHeader();
    u32 current_reg = 0;
    for (u32 rt = 0; rt < Tegra::Engines::Maxwell3D::Regs::NumRenderTargets; ++rt) {
        for (u32 component = 0; component < 4; ++component) {
            if (!header.ps.IsColorComponentOutputEnabled(rt, component)) {
                continue;
            }
            AddLine("MOV.F result_color{}.{}, {};", rt, Swizzle(component),
                    safe_get_register(current_reg));
            ++current_reg;
        }
    }
    if (header.ps.omap.depth) {
        AddLine("MOV.F result.depth.z, {};", safe_get_register(current_reg + 1));
    }

    AddLine("RET;");
}

std::string ARBDecompiler::Assign(Operation operation) {
    const Node& dest = operation[0];
    const Node& src = operation[1];

    std::string dest_name;
    if (const auto gpr = std::get_if<GprNode>(&*dest)) {
        if (gpr->GetIndex() == Register::ZeroIndex) {
            // Writing to Register::ZeroIndex is a no op
            return {};
        }
        dest_name = fmt::format("R{}.x", gpr->GetIndex());
    } else if (const auto abuf = std::get_if<AbufNode>(&*dest)) {
        const u32 element = abuf->GetElement();
        const char swizzle = Swizzle(element);
        switch (const Attribute::Index index = abuf->GetIndex()) {
        case Attribute::Index::Position:
            dest_name = fmt::format("result.position.{}", swizzle);
            break;
        case Attribute::Index::LayerViewportPointSize:
            switch (element) {
            case 0:
                UNIMPLEMENTED();
                return {};
            case 1:
            case 2:
                if (!device.HasNvViewportArray2()) {
                    LOG_ERROR(
                        Render_OpenGL,
                        "NV_viewport_array2 is missing. Maxwell gen 2 or better is required.");
                    return {};
                }
                dest_name = element == 1 ? "result.layer.x" : "result.viewport.x";
                break;
            case 3:
                dest_name = "result.pointsize.x";
                break;
            }
            break;
        case Attribute::Index::ClipDistances0123:
            dest_name = fmt::format("result.clip[{}].x", element);
            break;
        case Attribute::Index::ClipDistances4567:
            dest_name = fmt::format("result.clip[{}].x", element + 4);
            break;
        default:
            if (!IsGenericAttribute(index)) {
                UNREACHABLE();
                return {};
            }
            dest_name =
                fmt::format("result.attrib[{}].{}", GetGenericAttributeIndex(index), swizzle);
            break;
        }
    } else if (const auto lmem = std::get_if<LmemNode>(&*dest)) {
        const std::string address = Visit(lmem->GetAddress());
        AddLine("SHR.U {}, {}, 2;", address, address);
        dest_name = fmt::format("lmem[{}].x", address);
    } else if (const auto smem = std::get_if<SmemNode>(&*dest)) {
        AddLine("STS.U32 {}, shared_mem[{}];", Visit(src), Visit(smem->GetAddress()));
        ResetTemporaries();
        return {};
    } else if (const auto gmem = std::get_if<GmemNode>(&*dest)) {
        AddLine("IF NE.x;");
        AddLine("STORE.U32 {}, {};", Visit(src), GlobalMemoryPointer(*gmem));
        AddLine("ENDIF;");
        ResetTemporaries();
        return {};
    } else {
        UNREACHABLE();
        ResetTemporaries();
        return {};
    }

    AddLine("MOV.U {}, {};", dest_name, Visit(src));
    ResetTemporaries();
    return {};
}

std::string ARBDecompiler::Select(Operation operation) {
    std::string temporary = AllocTemporary();
    AddLine("CMP.S {}, {}, {}, {};", temporary, Visit(operation[0]), Visit(operation[1]),
            Visit(operation[2]));
    return temporary;
}

std::string ARBDecompiler::FClamp(Operation operation) {
    // 1.0f in hex, replace with std::bit_cast on C++20
    static constexpr u32 POSITIVE_ONE = 0x3f800000;

    std::string temporary = AllocTemporary();
    const Node& value = operation[0];
    const Node& low = operation[1];
    const Node& high = operation[2];
    const auto* const imm_low = std::get_if<ImmediateNode>(&*low);
    const auto* const imm_high = std::get_if<ImmediateNode>(&*high);
    if (imm_low && imm_high && imm_low->GetValue() == 0 && imm_high->GetValue() == POSITIVE_ONE) {
        AddLine("MOV.F32.SAT {}, {};", temporary, Visit(value));
    } else {
        AddLine("MIN.F {}, {}, {};", temporary, Visit(value), Visit(high));
        AddLine("MAX.F {}, {}, {};", temporary, temporary, Visit(low));
    }
    return temporary;
}

std::string ARBDecompiler::FCastHalf0(Operation operation) {
    const std::string temporary = AllocVectorTemporary();
    AddLine("UP2H.F {}.x, {};", temporary, Visit(operation[0]));
    return fmt::format("{}.x", temporary);
}

std::string ARBDecompiler::FCastHalf1(Operation operation) {
    const std::string temporary = AllocVectorTemporary();
    AddLine("UP2H.F {}.y, {};", temporary, Visit(operation[0]));
    AddLine("MOV {}.x, {}.y;", temporary, temporary);
    return fmt::format("{}.x", temporary);
}

std::string ARBDecompiler::FSqrt(Operation operation) {
    std::string temporary = AllocTemporary();
    AddLine("RSQ.F32 {}, {};", temporary, Visit(operation[0]));
    AddLine("RCP.F32 {}, {};", temporary, temporary);
    return temporary;
}

std::string ARBDecompiler::FSwizzleAdd(Operation operation) {
    const std::string temporary = AllocVectorTemporary();
    if (!device.HasWarpIntrinsics()) {
        LOG_ERROR(Render_OpenGL,
                  "NV_shader_thread_shuffle is missing. Kepler or better is required.");
        AddLine("ADD.F {}.x, {}, {};", temporary, Visit(operation[0]), Visit(operation[1]));
        return fmt::format("{}.x", temporary);
    }

    AddLine("AND.U {}.z, {}.threadid, 3;", temporary, StageInputName(stage));
    AddLine("SHL.U {}.z, {}.z, 1;", temporary, temporary);
    AddLine("SHR.U {}.z, {}, {}.z;", temporary, Visit(operation[2]), temporary);
    AddLine("AND.U {}.z, {}.z, 3;", temporary, temporary);
    AddLine("MUL.F32 {}.x, {}, FSWZA[{}.z];", temporary, Visit(operation[0]), temporary);
    AddLine("MUL.F32 {}.y, {}, FSWZB[{}.z];", temporary, Visit(operation[1]), temporary);
    AddLine("ADD.F32 {}.x, {}.x, {}.y;", temporary, temporary, temporary);
    return fmt::format("{}.x", temporary);
}

std::string ARBDecompiler::HAdd2(Operation operation) {
    const std::string tmp1 = AllocVectorTemporary();
    const std::string tmp2 = AllocVectorTemporary();
    AddLine("UP2H.F {}.xy, {};", tmp1, Visit(operation[0]));
    AddLine("UP2H.F {}.xy, {};", tmp2, Visit(operation[1]));
    AddLine("ADD.F16 {}, {}, {};", tmp1, tmp1, tmp2);
    AddLine("PK2H.F {}.x, {};", tmp1, tmp1);
    return fmt::format("{}.x", tmp1);
}

std::string ARBDecompiler::HMul2(Operation operation) {
    const std::string tmp1 = AllocVectorTemporary();
    const std::string tmp2 = AllocVectorTemporary();
    AddLine("UP2H.F {}.xy, {};", tmp1, Visit(operation[0]));
    AddLine("UP2H.F {}.xy, {};", tmp2, Visit(operation[1]));
    AddLine("MUL.F16 {}, {}, {};", tmp1, tmp1, tmp2);
    AddLine("PK2H.F {}.x, {};", tmp1, tmp1);
    return fmt::format("{}.x", tmp1);
}

std::string ARBDecompiler::HFma2(Operation operation) {
    const std::string tmp1 = AllocVectorTemporary();
    const std::string tmp2 = AllocVectorTemporary();
    const std::string tmp3 = AllocVectorTemporary();
    AddLine("UP2H.F {}.xy, {};", tmp1, Visit(operation[0]));
    AddLine("UP2H.F {}.xy, {};", tmp2, Visit(operation[1]));
    AddLine("UP2H.F {}.xy, {};", tmp3, Visit(operation[2]));
    AddLine("MAD.F16 {}, {}, {}, {};", tmp1, tmp1, tmp2, tmp3);
    AddLine("PK2H.F {}.x, {};", tmp1, tmp1);
    return fmt::format("{}.x", tmp1);
}

std::string ARBDecompiler::HAbsolute(Operation operation) {
    const std::string temporary = AllocVectorTemporary();
    AddLine("UP2H.F {}.xy, {};", temporary, Visit(operation[0]));
    AddLine("PK2H.F {}.x, |{}|;", temporary, temporary);
    return fmt::format("{}.x", temporary);
}

std::string ARBDecompiler::HNegate(Operation operation) {
    const std::string temporary = AllocVectorTemporary();
    AddLine("UP2H.F {}.xy, {};", temporary, Visit(operation[0]));
    AddLine("MOVC.S RC.x, {};", Visit(operation[1]));
    AddLine("MOV.F {}.x (NE.x), -{}.x;", temporary, temporary);
    AddLine("MOVC.S RC.x, {};", Visit(operation[2]));
    AddLine("MOV.F {}.y (NE.x), -{}.y;", temporary, temporary);
    AddLine("PK2H.F {}.x, {};", temporary, temporary);
    return fmt::format("{}.x", temporary);
}

std::string ARBDecompiler::HClamp(Operation operation) {
    const std::string tmp1 = AllocVectorTemporary();
    const std::string tmp2 = AllocVectorTemporary();
    AddLine("UP2H.F {}.xy, {};", tmp1, Visit(operation[0]));
    AddLine("MOV.U {}.x, {};", tmp2, Visit(operation[1]));
    AddLine("MOV.U {}.y, {}.x;", tmp2, tmp2);
    AddLine("MAX.F {}, {}, {};", tmp1, tmp1, tmp2);
    AddLine("MOV.U {}.x, {};", tmp2, Visit(operation[2]));
    AddLine("MOV.U {}.y, {}.x;", tmp2, tmp2);
    AddLine("MIN.F {}, {}, {};", tmp1, tmp1, tmp2);
    AddLine("PK2H.F {}.x, {};", tmp1, tmp1);
    return fmt::format("{}.x", tmp1);
}

std::string ARBDecompiler::HCastFloat(Operation operation) {
    const std::string temporary = AllocVectorTemporary();
    AddLine("MOV.F {}.y, {{0, 0, 0, 0}};", temporary);
    AddLine("MOV.F {}.x, {};", temporary, Visit(operation[0]));
    AddLine("PK2H.F {}.x, {};", temporary, temporary);
    return fmt::format("{}.x", temporary);
}

std::string ARBDecompiler::HUnpack(Operation operation) {
    std::string operand = Visit(operation[0]);
    switch (std::get<Tegra::Shader::HalfType>(operation.GetMeta())) {
    case Tegra::Shader::HalfType::H0_H1:
        return operand;
    case Tegra::Shader::HalfType::F32: {
        const std::string temporary = AllocVectorTemporary();
        AddLine("MOV.U {}.x, {};", temporary, operand);
        AddLine("MOV.U {}.y, {}.x;", temporary, temporary);
        AddLine("PK2H.F {}.x, {};", temporary, temporary);
        return fmt::format("{}.x", temporary);
    }
    case Tegra::Shader::HalfType::H0_H0: {
        const std::string temporary = AllocVectorTemporary();
        AddLine("UP2H.F {}.xy, {};", temporary, operand);
        AddLine("MOV.U {}.y, {}.x;", temporary, temporary);
        AddLine("PK2H.F {}.x, {};", temporary, temporary);
        return fmt::format("{}.x", temporary);
    }
    case Tegra::Shader::HalfType::H1_H1: {
        const std::string temporary = AllocVectorTemporary();
        AddLine("UP2H.F {}.xy, {};", temporary, operand);
        AddLine("MOV.U {}.x, {}.y;", temporary, temporary);
        AddLine("PK2H.F {}.x, {};", temporary, temporary);
        return fmt::format("{}.x", temporary);
    }
    }
    UNREACHABLE();
    return "{0, 0, 0, 0}.x";
}

std::string ARBDecompiler::HMergeF32(Operation operation) {
    const std::string temporary = AllocVectorTemporary();
    AddLine("UP2H.F {}.xy, {};", temporary, Visit(operation[0]));
    return fmt::format("{}.x", temporary);
}

std::string ARBDecompiler::HMergeH0(Operation operation) {
    const std::string temporary = AllocVectorTemporary();
    AddLine("UP2H.F {}.xy, {};", temporary, Visit(operation[0]));
    AddLine("UP2H.F {}.zw, {};", temporary, Visit(operation[1]));
    AddLine("MOV.U {}.x, {}.z;", temporary, temporary);
    AddLine("PK2H.F {}.x, {};", temporary, temporary);
    return fmt::format("{}.x", temporary);
}

std::string ARBDecompiler::HMergeH1(Operation operation) {
    const std::string temporary = AllocVectorTemporary();
    AddLine("UP2H.F {}.xy, {};", temporary, Visit(operation[0]));
    AddLine("UP2H.F {}.zw, {};", temporary, Visit(operation[1]));
    AddLine("MOV.U {}.y, {}.w;", temporary, temporary);
    AddLine("PK2H.F {}.x, {};", temporary, temporary);
    return fmt::format("{}.x", temporary);
}

std::string ARBDecompiler::HPack2(Operation operation) {
    const std::string temporary = AllocVectorTemporary();
    AddLine("MOV.U {}.x, {};", temporary, Visit(operation[0]));
    AddLine("MOV.U {}.y, {};", temporary, Visit(operation[1]));
    AddLine("PK2H.F {}.x, {};", temporary, temporary);
    return fmt::format("{}.x", temporary);
}

std::string ARBDecompiler::LogicalAssign(Operation operation) {
    const Node& dest = operation[0];
    const Node& src = operation[1];

    std::string target;

    if (const auto pred = std::get_if<PredicateNode>(&*dest)) {
        ASSERT_MSG(!pred->IsNegated(), "Negating logical assignment");

        const Tegra::Shader::Pred index = pred->GetIndex();
        switch (index) {
        case Tegra::Shader::Pred::NeverExecute:
        case Tegra::Shader::Pred::UnusedIndex:
            // Writing to these predicates is a no-op
            return {};
        }
        target = fmt::format("P{}.x", static_cast<u64>(index));
    } else if (const auto internal_flag = std::get_if<InternalFlagNode>(&*dest)) {
        const std::size_t index = static_cast<std::size_t>(internal_flag->GetFlag());
        target = fmt::format("{}.x", INTERNAL_FLAG_NAMES[index]);
    } else {
        UNREACHABLE();
        ResetTemporaries();
        return {};
    }

    AddLine("MOV.U {}, {};", target, Visit(src));
    ResetTemporaries();
    return {};
}

std::string ARBDecompiler::LogicalPick2(Operation operation) {
    std::string temporary = AllocTemporary();
    const u32 index = std::get<ImmediateNode>(*operation[1]).GetValue();
    AddLine("MOV.U {}, {}.{};", temporary, Visit(operation[0]), Swizzle(index));
    return temporary;
}

std::string ARBDecompiler::LogicalAnd2(Operation operation) {
    std::string temporary = AllocTemporary();
    const std::string op = Visit(operation[0]);
    AddLine("AND.U {}, {}.x, {}.y;", temporary, op, op);
    return temporary;
}

std::string ARBDecompiler::FloatOrdered(Operation operation) {
    std::string temporary = AllocTemporary();
    AddLine("MOVC.F32 RC.x, {};", Visit(operation[0]));
    AddLine("MOVC.F32 RC.y, {};", Visit(operation[1]));
    AddLine("MOV.S {}, -1;", temporary);
    AddLine("MOV.S {} (NAN.x), 0;", temporary);
    AddLine("MOV.S {} (NAN.y), 0;", temporary);
    return temporary;
}

std::string ARBDecompiler::FloatUnordered(Operation operation) {
    std::string temporary = AllocTemporary();
    AddLine("MOVC.F32 RC.x, {};", Visit(operation[0]));
    AddLine("MOVC.F32 RC.y, {};", Visit(operation[1]));
    AddLine("MOV.S {}, 0;", temporary);
    AddLine("MOV.S {} (NAN.x), -1;", temporary);
    AddLine("MOV.S {} (NAN.y), -1;", temporary);
    return temporary;
}

std::string ARBDecompiler::LogicalAddCarry(Operation operation) {
    std::string temporary = AllocTemporary();
    AddLine("ADDC.U RC, {}, {};", Visit(operation[0]), Visit(operation[1]));
    AddLine("MOV.S {}, 0;", temporary);
    AddLine("IF CF.x;");
    AddLine("MOV.S {}, -1;", temporary);
    AddLine("ENDIF;");
    return temporary;
}

std::string ARBDecompiler::Texture(Operation operation) {
    const auto& meta = std::get<MetaTexture>(operation.GetMeta());
    const u32 sampler_id = device.GetBaseBindings(stage).sampler + meta.sampler.index;
    const auto [coords, temporary, swizzle] = BuildCoords(operation);

    std::string_view opcode = "TEX";
    std::string extra;
    if (meta.bias) {
        ASSERT(!meta.lod);
        opcode = "TXB";

        if (swizzle < 4) {
            AddLine("MOV.F {}.w, {};", temporary, Visit(meta.bias));
        } else {
            const std::string bias = AllocTemporary();
            AddLine("MOV.F {}, {};", bias, Visit(meta.bias));
            extra = fmt::format(" {},", bias);
        }
    }
    if (meta.lod) {
        ASSERT(!meta.bias);
        opcode = "TXL";

        if (swizzle < 4) {
            AddLine("MOV.F {}.w, {};", temporary, Visit(meta.lod));
        } else {
            const std::string lod = AllocTemporary();
            AddLine("MOV.F {}, {};", lod, Visit(meta.lod));
            extra = fmt::format(" {},", lod);
        }
    }

    AddLine("{}.F {}, {},{} texture[{}], {}{};", opcode, temporary, coords, extra, sampler_id,
            TextureType(meta), BuildAoffi(operation));
    AddLine("MOV.U {}.x, {}.{};", temporary, temporary, Swizzle(meta.element));
    return fmt::format("{}.x", temporary);
}

std::string ARBDecompiler::TextureGather(Operation operation) {
    const auto& meta = std::get<MetaTexture>(operation.GetMeta());
    const u32 sampler_id = device.GetBaseBindings(stage).sampler + meta.sampler.index;
    const auto [coords, temporary, swizzle] = BuildCoords(operation);

    std::string comp;
    if (!meta.sampler.is_shadow) {
        const auto& immediate = std::get<ImmediateNode>(*meta.component);
        comp = fmt::format(".{}", Swizzle(immediate.GetValue()));
    }

    AddLine("TXG.F {}, {}, texture[{}]{}, {}{};", temporary, temporary, sampler_id, comp,
            TextureType(meta), BuildAoffi(operation));
    AddLine("MOV.U {}.x, {}.{};", temporary, coords, Swizzle(meta.element));
    return fmt::format("{}.x", temporary);
}

std::string ARBDecompiler::TextureQueryDimensions(Operation operation) {
    const auto& meta = std::get<MetaTexture>(operation.GetMeta());
    const std::string temporary = AllocVectorTemporary();
    const u32 sampler_id = device.GetBaseBindings(stage).sampler + meta.sampler.index;

    ASSERT(!meta.sampler.is_array);

    const std::string lod = operation.GetOperandsCount() > 0 ? Visit(operation[0]) : "0";
    AddLine("TXQ {}, {}, texture[{}], {};", temporary, lod, sampler_id, TextureType(meta));
    AddLine("MOV.U {}.x, {}.{};", temporary, temporary, Swizzle(meta.element));
    return fmt::format("{}.x", temporary);
}

std::string ARBDecompiler::TextureQueryLod(Operation operation) {
    const auto& meta = std::get<MetaTexture>(operation.GetMeta());
    const std::string temporary = AllocVectorTemporary();
    const u32 sampler_id = device.GetBaseBindings(stage).sampler + meta.sampler.index;

    ASSERT(!meta.sampler.is_array);

    const std::size_t count = operation.GetOperandsCount();
    for (std::size_t i = 0; i < count; ++i) {
        AddLine("MOV.F {}.{}, {};", temporary, Swizzle(i), Visit(operation[i]));
    }
    AddLine("LOD.F {}, {}, texture[{}], {};", temporary, temporary, sampler_id, TextureType(meta));
    AddLine("MUL.F32 {}, {}, {{256, 256, 0, 0}};", temporary, temporary);
    AddLine("TRUNC.S {}, {};", temporary, temporary);
    AddLine("MOV.U {}.x, {}.{};", temporary, temporary, Swizzle(meta.element));
    return fmt::format("{}.x", temporary);
}

std::string ARBDecompiler::TexelFetch(Operation operation) {
    const auto& meta = std::get<MetaTexture>(operation.GetMeta());
    const u32 sampler_id = device.GetBaseBindings(stage).sampler + meta.sampler.index;
    const auto [coords, temporary, swizzle] = BuildCoords(operation);

    if (!meta.sampler.is_buffer) {
        ASSERT(swizzle < 4);
        AddLine("MOV.F {}.w, {};", temporary, Visit(meta.lod));
    }
    AddLine("TXF.F {}, {}, texture[{}], {}{};", temporary, coords, sampler_id, TextureType(meta),
            BuildAoffi(operation));
    AddLine("MOV.U {}.x, {}.{};", temporary, temporary, Swizzle(meta.element));
    return fmt::format("{}.x", temporary);
}

std::string ARBDecompiler::TextureGradient(Operation operation) {
    const auto& meta = std::get<MetaTexture>(operation.GetMeta());
    const u32 sampler_id = device.GetBaseBindings(stage).sampler + meta.sampler.index;
    const std::string ddx = AllocVectorTemporary();
    const std::string ddy = AllocVectorTemporary();
    const std::string coord = std::get<1>(BuildCoords(operation));

    const std::size_t num_components = meta.derivates.size() / 2;
    for (std::size_t index = 0; index < num_components; ++index) {
        const char swizzle = Swizzle(index);
        AddLine("MOV.F {}.{}, {};", ddx, swizzle, Visit(meta.derivates[index * 2]));
        AddLine("MOV.F {}.{}, {};", ddy, swizzle, Visit(meta.derivates[index * 2 + 1]));
    }

    const std::string_view result = coord;
    AddLine("TXD.F {}, {}, {}, {}, texture[{}], {}{};", result, coord, ddx, ddy, sampler_id,
            TextureType(meta), BuildAoffi(operation));
    AddLine("MOV.F {}.x, {}.{};", result, result, Swizzle(meta.element));
    return fmt::format("{}.x", result);
}

std::string ARBDecompiler::ImageLoad(Operation operation) {
    const auto& meta = std::get<MetaImage>(operation.GetMeta());
    const u32 image_id = device.GetBaseBindings(stage).image + meta.image.index;
    const std::size_t count = operation.GetOperandsCount();
    const std::string_view type = ImageType(meta.image.type);

    const std::string temporary = AllocVectorTemporary();
    for (std::size_t i = 0; i < count; ++i) {
        AddLine("MOV.S {}.{}, {};", temporary, Swizzle(i), Visit(operation[i]));
    }
    AddLine("LOADIM.F {}, {}, image[{}], {};", temporary, temporary, image_id, type);
    AddLine("MOV.F {}.x, {}.{};", temporary, temporary, Swizzle(meta.element));
    return fmt::format("{}.x", temporary);
}

std::string ARBDecompiler::ImageStore(Operation operation) {
    const auto& meta = std::get<MetaImage>(operation.GetMeta());
    const u32 image_id = device.GetBaseBindings(stage).image + meta.image.index;
    const std::size_t num_coords = operation.GetOperandsCount();
    const std::size_t num_values = meta.values.size();
    const std::string_view type = ImageType(meta.image.type);

    const std::string coord = AllocVectorTemporary();
    const std::string value = AllocVectorTemporary();
    for (std::size_t i = 0; i < num_coords; ++i) {
        AddLine("MOV.S {}.{}, {};", coord, Swizzle(i), Visit(operation[i]));
    }
    for (std::size_t i = 0; i < num_values; ++i) {
        AddLine("MOV.F {}.{}, {};", value, Swizzle(i), Visit(meta.values[i]));
    }
    AddLine("STOREIM.F image[{}], {}, {}, {};", image_id, value, coord, type);
    return {};
}

std::string ARBDecompiler::Branch(Operation operation) {
    const auto target = std::get<ImmediateNode>(*operation[0]);
    AddLine("MOV.U PC.x, {};", target.GetValue());
    AddLine("CONT;");
    return {};
}

std::string ARBDecompiler::BranchIndirect(Operation operation) {
    AddLine("MOV.U PC.x, {};", Visit(operation[0]));
    AddLine("CONT;");
    return {};
}

std::string ARBDecompiler::PushFlowStack(Operation operation) {
    const auto stack = std::get<MetaStackClass>(operation.GetMeta());
    const u32 target = std::get<ImmediateNode>(*operation[0]).GetValue();
    const std::string_view stack_name = StackName(stack);
    AddLine("MOV.U {}[{}_TOP.x].x, {};", stack_name, stack_name, target);
    AddLine("ADD.S {}_TOP.x, {}_TOP.x, 1;", stack_name, stack_name);
    return {};
}

std::string ARBDecompiler::PopFlowStack(Operation operation) {
    const auto stack = std::get<MetaStackClass>(operation.GetMeta());
    const std::string_view stack_name = StackName(stack);
    AddLine("SUB.S {}_TOP.x, {}_TOP.x, 1;", stack_name, stack_name);
    AddLine("MOV.U PC.x, {}[{}_TOP.x].x;", stack_name, stack_name);
    AddLine("CONT;");
    return {};
}

std::string ARBDecompiler::Exit(Operation) {
    Exit();
    return {};
}

std::string ARBDecompiler::Discard(Operation) {
    AddLine("KIL TR;");
    return {};
}

std::string ARBDecompiler::EmitVertex(Operation) {
    AddLine("EMIT;");
    return {};
}

std::string ARBDecompiler::EndPrimitive(Operation) {
    AddLine("ENDPRIM;");
    return {};
}

std::string ARBDecompiler::InvocationId(Operation) {
    return "primitive.invocation";
}

std::string ARBDecompiler::YNegate(Operation) {
    LOG_WARNING(Render_OpenGL, "(STUBBED)");
    std::string temporary = AllocTemporary();
    AddLine("MOV.F {}, 1;", temporary);
    return temporary;
}

std::string ARBDecompiler::ThreadId(Operation) {
    return fmt::format("{}.threadid", StageInputName(stage));
}

std::string ARBDecompiler::ShuffleIndexed(Operation operation) {
    if (!device.HasWarpIntrinsics()) {
        LOG_ERROR(Render_OpenGL,
                  "NV_shader_thread_shuffle is missing. Kepler or better is required.");
        return Visit(operation[0]);
    }
    const std::string temporary = AllocVectorTemporary();
    AddLine("SHFIDX.U {}, {}, {}, {{31, 0, 0, 0}};", temporary, Visit(operation[0]),
            Visit(operation[1]));
    AddLine("MOV.U {}.x, {}.y;", temporary, temporary);
    return fmt::format("{}.x", temporary);
}

std::string ARBDecompiler::Barrier(Operation) {
    AddLine("BAR;");
    return {};
}

std::string ARBDecompiler::MemoryBarrierGroup(Operation) {
    AddLine("MEMBAR.CTA;");
    return {};
}

std::string ARBDecompiler::MemoryBarrierGlobal(Operation) {
    AddLine("MEMBAR;");
    return {};
}

} // Anonymous namespace

std::string DecompileAssemblyShader(const Device& device, const VideoCommon::Shader::ShaderIR& ir,
                                    const VideoCommon::Shader::Registry& registry,
                                    Tegra::Engines::ShaderType stage, std::string_view identifier) {
    return ARBDecompiler(device, ir, registry, stage, identifier).Code();
}

} // namespace OpenGL