code:
// SPDX-FileCopyrightText: Copyright 2021 yuzu Emulator Project
// SPDX-License-Identifier: GPL-2.0-or-later

#include "common/bit_field.h"
#include "common/common_types.h"
#include "shader_recompiler/frontend/ir/modifiers.h"
#include "shader_recompiler/frontend/maxwell/translate/impl/impl.h"

namespace Shader::Maxwell {
namespace {

enum class TextureType : u64 {
    _1D,
    ARRAY_1D,
    _2D,
    ARRAY_2D,
    _3D,
    ARRAY_3D,
    CUBE,
    ARRAY_CUBE,
};

Shader::TextureType GetType(TextureType type) {
    switch (type) {
    case TextureType::_1D:
        return Shader::TextureType::Color1D;
    case TextureType::ARRAY_1D:
        return Shader::TextureType::ColorArray1D;
    case TextureType::_2D:
        return Shader::TextureType::Color2D;
    case TextureType::ARRAY_2D:
        return Shader::TextureType::ColorArray2D;
    case TextureType::_3D:
        return Shader::TextureType::Color3D;
    case TextureType::ARRAY_3D:
        throw NotImplementedException("3D array texture type");
    case TextureType::CUBE:
        return Shader::TextureType::ColorCube;
    case TextureType::ARRAY_CUBE:
        return Shader::TextureType::ColorArrayCube;
    }
    throw NotImplementedException("Invalid texture type {}", type);
}

IR::Value MakeOffset(TranslatorVisitor& v, IR::Reg reg, bool has_lod_clamp) {
    const IR::U32 value{v.X(reg)};
    const u32 base{has_lod_clamp ? 12U : 16U};
    return v.ir.CompositeConstruct(
        v.ir.BitFieldExtract(value, v.ir.Imm32(base), v.ir.Imm32(4), true),
        v.ir.BitFieldExtract(value, v.ir.Imm32(base + 4), v.ir.Imm32(4), true));
}

void Impl(TranslatorVisitor& v, u64 insn, bool is_bindless) {
    union {
        u64 raw;
        BitField<49, 1, u64> nodep;
        BitField<35, 1, u64> aoffi;
        BitField<50, 1, u64> lc;
        BitField<51, 3, IR::Pred> sparse_pred;
        BitField<0, 8, IR::Reg> dest_reg;
        BitField<8, 8, IR::Reg> coord_reg;
        BitField<20, 8, IR::Reg> derivative_reg;
        BitField<28, 3, TextureType> type;
        BitField<31, 4, u64> mask;
        BitField<36, 13, u64> cbuf_offset;
    } const txd{insn};

    const bool has_lod_clamp = txd.lc != 0;
    if (has_lod_clamp) {
        throw NotImplementedException("TXD.LC - CLAMP is not implemented");
    }

    IR::Value coords;
    u32 num_derivatives{};
    IR::Reg base_reg{txd.coord_reg};
    IR::Reg last_reg;
    IR::Value handle;
    if (is_bindless) {
        handle = v.X(base_reg++);
    } else {
        handle = v.ir.Imm32(static_cast<u32>(txd.cbuf_offset.Value() * 4));
    }

    const auto read_array{[&]() -> IR::F32 {
        const IR::U32 base{v.ir.Imm32(0)};
        const IR::U32 count{v.ir.Imm32(has_lod_clamp ? 12 : 16)};
        const IR::U32 array_index{v.ir.BitFieldExtract(v.X(last_reg), base, count)};
        return v.ir.ConvertUToF(32, 16, array_index);
    }};
    switch (txd.type) {
    case TextureType::_1D: {
        coords = v.F(base_reg);
        num_derivatives = 1;
        last_reg = base_reg + 1;
        break;
    }
    case TextureType::ARRAY_1D: {
        last_reg = base_reg + 1;
        coords = v.ir.CompositeConstruct(v.F(base_reg), read_array());
        num_derivatives = 1;
        break;
    }
    case TextureType::_2D: {
        last_reg = base_reg + 2;
        coords = v.ir.CompositeConstruct(v.F(base_reg), v.F(base_reg + 1));
        num_derivatives = 2;
        break;
    }
    case TextureType::ARRAY_2D: {
        last_reg = base_reg + 2;
        coords = v.ir.CompositeConstruct(v.F(base_reg), v.F(base_reg + 1), read_array());
        num_derivatives = 2;
        break;
    }
    default:
        throw NotImplementedException("Invalid texture type");
    }

    const IR::Reg derivative_reg{txd.derivative_reg};
    IR::Value derivatives;
    switch (num_derivatives) {
    case 1: {
        derivatives = v.ir.CompositeConstruct(v.F(derivative_reg), v.F(derivative_reg + 1));
        break;
    }
    case 2: {
        derivatives = v.ir.CompositeConstruct(v.F(derivative_reg), v.F(derivative_reg + 1),
                                              v.F(derivative_reg + 2), v.F(derivative_reg + 3));
        break;
    }
    default:
        throw NotImplementedException("Invalid texture type");
    }

    IR::Value offset;
    if (txd.aoffi != 0) {
        offset = MakeOffset(v, last_reg, has_lod_clamp);
    }

    IR::F32 lod_clamp;
    if (has_lod_clamp) {
        // Lod Clamp is a Fixed Point 4.8, we need to transform it to float.
        // to convert a fixed point, float(value) / float(1 << fixed_point)
        // in this case the fixed_point is 8.
        const IR::F32 conv4_8fixp_f{v.ir.Imm32(static_cast<f32>(1U << 8))};
        const IR::F32 fixp_lc{v.ir.ConvertUToF(
            32, 16, v.ir.BitFieldExtract(v.X(last_reg), v.ir.Imm32(20), v.ir.Imm32(12)))};
        lod_clamp = v.ir.FPMul(fixp_lc, conv4_8fixp_f);
    }

    IR::TextureInstInfo info{};
    info.type.Assign(GetType(txd.type));
    info.num_derivatives.Assign(num_derivatives);
    info.has_lod_clamp.Assign(has_lod_clamp ? 1 : 0);
    const IR::Value sample{
        v.ir.ImageGradient(handle, coords, derivatives, offset, lod_clamp, info)};

    IR::Reg dest_reg{txd.dest_reg};
    for (size_t element = 0; element < 4; ++element) {
        if (((txd.mask >> element) & 1) == 0) {
            continue;
        }
        v.F(dest_reg, IR::F32{v.ir.CompositeExtract(sample, element)});
        ++dest_reg;
    }
    if (txd.sparse_pred != IR::Pred::PT) {
        v.ir.SetPred(txd.sparse_pred, v.ir.LogicalNot(v.ir.GetSparseFromOp(sample)));
    }
}
} // Anonymous namespace

void TranslatorVisitor::TXD(u64 insn) {
    Impl(*this, insn, false);
}

void TranslatorVisitor::TXD_b(u64 insn) {
    Impl(*this, insn, true);
}

} // namespace Shader::Maxwell


code:
// SPDX-FileCopyrightText: Copyright 2021 yuzu Emulator Project
// SPDX-License-Identifier: GPL-2.0-or-later

#include "common/bit_field.h"
#include "common/common_types.h"
#include "shader_recompiler/frontend/ir/modifiers.h"
#include "shader_recompiler/frontend/maxwell/translate/impl/impl.h"

namespace Shader::Maxwell {
namespace {

enum class TextureType : u64 {
    _1D,
    ARRAY_1D,
    _2D,
    ARRAY_2D,
    _3D,
    ARRAY_3D,
    CUBE,
    ARRAY_CUBE,
};

Shader::TextureType GetType(TextureType type) {
    switch (type) {
    case TextureType::_1D:
        return Shader::TextureType::Color1D;
    case TextureType::ARRAY_1D:
        return Shader::TextureType::ColorArray1D;
    case TextureType::_2D:
        return Shader::TextureType::Color2D;
    case TextureType::ARRAY_2D:
        return Shader::TextureType::ColorArray2D;
    case TextureType::_3D:
        return Shader::TextureType::Color3D;
    case TextureType::ARRAY_3D:
        throw NotImplementedException("3D array texture type");
    case TextureType::CUBE:
        return Shader::TextureType::ColorCube;
    case TextureType::ARRAY_CUBE:
        return Shader::TextureType::ColorArrayCube;
    }
    throw NotImplementedException("Invalid texture type {}", type);
}

IR::Value MakeCoords(TranslatorVisitor& v, IR::Reg reg, TextureType type) {
    const auto read_array{
        [&]() -> IR::U32 { return v.ir.BitFieldExtract(v.X(reg), v.ir.Imm32(0), v.ir.Imm32(16)); }};
    switch (type) {
    case TextureType::_1D:
        return v.X(reg);
    case TextureType::ARRAY_1D:
        return v.ir.CompositeConstruct(v.X(reg + 1), read_array());
    case TextureType::_2D:
        return v.ir.CompositeConstruct(v.X(reg), v.X(reg + 1));
    case TextureType::ARRAY_2D:
        return v.ir.CompositeConstruct(v.X(reg + 1), v.X(reg + 2), read_array());
    case TextureType::_3D:
        return v.ir.CompositeConstruct(v.X(reg), v.X(reg + 1), v.X(reg + 2));
    case TextureType::ARRAY_3D:
        throw NotImplementedException("3D array texture type");
    case TextureType::CUBE:
        return v.ir.CompositeConstruct(v.X(reg), v.X(reg + 1), v.X(reg + 2));
    case TextureType::ARRAY_CUBE:
        return v.ir.CompositeConstruct(v.X(reg + 1), v.X(reg + 2), v.X(reg + 3), read_array());
    }
    throw NotImplementedException("Invalid texture type {}", type);
}

IR::Value MakeOffset(TranslatorVisitor& v, IR::Reg& reg, TextureType type) {
    const IR::U32 value{v.X(reg++)};
    switch (type) {
    case TextureType::_1D:
    case TextureType::ARRAY_1D:
        return v.ir.BitFieldExtract(value, v.ir.Imm32(0), v.ir.Imm32(4), true);
    case TextureType::_2D:
    case TextureType::ARRAY_2D:
        return v.ir.CompositeConstruct(
            v.ir.BitFieldExtract(value, v.ir.Imm32(0), v.ir.Imm32(4), true),
            v.ir.BitFieldExtract(value, v.ir.Imm32(4), v.ir.Imm32(4), true));
    case TextureType::_3D:
    case TextureType::ARRAY_3D:
        return v.ir.CompositeConstruct(
            v.ir.BitFieldExtract(value, v.ir.Imm32(0), v.ir.Imm32(4), true),
            v.ir.BitFieldExtract(value, v.ir.Imm32(4), v.ir.Imm32(4), true),
            v.ir.BitFieldExtract(value, v.ir.Imm32(8), v.ir.Imm32(4), true));
    case TextureType::CUBE:
    case TextureType::ARRAY_CUBE:
        throw NotImplementedException("Illegal offset on CUBE sample");
    }
    throw NotImplementedException("Invalid texture type {}", type);
}

void Impl(TranslatorVisitor& v, u64 insn, bool is_bindless) {
    union {
        u64 raw;
        BitField<49, 1, u64> nodep;
        BitField<55, 1, u64> lod;
        BitField<50, 1, u64> multisample;
        BitField<35, 1, u64> aoffi;
        BitField<54, 1, u64> clamp;
        BitField<51, 3, IR::Pred> sparse_pred;
        BitField<0, 8, IR::Reg> dest_reg;
        BitField<8, 8, IR::Reg> coord_reg;
        BitField<20, 8, IR::Reg> meta_reg;
        BitField<28, 3, TextureType> type;
        BitField<31, 4, u64> mask;
        BitField<36, 13, u64> cbuf_offset;
    } const tld{insn};

    const IR::Value coords{MakeCoords(v, tld.coord_reg, tld.type)};

    IR::Reg meta_reg{tld.meta_reg};
    IR::Value handle;
    IR::Value offset;
    IR::U32 lod;
    IR::U32 multisample;
    if (is_bindless) {
        handle = v.X(meta_reg++);
    } else {
        handle = v.ir.Imm32(static_cast<u32>(tld.cbuf_offset.Value() * 4));
    }
    if (tld.lod != 0) {
        lod = v.X(meta_reg++);
    } else {
        lod = v.ir.Imm32(0U);
    }
    if (tld.aoffi != 0) {
        offset = MakeOffset(v, meta_reg, tld.type);
    }
    if (tld.multisample != 0) {
        multisample = v.X(meta_reg++);
    }
    if (tld.clamp != 0) {
        throw NotImplementedException("TLD.CL - CLAMP is not implemented");
    }
    IR::TextureInstInfo info{};
    info.type.Assign(GetType(tld.type));
    const IR::Value sample{v.ir.ImageFetch(handle, coords, offset, lod, multisample, info)};

    IR::Reg dest_reg{tld.dest_reg};
    for (size_t element = 0; element < 4; ++element) {
        if (((tld.mask >> element) & 1) == 0) {
            continue;
        }
        v.F(dest_reg, IR::F32{v.ir.CompositeExtract(sample, element)});
        ++dest_reg;
    }
    if (tld.sparse_pred != IR::Pred::PT) {
        v.ir.SetPred(tld.sparse_pred, v.ir.LogicalNot(v.ir.GetSparseFromOp(sample)));
    }
}
} // Anonymous namespace

void TranslatorVisitor::TLD(u64 insn) {
    Impl(*this, insn, false);
}

void TranslatorVisitor::TLD_b(u64 insn) {
    Impl(*this, insn, true);
}

} // namespace Shader::Maxwell

repo_name: litucks/torzu

license: mit


code:
// SPDX-FileCopyrightText: Copyright 2021 yuzu Emulator Project
// SPDX-License-Identifier: GPL-2.0-or-later

#include <array>

#include "common/bit_field.h"
#include "common/common_types.h"
#include "shader_recompiler/frontend/ir/modifiers.h"
#include "shader_recompiler/frontend/maxwell/translate/impl/impl.h"

namespace Shader::Maxwell {
namespace {
enum class Precision : u64 {
    F16,
    F32,
};

constexpr unsigned R = 1;
constexpr unsigned G = 2;
constexpr unsigned B = 4;
constexpr unsigned A = 8;

constexpr std::array RG_LUT{
    R,     //
    G,     //
    B,     //
    A,     //
    R | G, //
    R | A, //
    G | A, //
    B | A, //
};

constexpr std::array RGBA_LUT{
    R | G | B,     //
    R | G | A,     //
    R | B | A,     //
    G | B | A,     //
    R | G | B | A, //
};

union Encoding {
    u64 raw;
    BitField<59, 1, Precision> precision;
    BitField<54, 1, u64> aoffi;
    BitField<53, 1, u64> lod;
    BitField<55, 1, u64> ms;
    BitField<49, 1, u64> nodep;
    BitField<28, 8, IR::Reg> dest_reg_b;
    BitField<0, 8, IR::Reg> dest_reg_a;
    BitField<8, 8, IR::Reg> src_reg_a;
    BitField<20, 8, IR::Reg> src_reg_b;
    BitField<36, 13, u64> cbuf_offset;
    BitField<50, 3, u64> swizzle;
    BitField<53, 4, u64> encoding;
};

void CheckAlignment(IR::Reg reg, size_t alignment) {
    if (!IR::IsAligned(reg, alignment)) {
        throw NotImplementedException("Unaligned source register {}", reg);
    }
}

IR::Value MakeOffset(TranslatorVisitor& v, IR::Reg reg) {
    const IR::U32 value{v.X(reg)};
    return v.ir.CompositeConstruct(v.ir.BitFieldExtract(value, v.ir.Imm32(0), v.ir.Imm32(4), true),
                                   v.ir.BitFieldExtract(value, v.ir.Imm32(4), v.ir.Imm32(4), true));
}

IR::Value Sample(TranslatorVisitor& v, u64 insn) {
    const Encoding tlds{insn};
    const IR::U32 handle{v.ir.Imm32(static_cast<u32>(tlds.cbuf_offset * 4))};
    const IR::Reg reg_a{tlds.src_reg_a};
    const IR::Reg reg_b{tlds.src_reg_b};
    IR::Value coords;
    IR::U32 lod{v.ir.Imm32(0U)};
    IR::Value offsets;
    IR::U32 multisample;
    Shader::TextureType texture_type{};
    switch (tlds.encoding) {
    case 0:
        texture_type = Shader::TextureType::Color1D;
        coords = v.X(reg_a);
        break;
    case 1:
        texture_type = Shader::TextureType::Color1D;
        coords = v.X(reg_a);
        lod = v.X(reg_b);
        break;
    case 2:
        texture_type = Shader::TextureType::Color2D;
        coords = v.ir.CompositeConstruct(v.X(reg_a), v.X(reg_b));
        break;
    case 4:
        CheckAlignment(reg_a, 2);
        texture_type = Shader::TextureType::Color2D;
        coords = v.ir.CompositeConstruct(v.X(reg_a), v.X(reg_a + 1));
        offsets = MakeOffset(v, reg_b);
        break;
    case 5:
        CheckAlignment(reg_a, 2);
        texture_type = Shader::TextureType::Color2D;
        coords = v.ir.CompositeConstruct(v.X(reg_a), v.X(reg_a + 1));
        lod = v.X(reg_b);
        break;
    case 6:
        CheckAlignment(reg_a, 2);
        texture_type = Shader::TextureType::Color2D;
        coords = v.ir.CompositeConstruct(v.X(reg_a), v.X(reg_a + 1));
        multisample = v.X(reg_b);
        break;
    case 7:
        CheckAlignment(reg_a, 2);
        texture_type = Shader::TextureType::Color3D;
        coords = v.ir.CompositeConstruct(v.X(reg_a), v.X(reg_a + 1), v.X(reg_b));
        break;
    case 8: {
        CheckAlignment(reg_b, 2);
        const IR::U32 array{v.ir.BitFieldExtract(v.X(reg_a), v.ir.Imm32(0), v.ir.Imm32(16))};
        texture_type = Shader::TextureType::ColorArray2D;
        coords = v.ir.CompositeConstruct(v.X(reg_b), v.X(reg_b + 1), array);
        break;
    }
    case 12:
        CheckAlignment(reg_a, 2);
        CheckAlignment(reg_b, 2);
        texture_type = Shader::TextureType::Color2D;
        coords = v.ir.CompositeConstruct(v.X(reg_a), v.X(reg_a + 1));
        lod = v.X(reg_b);
        offsets = MakeOffset(v, reg_b + 1);
        break;
    default:
        throw NotImplementedException("Illegal encoding {}", tlds.encoding.Value());
    }
    IR::TextureInstInfo info{};
    if (tlds.precision == Precision::F16) {
        info.relaxed_precision.Assign(1);
    }
    info.type.Assign(texture_type);
    return v.ir.ImageFetch(handle, coords, offsets, lod, multisample, info);
}

unsigned Swizzle(u64 insn) {
    const Encoding tlds{insn};
    const size_t encoding{tlds.swizzle};
    if (tlds.dest_reg_b == IR::Reg::RZ) {
        if (encoding >= RG_LUT.size()) {
            throw NotImplementedException("Illegal RG encoding {}", encoding);
        }
        return RG_LUT[encoding];
    } else {
        if (encoding >= RGBA_LUT.size()) {
            throw NotImplementedException("Illegal RGBA encoding {}", encoding);
        }
        return RGBA_LUT[encoding];
    }
}

IR::F32 Extract(TranslatorVisitor& v, const IR::Value& sample, unsigned component) {
    return IR::F32{v.ir.CompositeExtract(sample, component)};
}

IR::Reg RegStoreComponent32(u64 insn, unsigned index) {
    const Encoding tlds{insn};
    switch (index) {
    case 0:
        return tlds.dest_reg_a;
    case 1:
        CheckAlignment(tlds.dest_reg_a, 2);
        return tlds.dest_reg_a + 1;
    case 2:
        return tlds.dest_reg_b;
    case 3:
        CheckAlignment(tlds.dest_reg_b, 2);
        return tlds.dest_reg_b + 1;
    }
    throw LogicError("Invalid store index {}", index);
}

void Store32(TranslatorVisitor& v, u64 insn, const IR::Value& sample) {
    const unsigned swizzle{Swizzle(insn)};
    unsigned store_index{0};
    for (unsigned component = 0; component < 4; ++component) {
        if (((swizzle >> component) & 1) == 0) {
            continue;
        }
        const IR::Reg dest{RegStoreComponent32(insn, store_index)};
        v.F(dest, Extract(v, sample, component));
        ++store_index;
    }
}

IR::U32 Pack(TranslatorVisitor& v, const IR::F32& lhs, const IR::F32& rhs) {
    return v.ir.PackHalf2x16(v.ir.CompositeConstruct(lhs, rhs));
}

void Store16(TranslatorVisitor& v, u64 insn, const IR::Value& sample) {
    const unsigned swizzle{Swizzle(insn)};
    unsigned store_index{0};
    std::array<IR::F32, 4> swizzled;
    for (unsigned component = 0; component < 4; ++component) {
        if (((swizzle >> component) & 1) == 0) {
            continue;
        }
        swizzled[store_index] = Extract(v, sample, component);
        ++store_index;
    }
    const IR::F32 zero{v.ir.Imm32(0.0f)};
    const Encoding tlds{insn};
    switch (store_index) {
    case 1:
        v.X(tlds.dest_reg_a, Pack(v, swizzled[0], zero));
        break;
    case 2:
    case 3:
    case 4:
        v.X(tlds.dest_reg_a, Pack(v, swizzled[0], swizzled[1]));
        switch (store_index) {
        case 2:
            break;
        case 3:
            v.X(tlds.dest_reg_b, Pack(v, swizzled[2], zero));
            break;
        case 4:
            v.X(tlds.dest_reg_b, Pack(v, swizzled[2], swizzled[3]));
            break;
        }
        break;
    }
}
} // Anonymous namespace

void TranslatorVisitor::TLDS(u64 insn) {
    const IR::Value sample{Sample(*this, insn)};
    if (Encoding{insn}.precision == Precision::F32) {
        Store32(*this, insn, sample);
    } else {
        Store16(*this, insn, sample);
    }
}
} // namespace Shader::Maxwell


code:
// SPDX-FileCopyrightText: Copyright 2021 yuzu Emulator Project
// SPDX-License-Identifier: GPL-2.0-or-later

#include "common/bit_field.h"
#include "common/common_types.h"
#include "shader_recompiler/frontend/ir/modifiers.h"
#include "shader_recompiler/frontend/maxwell/translate/impl/impl.h"

namespace Shader::Maxwell {
namespace {

enum class TextureType : u64 {
    _1D,
    ARRAY_1D,
    _2D,
    ARRAY_2D,
    _3D,
    ARRAY_3D,
    CUBE,
    ARRAY_CUBE,
};

Shader::TextureType GetType(TextureType type) {
    switch (type) {
    case TextureType::_1D:
        return Shader::TextureType::Color1D;
    case TextureType::ARRAY_1D:
        return Shader::TextureType::ColorArray1D;
    case TextureType::_2D:
        return Shader::TextureType::Color2D;
    case TextureType::ARRAY_2D:
        return Shader::TextureType::ColorArray2D;
    case TextureType::_3D:
        return Shader::TextureType::Color3D;
    case TextureType::ARRAY_3D:
        throw NotImplementedException("3D array texture type");
    case TextureType::CUBE:
        return Shader::TextureType::ColorCube;
    case TextureType::ARRAY_CUBE:
        return Shader::TextureType::ColorArrayCube;
    }
    throw NotImplementedException("Invalid texture type {}", type);
}

IR::Value MakeCoords(TranslatorVisitor& v, IR::Reg reg, TextureType type) {
    // The ISA reads an array component here, but this is not needed on high level shading languages
    // We are dropping this information.
    switch (type) {
    case TextureType::_1D:
        return v.F(reg);
    case TextureType::ARRAY_1D:
        return v.F(reg + 1);
    case TextureType::_2D:
        return v.ir.CompositeConstruct(v.F(reg), v.F(reg + 1));
    case TextureType::ARRAY_2D:
        return v.ir.CompositeConstruct(v.F(reg + 1), v.F(reg + 2));
    case TextureType::_3D:
        return v.ir.CompositeConstruct(v.F(reg), v.F(reg + 1), v.F(reg + 2));
    case TextureType::ARRAY_3D:
        throw NotImplementedException("3D array texture type");
    case TextureType::CUBE:
        return v.ir.CompositeConstruct(v.F(reg), v.F(reg + 1), v.F(reg + 2));
    case TextureType::ARRAY_CUBE:
        return v.ir.CompositeConstruct(v.F(reg + 1), v.F(reg + 2), v.F(reg + 3));
    }
    throw NotImplementedException("Invalid texture type {}", type);
}

void Impl(TranslatorVisitor& v, u64 insn, bool is_bindless) {
    union {
        u64 raw;
        BitField<49, 1, u64> nodep;
        BitField<35, 1, u64> ndv;
        BitField<0, 8, IR::Reg> dest_reg;
        BitField<8, 8, IR::Reg> coord_reg;
        BitField<20, 8, IR::Reg> meta_reg;
        BitField<28, 3, TextureType> type;
        BitField<31, 4, u64> mask;
        BitField<36, 13, u64> cbuf_offset;
    } const tmml{insn};

    if ((tmml.mask & 0b1100) != 0) {
        throw NotImplementedException("TMML BA results are not implemented");
    }
    const IR::Value coords{MakeCoords(v, tmml.coord_reg, tmml.type)};

    IR::U32 handle;
    IR::Reg meta_reg{tmml.meta_reg};
    if (is_bindless) {
        handle = v.X(meta_reg++);
    } else {
        handle = v.ir.Imm32(static_cast<u32>(tmml.cbuf_offset.Value() * 4));
    }
    IR::TextureInstInfo info{};
    info.type.Assign(GetType(tmml.type));
    const IR::Value sample{v.ir.ImageQueryLod(handle, coords, info)};

    IR::Reg dest_reg{tmml.dest_reg};
    for (size_t element = 0; element < 4; ++element) {
        if (((tmml.mask >> element) & 1) == 0) {
            continue;
        }
        IR::F32 value{v.ir.CompositeExtract(sample, element)};
        if (element < 2) {
            IR::U32 casted_value = v.ir.ConvertFToU(32, value);
            v.X(dest_reg, v.ir.ShiftLeftLogical(casted_value, v.ir.Imm32(8)));
        } else {
            v.F(dest_reg, value);
        }
        ++dest_reg;
    }
}
} // Anonymous namespace

void TranslatorVisitor::TMML(u64 insn) {
    Impl(*this, insn, false);
}

void TranslatorVisitor::TMML_b(u64 insn) {
    Impl(*this, insn, true);
}

} // namespace Shader::Maxwell


code:
// SPDX-FileCopyrightText: Copyright 2021 yuzu Emulator Project
// SPDX-License-Identifier: GPL-2.0-or-later

#include <optional>

#include "common/bit_field.h"
#include "common/common_types.h"
#include "shader_recompiler/frontend/maxwell/translate/impl/impl.h"

namespace Shader::Maxwell {
namespace {
enum class Mode : u64 {
    Dimension = 1,
    TextureType = 2,
    SamplePos = 5,
};

IR::Value Query(TranslatorVisitor& v, const IR::U32& handle, Mode mode, IR::Reg src_reg, u64 mask) {
    switch (mode) {
    case Mode::Dimension: {
        const bool needs_num_mips{((mask >> 3) & 1) != 0};
        const IR::U1 skip_mips{v.ir.Imm1(!needs_num_mips)};
        const IR::U32 lod{v.X(src_reg)};
        return v.ir.ImageQueryDimension(handle, lod, skip_mips);
    }
    case Mode::TextureType:
    case Mode::SamplePos:
    default:
        throw NotImplementedException("Mode {}", mode);
    }
}

void Impl(TranslatorVisitor& v, u64 insn, std::optional<u32> cbuf_offset) {
    union {
        u64 raw;
        BitField<49, 1, u64> nodep;
        BitField<0, 8, IR::Reg> dest_reg;
        BitField<8, 8, IR::Reg> src_reg;
        BitField<22, 3, Mode> mode;
        BitField<31, 4, u64> mask;
    } const txq{insn};

    IR::Reg src_reg{txq.src_reg};
    IR::U32 handle;
    if (cbuf_offset) {
        handle = v.ir.Imm32(*cbuf_offset);
    } else {
        handle = v.X(src_reg);
        ++src_reg;
    }
    const IR::Value query{Query(v, handle, txq.mode, src_reg, txq.mask)};
    IR::Reg dest_reg{txq.dest_reg};
    for (int element = 0; element < 4; ++element) {
        if (((txq.mask >> element) & 1) == 0) {
            continue;
        }
        v.X(dest_reg, IR::U32{v.ir.CompositeExtract(query, static_cast<size_t>(element))});
        ++dest_reg;
    }
}
} // Anonymous namespace

void TranslatorVisitor::TXQ(u64 insn) {
    union {
        u64 raw;
        BitField<36, 13, u64> cbuf_offset;
    } const txq{insn};

    Impl(*this, insn, static_cast<u32>(txq.cbuf_offset * 4));
}

void TranslatorVisitor::TXQ_b(u64 insn) {
    Impl(*this, insn, std::nullopt);
}

} // namespace Shader::Maxwell


code:
// SPDX-FileCopyrightText: Copyright 2021 yuzu Emulator Project
// SPDX-License-Identifier: GPL-2.0-or-later

#include "shader_recompiler/exception.h"
#include "shader_recompiler/frontend/maxwell/translate/impl/video_helper.h"

namespace Shader::Maxwell {

IR::U32 ExtractVideoOperandValue(IR::IREmitter& ir, const IR::U32& value, VideoWidth width,
                                 u32 selector, bool is_signed) {
    switch (width) {
    case VideoWidth::Byte:
    case VideoWidth::Unknown:
        return ir.BitFieldExtract(value, ir.Imm32(selector * 8), ir.Imm32(8), is_signed);
    case VideoWidth::Short:
        return ir.BitFieldExtract(value, ir.Imm32(selector * 16), ir.Imm32(16), is_signed);
    case VideoWidth::Word:
        return value;
    default:
        throw NotImplementedException("Unknown VideoWidth {}", width);
    }
}

VideoWidth GetVideoSourceWidth(VideoWidth width, bool is_immediate) {
    // immediates must be 16-bit format.
    return is_immediate ? VideoWidth::Short : width;
}

} // namespace Shader::Maxwell


code:
// SPDX-FileCopyrightText: Copyright 2021 yuzu Emulator Project
// SPDX-License-Identifier: GPL-2.0-or-later

#pragma once

#include "common/common_types.h"
#include "shader_recompiler/frontend/maxwell/translate/impl/impl.h"

namespace Shader::Maxwell {
enum class VideoWidth : u64 {
    Byte,
    Unknown,
    Short,
    Word,
};

[[nodiscard]] IR::U32 ExtractVideoOperandValue(IR::IREmitter& ir, const IR::U32& value,
                                               VideoWidth width, u32 selector, bool is_signed);

[[nodiscard]] VideoWidth GetVideoSourceWidth(VideoWidth width, bool is_immediate);

} // namespace Shader::Maxwell


code:
// SPDX-FileCopyrightText: Copyright 2021 yuzu Emulator Project
// SPDX-License-Identifier: GPL-2.0-or-later

#include "common/common_types.h"
#include "shader_recompiler/exception.h"
#include "shader_recompiler/frontend/maxwell/translate/impl/impl.h"
#include "shader_recompiler/frontend/maxwell/translate/impl/video_helper.h"

namespace Shader::Maxwell {
namespace {
enum class VideoMinMaxOps : u64 {
    MRG_16H,
    MRG_16L,
    MRG_8B0,
    MRG_8B2,
    ACC,
    MIN,
    MAX,
};

[[nodiscard]] IR::U32 ApplyVideoMinMaxOp(IR::IREmitter& ir, const IR::U32& lhs, const IR::U32& rhs,
                                         VideoMinMaxOps op, bool is_signed) {
    switch (op) {
    case VideoMinMaxOps::MIN:
        return ir.IMin(lhs, rhs, is_signed);
    case VideoMinMaxOps::MAX:
        return ir.IMax(lhs, rhs, is_signed);
    default:
        throw NotImplementedException("VMNMX op {}", op);
    }
}
} // Anonymous namespace

void TranslatorVisitor::VMNMX(u64 insn) {
    union {
        u64 raw;
        BitField<0, 8, IR::Reg> dest_reg;
        BitField<20, 16, u64> src_b_imm;
        BitField<28, 2, u64> src_b_selector;
        BitField<29, 2, VideoWidth> src_b_width;
        BitField<36, 2, u64> src_a_selector;
        BitField<37, 2, VideoWidth> src_a_width;
        BitField<47, 1, u64> cc;
        BitField<48, 1, u64> src_a_sign;
        BitField<49, 1, u64> src_b_sign;
        BitField<50, 1, u64> is_src_b_reg;
        BitField<51, 3, VideoMinMaxOps> op;
        BitField<54, 1, u64> dest_sign;
        BitField<55, 1, u64> sat;
        BitField<56, 1, u64> mx;
    } const vmnmx{insn};

    if (vmnmx.cc != 0) {
        throw NotImplementedException("VMNMX CC");
    }
    if (vmnmx.sat != 0) {
        throw NotImplementedException("VMNMX SAT");
    }

    const bool is_b_imm{vmnmx.is_src_b_reg == 0};
    const IR::U32 src_a{GetReg8(insn)};
    const IR::U32 src_b{is_b_imm ? ir.Imm32(static_cast<u32>(vmnmx.src_b_imm)) : GetReg20(insn)};
    const IR::U32 src_c{GetReg39(insn)};

    const VideoWidth a_width{vmnmx.src_a_width};
    const VideoWidth b_width{GetVideoSourceWidth(vmnmx.src_b_width, is_b_imm)};

    const u32 a_selector{static_cast<u32>(vmnmx.src_a_selector)};
    // Immediate values can't have a selector
    const u32 b_selector{is_b_imm ? 0U : static_cast<u32>(vmnmx.src_b_selector)};

    const bool src_a_signed{vmnmx.src_a_sign != 0};
    const bool src_b_signed{vmnmx.src_b_sign != 0};
    const IR::U32 op_a{ExtractVideoOperandValue(ir, src_a, a_width, a_selector, src_a_signed)};
    const IR::U32 op_b{ExtractVideoOperandValue(ir, src_b, b_width, b_selector, src_b_signed)};

    // First operation's sign is only dependent on operand b's sign
    const bool op_1_signed{src_b_signed};

    const IR::U32 lhs{vmnmx.mx != 0 ? ir.IMax(op_a, op_b, op_1_signed)
                                    : ir.IMin(op_a, op_b, op_1_signed)};
    X(vmnmx.dest_reg, ApplyVideoMinMaxOp(ir, lhs, src_c, vmnmx.op, vmnmx.dest_sign != 0));
}

} // namespace Shader::Maxwell


code:
// SPDX-FileCopyrightText: Copyright 2021 yuzu Emulator Project
// SPDX-License-Identifier: GPL-2.0-or-later

#include "common/common_types.h"
#include "shader_recompiler/exception.h"
#include "shader_recompiler/frontend/maxwell/translate/impl/impl.h"
#include "shader_recompiler/frontend/maxwell/translate/impl/video_helper.h"

namespace Shader::Maxwell {
void TranslatorVisitor::VMAD(u64 insn) {
    union {
        u64 raw;
        BitField<0, 8, IR::Reg> dest_reg;
        BitField<20, 16, u64> src_b_imm;
        BitField<28, 2, u64> src_b_selector;
        BitField<29, 2, VideoWidth> src_b_width;
        BitField<36, 2, u64> src_a_selector;
        BitField<37, 2, VideoWidth> src_a_width;
        BitField<47, 1, u64> cc;
        BitField<48, 1, u64> src_a_sign;
        BitField<49, 1, u64> src_b_sign;
        BitField<50, 1, u64> is_src_b_reg;
        BitField<51, 2, u64> scale;
        BitField<53, 1, u64> src_c_neg;
        BitField<54, 1, u64> src_a_neg;
        BitField<55, 1, u64> sat;
    } const vmad{insn};

    if (vmad.cc != 0) {
        throw NotImplementedException("VMAD CC");
    }
    if (vmad.sat != 0) {
        throw NotImplementedException("VMAD SAT");
    }
    if (vmad.scale != 0) {
        throw NotImplementedException("VMAD SCALE");
    }
    if (vmad.src_a_neg != 0 && vmad.src_c_neg != 0) {
        throw NotImplementedException("VMAD PO");
    }
    if (vmad.src_a_neg != 0 || vmad.src_c_neg != 0) {
        throw NotImplementedException("VMAD NEG");
    }
    const bool is_b_imm{vmad.is_src_b_reg == 0};
    const IR::U32 src_a{GetReg8(insn)};
    const IR::U32 src_b{is_b_imm ? ir.Imm32(static_cast<u32>(vmad.src_b_imm)) : GetReg20(insn)};
    const IR::U32 src_c{GetReg39(insn)};

    const u32 a_selector{static_cast<u32>(vmad.src_a_selector)};
    // Immediate values can't have a selector
    const u32 b_selector{is_b_imm ? 0U : static_cast<u32>(vmad.src_b_selector)};
    const VideoWidth a_width{vmad.src_a_width};
    const VideoWidth b_width{GetVideoSourceWidth(vmad.src_b_width, is_b_imm)};

    const bool src_a_signed{vmad.src_a_sign != 0};
    const bool src_b_signed{vmad.src_b_sign != 0};
    const IR::U32 op_a{ExtractVideoOperandValue(ir, src_a, a_width, a_selector, src_a_signed)};
    const IR::U32 op_b{ExtractVideoOperandValue(ir, src_b, b_width, b_selector, src_b_signed)};

    X(vmad.dest_reg, ir.IAdd(ir.IMul(op_a, op_b), src_c));
}

} // namespace Shader::Maxwell


code:
// SPDX-FileCopyrightText: Copyright 2021 yuzu Emulator Project
// SPDX-License-Identifier: GPL-2.0-or-later

#include "common/common_types.h"
#include "shader_recompiler/exception.h"
#include "shader_recompiler/frontend/maxwell/translate/impl/common_funcs.h"
#include "shader_recompiler/frontend/maxwell/translate/impl/impl.h"
#include "shader_recompiler/frontend/maxwell/translate/impl/video_helper.h"

namespace Shader::Maxwell {
namespace {
enum class VsetpCompareOp : u64 {
    False = 0,
    LessThan,
    Equal,
    LessThanEqual,
    GreaterThan = 16,
    NotEqual,
    GreaterThanEqual,
    True,
};

CompareOp VsetpToShaderCompareOp(VsetpCompareOp op) {
    switch (op) {
    case VsetpCompareOp::False:
        return CompareOp::False;
    case VsetpCompareOp::LessThan:
        return CompareOp::LessThan;
    case VsetpCompareOp::Equal:
        return CompareOp::Equal;
    case VsetpCompareOp::LessThanEqual:
        return CompareOp::LessThanEqual;
    case VsetpCompareOp::GreaterThan:
        return CompareOp::GreaterThan;
    case VsetpCompareOp::NotEqual:
        return CompareOp::NotEqual;
    case VsetpCompareOp::GreaterThanEqual:
        return CompareOp::GreaterThanEqual;
    case VsetpCompareOp::True:
        return CompareOp::True;
    default:
        throw NotImplementedException("Invalid compare op {}", op);
    }
}
} // Anonymous namespace

void TranslatorVisitor::VSETP(u64 insn) {
    union {
        u64 raw;
        BitField<0, 3, IR::Pred> dest_pred_b;
        BitField<3, 3, IR::Pred> dest_pred_a;
        BitField<20, 16, u64> src_b_imm;
        BitField<28, 2, u64> src_b_selector;
        BitField<29, 2, VideoWidth> src_b_width;
        BitField<36, 2, u64> src_a_selector;
        BitField<37, 2, VideoWidth> src_a_width;
        BitField<39, 3, IR::Pred> bop_pred;
        BitField<42, 1, u64> neg_bop_pred;
        BitField<43, 5, VsetpCompareOp> compare_op;
        BitField<45, 2, BooleanOp> bop;
        BitField<48, 1, u64> src_a_sign;
        BitField<49, 1, u64> src_b_sign;
        BitField<50, 1, u64> is_src_b_reg;
    } const vsetp{insn};

    const bool is_b_imm{vsetp.is_src_b_reg == 0};
    const IR::U32 src_a{GetReg8(insn)};
    const IR::U32 src_b{is_b_imm ? ir.Imm32(static_cast<u32>(vsetp.src_b_imm)) : GetReg20(insn)};

    const u32 a_selector{static_cast<u32>(vsetp.src_a_selector)};
    const u32 b_selector{static_cast<u32>(vsetp.src_b_selector)};
    const VideoWidth a_width{vsetp.src_a_width};
    const VideoWidth b_width{GetVideoSourceWidth(vsetp.src_b_width, is_b_imm)};

    const bool src_a_signed{vsetp.src_a_sign != 0};
    const bool src_b_signed{vsetp.src_b_sign != 0};
    const IR::U32 op_a{ExtractVideoOperandValue(ir, src_a, a_width, a_selector, src_a_signed)};
    const IR::U32 op_b{ExtractVideoOperandValue(ir, src_b, b_width, b_selector, src_b_signed)};

    // Compare operation's sign is only dependent on operand b's sign
    const bool compare_signed{src_b_signed};
    const CompareOp compare_op{VsetpToShaderCompareOp(vsetp.compare_op)};
    const IR::U1 comparison{IntegerCompare(ir, op_a, op_b, compare_op, compare_signed)};
    const IR::U1 bop_pred{ir.GetPred(vsetp.bop_pred, vsetp.neg_bop_pred != 0)};
    const IR::U1 result_a{PredicateCombine(ir, comparison, bop_pred, vsetp.bop)};
    const IR::U1 result_b{PredicateCombine(ir, ir.LogicalNot(comparison), bop_pred, vsetp.bop)};
    ir.SetPred(vsetp.dest_pred_a, result_a);
    ir.SetPred(vsetp.dest_pred_b, result_b);
}

} // namespace Shader::Maxwell


code:
// SPDX-FileCopyrightText: Copyright 2021 yuzu Emulator Project
// SPDX-License-Identifier: GPL-2.0-or-later

#include "common/bit_field.h"
#include "common/common_types.h"
#include "shader_recompiler/frontend/maxwell/translate/impl/impl.h"

namespace Shader::Maxwell {
namespace {
enum class VoteOp : u64 {
    ALL,
    ANY,
    EQ,
};

[[nodiscard]] IR::U1 VoteOperation(IR::IREmitter& ir, const IR::U1& pred, VoteOp vote_op) {
    switch (vote_op) {
    case VoteOp::ALL:
        return ir.VoteAll(pred);
    case VoteOp::ANY:
        return ir.VoteAny(pred);
    case VoteOp::EQ:
        return ir.VoteEqual(pred);
    default:
        throw NotImplementedException("Invalid VOTE op {}", vote_op);
    }
}

void Vote(TranslatorVisitor& v, u64 insn) {
    union {
        u64 insn;
        BitField<0, 8, IR::Reg> dest_reg;
        BitField<39, 3, IR::Pred> pred_a;
        BitField<42, 1, u64> neg_pred_a;
        BitField<45, 3, IR::Pred> pred_b;
        BitField<48, 2, VoteOp> vote_op;
    } const vote{insn};

    const IR::U1 vote_pred{v.ir.GetPred(vote.pred_a, vote.neg_pred_a != 0)};
    v.ir.SetPred(vote.pred_b, VoteOperation(v.ir, vote_pred, vote.vote_op));
    v.X(vote.dest_reg, v.ir.SubgroupBallot(vote_pred));
}
} // Anonymous namespace

void TranslatorVisitor::VOTE(u64 insn) {
    Vote(*this, insn);
}

void TranslatorVisitor::VOTE_vtg(u64) {
    LOG_WARNING(Shader, "(STUBBED) called");
}

} // namespace Shader::Maxwell


code:
// SPDX-FileCopyrightText: Copyright 2021 yuzu Emulator Project
// SPDX-License-Identifier: GPL-2.0-or-later

#include "common/bit_field.h"
#include "common/common_types.h"
#include "shader_recompiler/frontend/maxwell/translate/impl/impl.h"

namespace Shader::Maxwell {
namespace {
enum class ShuffleMode : u64 {
    IDX,
    UP,
    DOWN,
    BFLY,
};

[[nodiscard]] IR::U32 ShuffleOperation(IR::IREmitter& ir, const IR::U32& value,
                                       const IR::U32& index, const IR::U32& mask,
                                       ShuffleMode shfl_op) {
    const IR::U32 clamp{ir.BitFieldExtract(mask, ir.Imm32(0), ir.Imm32(5))};
    const IR::U32 seg_mask{ir.BitFieldExtract(mask, ir.Imm32(8), ir.Imm32(5))};
    switch (shfl_op) {
    case ShuffleMode::IDX:
        return ir.ShuffleIndex(value, index, clamp, seg_mask);
    case ShuffleMode::UP:
        return ir.ShuffleUp(value, index, clamp, seg_mask);
    case ShuffleMode::DOWN:
        return ir.ShuffleDown(value, index, clamp, seg_mask);
    case ShuffleMode::BFLY:
        return ir.ShuffleButterfly(value, index, clamp, seg_mask);
    default:
        throw NotImplementedException("Invalid SHFL op {}", shfl_op);
    }
}

void Shuffle(TranslatorVisitor& v, u64 insn, const IR::U32& index, const IR::U32& mask) {
    union {
        u64 insn;
        BitField<0, 8, IR::Reg> dest_reg;
        BitField<8, 8, IR::Reg> src_reg;
        BitField<30, 2, ShuffleMode> mode;
        BitField<48, 3, IR::Pred> pred;
    } const shfl{insn};

    const IR::U32 result{ShuffleOperation(v.ir, v.X(shfl.src_reg), index, mask, shfl.mode)};
    v.ir.SetPred(shfl.pred, v.ir.GetInBoundsFromOp(result));
    v.X(shfl.dest_reg, result);
}
} // Anonymous namespace

void TranslatorVisitor::SHFL(u64 insn) {
    union {
        u64 insn;
        BitField<20, 5, u64> src_a_imm;
        BitField<28, 1, u64> src_a_flag;
        BitField<29, 1, u64> src_b_flag;
        BitField<34, 13, u64> src_b_imm;
    } const flags{insn};
    const IR::U32 src_a{flags.src_a_flag != 0 ? ir.Imm32(static_cast<u32>(flags.src_a_imm))
                                              : GetReg20(insn)};
    const IR::U32 src_b{flags.src_b_flag != 0 ? ir.Imm32(static_cast<u32>(flags.src_b_imm))
                                              : GetReg39(insn)};
    Shuffle(*this, insn, src_a, src_b);
}

} // namespace Shader::Maxwell


code:
// SPDX-FileCopyrightText: Copyright 2021 yuzu Emulator Project
// SPDX-License-Identifier: GPL-2.0-or-later

#include "shader_recompiler/environment.h"
#include "shader_recompiler/frontend/ir/basic_block.h"
#include "shader_recompiler/frontend/maxwell/decode.h"
#include "shader_recompiler/frontend/maxwell/location.h"
#include "shader_recompiler/frontend/maxwell/translate/impl/impl.h"
#include "shader_recompiler/frontend/maxwell/translate/translate.h"

namespace Shader::Maxwell {

template <auto method>
static void Invoke(TranslatorVisitor& visitor, Location pc, u64 insn) {
    using MethodType = decltype(method);
    if constexpr (std::is_invocable_r_v<void, MethodType, TranslatorVisitor&, Location, u64>) {
        (visitor.*method)(pc, insn);
    } else if constexpr (std::is_invocable_r_v<void, MethodType, TranslatorVisitor&, u64>) {
        (visitor.*method)(insn);
    } else {
        (visitor.*method)();
    }
}

void Translate(Environment& env, IR::Block* block, u32 location_begin, u32 location_end) {
    if (location_begin == location_end) {
        return;
    }
    TranslatorVisitor visitor{env, *block};
    for (Location pc = location_begin; pc != location_end; ++pc) {
        const u64 insn{env.ReadInstruction(pc.Offset())};
        try {
            const Opcode opcode{Decode(insn)};
            switch (opcode) {
#define INST(name, cute, mask)                                                                     \
    case Opcode::name:                                                                             \
        Invoke<&TranslatorVisitor::name>(visitor, pc, insn);                                       \
        break;
#include "shader_recompiler/frontend/maxwell/maxwell.inc"
#undef OPCODE
            default:
                throw LogicError("Invalid opcode {}", opcode);
            }
        } catch (Exception& exception) {
            exception.Prepend(fmt::format("Translate {}: ", Decode(insn)));
            throw;
        }
    }
}

} // namespace Shader::Maxwell


code:
// SPDX-FileCopyrightText: Copyright 2021 yuzu Emulator Project
// SPDX-License-Identifier: GPL-2.0-or-later

#pragma once

#include "shader_recompiler/environment.h"
#include "shader_recompiler/frontend/ir/basic_block.h"

namespace Shader::Maxwell {

void Translate(Environment& env, IR::Block* block, u32 location_begin, u32 location_end);

} // namespace Shader::Maxwell


code:
// SPDX-FileCopyrightText: Copyright 2021 yuzu Emulator Project
// SPDX-License-Identifier: GPL-2.0-or-later

#include <algorithm>
#include <memory>
#include <vector>
#include <queue>

#include "common/settings.h"
#include "shader_recompiler/exception.h"
#include "shader_recompiler/frontend/ir/basic_block.h"
#include "shader_recompiler/frontend/ir/ir_emitter.h"
#include "shader_recompiler/frontend/ir/post_order.h"
#include "shader_recompiler/frontend/maxwell/structured_control_flow.h"
#include "shader_recompiler/frontend/maxwell/translate/translate.h"
#include "shader_recompiler/frontend/maxwell/translate_program.h"
#include "shader_recompiler/host_translate_info.h"
#include "shader_recompiler/ir_opt/passes.h"

namespace Shader::Maxwell {
namespace {
IR::BlockList GenerateBlocks(const IR::AbstractSyntaxList& syntax_list) {
    size_t num_syntax_blocks{};
    for (const auto& node : syntax_list) {
        if (node.type == IR::AbstractSyntaxNode::Type::Block) {
            ++num_syntax_blocks;
        }
    }
    IR::BlockList blocks;
    blocks.reserve(num_syntax_blocks);
    u32 order_index{};
    for (const auto& node : syntax_list) {
        if (node.type == IR::AbstractSyntaxNode::Type::Block) {
            blocks.push_back(node.data.block);
            blocks.back()->SetOrder(order_index++);
        }
    }
    return blocks;
}

void RemoveUnreachableBlocks(IR::Program& program) {
    // Some blocks might be unreachable if a function call exists unconditionally
    // If this happens the number of blocks and post order blocks will mismatch
    if (program.blocks.size() == program.post_order_blocks.size()) {
        return;
    }
    const auto begin{program.blocks.begin() + 1};
    const auto end{program.blocks.end()};
    const auto pred{[](IR::Block* block) { return block->ImmPredecessors().empty(); }};
    program.blocks.erase(std::remove_if(begin, end, pred), end);
}

void CollectInterpolationInfo(Environment& env, IR::Program& program) {
    if (program.stage != Stage::Fragment) {
        return;
    }
    const ProgramHeader& sph{env.SPH()};
    for (size_t index = 0; index < IR::NUM_GENERICS; ++index) {
        std::optional<PixelImap> imap;
        for (const PixelImap value : sph.ps.GenericInputMap(static_cast<u32>(index))) {
            if (value == PixelImap::Unused) {
                continue;
            }
            if (imap && imap != value) {
                throw NotImplementedException("Per component interpolation");
            }
            imap = value;
        }
        if (!imap) {
            continue;
        }
        program.info.interpolation[index] = [&] {
            switch (*imap) {
            case PixelImap::Unused:
            case PixelImap::Perspective:
                return Interpolation::Smooth;
            case PixelImap::Constant:
                return Interpolation::Flat;
            case PixelImap::ScreenLinear:
                return Interpolation::NoPerspective;
            }
            throw NotImplementedException("Unknown interpolation {}", *imap);
        }();
    }
}

void AddNVNStorageBuffers(IR::Program& program) {
    if (!program.info.uses_global_memory) {
        return;
    }
    const u32 driver_cbuf{0};
    const u32 descriptor_size{0x10};
    const u32 num_buffers{16};
    const u32 base{[&] {
        switch (program.stage) {
        case Stage::VertexA:
        case Stage::VertexB:
            return 0x110u;
        case Stage::TessellationControl:
            return 0x210u;
        case Stage::TessellationEval:
            return 0x310u;
        case Stage::Geometry:
            return 0x410u;
        case Stage::Fragment:
            return 0x510u;
        case Stage::Compute:
            return 0x310u;
        }
        throw InvalidArgument("Invalid stage {}", program.stage);
    }()};
    auto& descs{program.info.storage_buffers_descriptors};
    for (u32 index = 0; index < num_buffers; ++index) {
        if (!program.info.nvn_buffer_used[index]) {
            continue;
        }
        const u32 offset{base + index * descriptor_size};
        const auto it{std::ranges::find(descs, offset, &StorageBufferDescriptor::cbuf_offset)};
        if (it != descs.end()) {
            it->is_written |= program.info.stores_global_memory;
            continue;
        }
        descs.push_back({
            .cbuf_index = driver_cbuf,
            .cbuf_offset = offset,
            .count = 1,
            .is_written = program.info.stores_global_memory,
        });
    }
}

bool IsLegacyAttribute(IR::Attribute attribute) {
    return (attribute >= IR::Attribute::ColorFrontDiffuseR &&
            attribute <= IR::Attribute::ColorBackSpecularA) ||
           attribute == IR::Attribute::FogCoordinate ||
           (attribute >= IR::Attribute::FixedFncTexture0S &&
            attribute <= IR::Attribute::FixedFncTexture9Q);
}

std::map<IR::Attribute, IR::Attribute> GenerateLegacyToGenericMappings(
    const VaryingState& state, std::queue<IR::Attribute> unused_generics,
    const std::map<IR::Attribute, IR::Attribute>& previous_stage_mapping) {
    std::map<IR::Attribute, IR::Attribute> mapping;
    auto update_mapping = [&mapping, &unused_generics, previous_stage_mapping](IR::Attribute attr,
                                                                               size_t count) {
        if (previous_stage_mapping.find(attr) != previous_stage_mapping.end()) {
            for (size_t i = 0; i < count; ++i) {
                mapping.insert({attr + i, previous_stage_mapping.at(attr + i)});
            }
        } else {
            for (size_t i = 0; i < count; ++i) {
                mapping.insert({attr + i, unused_generics.front() + i});
            }
            unused_generics.pop();
        }
    };
    for (size_t index = 0; index < 4; ++index) {
        auto attr = IR::Attribute::ColorFrontDiffuseR + index * 4;
        if (state.AnyComponent(attr)) {
            update_mapping(attr, 4);
        }
    }
    if (state[IR::Attribute::FogCoordinate]) {
        update_mapping(IR::Attribute::FogCoordinate, 1);
    }
    for (size_t index = 0; index < IR::NUM_FIXEDFNCTEXTURE; ++index) {
        auto attr = IR::Attribute::FixedFncTexture0S + index * 4;
        if (state.AnyComponent(attr)) {
            update_mapping(attr, 4);
        }
    }
    return mapping;
}

void EmitGeometryPassthrough(IR::IREmitter& ir, const IR::Program& program,
                             const Shader::VaryingState& passthrough_mask,
                             bool passthrough_position,
                             std::optional<IR::Attribute> passthrough_layer_attr) {
    for (u32 i = 0; i < program.output_vertices; i++) {
        // Assign generics from input
        for (u32 j = 0; j < 32; j++) {
            if (!passthrough_mask.Generic(j)) {
                continue;
            }

            const IR::Attribute attr = IR::Attribute::Generic0X + (j * 4);
            ir.SetAttribute(attr + 0, ir.GetAttribute(attr + 0, ir.Imm32(i)), ir.Imm32(0));
            ir.SetAttribute(attr + 1, ir.GetAttribute(attr + 1, ir.Imm32(i)), ir.Imm32(0));
            ir.SetAttribute(attr + 2, ir.GetAttribute(attr + 2, ir.Imm32(i)), ir.Imm32(0));
            ir.SetAttribute(attr + 3, ir.GetAttribute(attr + 3, ir.Imm32(i)), ir.Imm32(0));
        }

        if (passthrough_position) {
            // Assign position from input
            const IR::Attribute attr = IR::Attribute::PositionX;
            ir.SetAttribute(attr + 0, ir.GetAttribute(attr + 0, ir.Imm32(i)), ir.Imm32(0));
            ir.SetAttribute(attr + 1, ir.GetAttribute(attr + 1, ir.Imm32(i)), ir.Imm32(0));
            ir.SetAttribute(attr + 2, ir.GetAttribute(attr + 2, ir.Imm32(i)), ir.Imm32(0));
            ir.SetAttribute(attr + 3, ir.GetAttribute(attr + 3, ir.Imm32(i)), ir.Imm32(0));
        }

        if (passthrough_layer_attr) {
            // Assign layer
            ir.SetAttribute(IR::Attribute::Layer, ir.GetAttribute(*passthrough_layer_attr),
                            ir.Imm32(0));
        }

        // Emit vertex
        ir.EmitVertex(ir.Imm32(0));
    }
    ir.EndPrimitive(ir.Imm32(0));
}

u32 GetOutputTopologyVertices(OutputTopology output_topology) {
    switch (output_topology) {
    case OutputTopology::PointList:
        return 1;
    case OutputTopology::LineStrip:
        return 2;
    default:
        return 3;
    }
}

void LowerGeometryPassthrough(const IR::Program& program, const HostTranslateInfo& host_info) {
    for (IR::Block* const block : program.blocks) {
        for (IR::Inst& inst : block->Instructions()) {
            if (inst.GetOpcode() == IR::Opcode::Epilogue) {
                IR::IREmitter ir{*block, IR::Block::InstructionList::s_iterator_to(inst)};
                EmitGeometryPassthrough(
                    ir, program, program.info.passthrough,
                    program.info.passthrough.AnyComponent(IR::Attribute::PositionX), {});
            }
        }
    }
}

} // Anonymous namespace

IR::Program TranslateProgram(ObjectPool<IR::Inst>& inst_pool, ObjectPool<IR::Block>& block_pool,
                             Environment& env, Flow::CFG& cfg, const HostTranslateInfo& host_info) {
    IR::Program program;
    program.syntax_list = BuildASL(inst_pool, block_pool, env, cfg, host_info);
    program.blocks = GenerateBlocks(program.syntax_list);
    program.post_order_blocks = PostOrder(program.syntax_list.front());
    program.stage = env.ShaderStage();
    program.local_memory_size = env.LocalMemorySize();
    switch (program.stage) {
    case Stage::TessellationControl: {
        const ProgramHeader& sph{env.SPH()};
        program.invocations = sph.common2.threads_per_input_primitive;
        break;
    }
    case Stage::Geometry: {
        const ProgramHeader& sph{env.SPH()};
        program.output_topology = sph.common3.output_topology;
        program.output_vertices = sph.common4.max_output_vertices;
        program.invocations = sph.common2.threads_per_input_primitive;
        program.is_geometry_passthrough = sph.common0.geometry_passthrough != 0;
        if (program.is_geometry_passthrough) {
            const auto& mask{env.GpPassthroughMask()};
            for (size_t i = 0; i < mask.size() * 32; ++i) {
                program.info.passthrough.mask[i] = ((mask[i / 32] >> (i % 32)) & 1) == 0;
            }

            if (!host_info.support_geometry_shader_passthrough) {
                program.output_vertices = GetOutputTopologyVertices(program.output_topology);
                LowerGeometryPassthrough(program, host_info);
            }
        }
        break;
    }
    case Stage::Compute:
        program.workgroup_size = env.WorkgroupSize();
        program.shared_memory_size = env.SharedMemorySize();
        break;
    default:
        break;
    }
    RemoveUnreachableBlocks(program);

    // Replace instructions before the SSA rewrite
    if (!host_info.support_float64) {
        Optimization::LowerFp64ToFp32(program);
    }
    if (!host_info.support_float16) {
        Optimization::LowerFp16ToFp32(program);
    }
    if (!host_info.support_int64) {
        Optimization::LowerInt64ToInt32(program);
    }
    if (!host_info.support_conditional_barrier) {
        Optimization::ConditionalBarrierPass(program);
    }
    Optimization::SsaRewritePass(program);

    Optimization::ConstantPropagationPass(env, program);

    Optimization::PositionPass(env, program);

    Optimization::GlobalMemoryToStorageBufferPass(program, host_info);
    Optimization::TexturePass(env, program, host_info);

    if (Settings::values.resolution_info.active) {
        Optimization::RescalingPass(program);
    }
    Optimization::DeadCodeEliminationPass(program);
    if (Settings::values.renderer_debug) {
        Optimization::VerificationPass(program);
    }
    Optimization::CollectShaderInfoPass(env, program);
    Optimization::LayerPass(program, host_info);
    Optimization::VendorWorkaroundPass(program);

    CollectInterpolationInfo(env, program);
    AddNVNStorageBuffers(program);
    return program;
}

IR::Program MergeDualVertexPrograms(IR::Program& vertex_a, IR::Program& vertex_b,
                                    Environment& env_vertex_b) {
    IR::Program result{};
    Optimization::VertexATransformPass(vertex_a);
    Optimization::VertexBTransformPass(vertex_b);
    for (const auto& term : vertex_a.syntax_list) {
        if (term.type != IR::AbstractSyntaxNode::Type::Return) {
            result.syntax_list.push_back(term);
        }
    }
    result.syntax_list.insert(result.syntax_list.end(), vertex_b.syntax_list.begin(),
                              vertex_b.syntax_list.end());
    result.blocks = GenerateBlocks(result.syntax_list);
    result.post_order_blocks = vertex_b.post_order_blocks;
    for (const auto& block : vertex_a.post_order_blocks) {
        result.post_order_blocks.push_back(block);
    }
    result.stage = Stage::VertexB;
    result.info = vertex_a.info;
    result.local_memory_size = std::max(vertex_a.local_memory_size, vertex_b.local_memory_size);
    result.info.loads.mask |= vertex_b.info.loads.mask;
    result.info.stores.mask |= vertex_b.info.stores.mask;

    Optimization::JoinTextureInfo(result.info, vertex_b.info);
    Optimization::JoinStorageInfo(result.info, vertex_b.info);
    Optimization::DeadCodeEliminationPass(result);
    if (Settings::values.renderer_debug) {
        Optimization::VerificationPass(result);
    }
    Optimization::CollectShaderInfoPass(env_vertex_b, result);
    return result;
}

void ConvertLegacyToGeneric(IR::Program& program, const Shader::RuntimeInfo& runtime_info) {
    auto& stores = program.info.stores;
    if (stores.Legacy()) {
        std::queue<IR::Attribute> unused_output_generics{};
        for (size_t index = 0; index < IR::NUM_GENERICS; ++index) {
            if (!stores.Generic(index)) {
                unused_output_generics.push(IR::Attribute::Generic0X + index * 4);
            }
        }
        program.info.legacy_stores_mapping =
            GenerateLegacyToGenericMappings(stores, unused_output_generics, {});
        for (IR::Block* const block : program.post_order_blocks) {
            for (IR::Inst& inst : block->Instructions()) {
                switch (inst.GetOpcode()) {
                case IR::Opcode::SetAttribute: {
                    const auto attr = inst.Arg(0).Attribute();
                    if (IsLegacyAttribute(attr)) {
                        stores.Set(program.info.legacy_stores_mapping[attr], true);
                        inst.SetArg(0, Shader::IR::Value(program.info.legacy_stores_mapping[attr]));
                    }
                    break;
                }
                default:
                    break;
                }
            }
        }
    }

    auto& loads = program.info.loads;
    if (loads.Legacy()) {
        std::queue<IR::Attribute> unused_input_generics{};
        for (size_t index = 0; index < IR::NUM_GENERICS; ++index) {
            const AttributeType input_type{runtime_info.generic_input_types[index]};
            if (!runtime_info.previous_stage_stores.Generic(index) || !loads.Generic(index) ||
                input_type == AttributeType::Disabled) {
                unused_input_generics.push(IR::Attribute::Generic0X + index * 4);
            }
        }
        auto mappings = GenerateLegacyToGenericMappings(
            loads, unused_input_generics, runtime_info.previous_stage_legacy_stores_mapping);
        for (IR::Block* const block : program.post_order_blocks) {
            for (IR::Inst& inst : block->Instructions()) {
                switch (inst.GetOpcode()) {
                case IR::Opcode::GetAttribute: {
                    const auto attr = inst.Arg(0).Attribute();
                    if (IsLegacyAttribute(attr)) {
                        loads.Set(mappings[attr], true);
                        inst.SetArg(0, Shader::IR::Value(mappings[attr]));
                    }
                    break;
                }
                default:
                    break;
                }
            }
        }
    }
}

IR::Program GenerateGeometryPassthrough(ObjectPool<IR::Inst>& inst_pool,
                                        ObjectPool<IR::Block>& block_pool,
                                        const HostTranslateInfo& host_info,
                                        IR::Program& source_program,
                                        Shader::OutputTopology output_topology) {
    IR::Program program;
    program.stage = Stage::Geometry;
    program.output_topology = output_topology;
    program.output_vertices = GetOutputTopologyVertices(output_topology);

    program.is_geometry_passthrough = false;
    program.info.loads.mask = source_program.info.stores.mask;
    program.info.stores.mask = source_program.info.stores.mask;
    program.info.stores.Set(IR::Attribute::Layer, true);
    program.info.stores.Set(source_program.info.emulated_layer, false);

    IR::Block* current_block = block_pool.Create(inst_pool);
    auto& node{program.syntax_list.emplace_back()};
    node.type = IR::AbstractSyntaxNode::Type::Block;
    node.data.block = current_block;

    IR::IREmitter ir{*current_block};
    EmitGeometryPassthrough(ir, program, program.info.stores, true,
                            source_program.info.emulated_layer);

    IR::Block* return_block{block_pool.Create(inst_pool)};
    IR::IREmitter{*return_block}.Epilogue();
    current_block->AddBranch(return_block);

    auto& merge{program.syntax_list.emplace_back()};
    merge.type = IR::AbstractSyntaxNode::Type::Block;
    merge.data.block = return_block;
    program.syntax_list.emplace_back().type = IR::AbstractSyntaxNode::Type::Return;

    program.blocks = GenerateBlocks(program.syntax_list);
    program.post_order_blocks = PostOrder(program.syntax_list.front());
    Optimization::SsaRewritePass(program);

    return program;
}

} // namespace Shader::Maxwell


code:
// SPDX-FileCopyrightText: Copyright 2021 yuzu Emulator Project
// SPDX-License-Identifier: GPL-2.0-or-later

#pragma once

#include "shader_recompiler/environment.h"
#include "shader_recompiler/frontend/ir/basic_block.h"
#include "shader_recompiler/frontend/ir/program.h"
#include "shader_recompiler/frontend/maxwell/control_flow.h"
#include "shader_recompiler/object_pool.h"
#include "shader_recompiler/runtime_info.h"

namespace Shader {
struct HostTranslateInfo;
}

namespace Shader::Maxwell {

[[nodiscard]] IR::Program TranslateProgram(ObjectPool<IR::Inst>& inst_pool,
                                           ObjectPool<IR::Block>& block_pool, Environment& env,
                                           Flow::CFG& cfg, const HostTranslateInfo& host_info);

[[nodiscard]] IR::Program MergeDualVertexPrograms(IR::Program& vertex_a, IR::Program& vertex_b,
                                                  Environment& env_vertex_b);

void ConvertLegacyToGeneric(IR::Program& program, const RuntimeInfo& runtime_info);

// Maxwell v1 and older Nvidia cards don't support setting gl_Layer from non-geometry stages.
// This creates a workaround by setting the layer as a generic output and creating a
// passthrough geometry shader that reads the generic and sets the layer.
[[nodiscard]] IR::Program GenerateGeometryPassthrough(ObjectPool<IR::Inst>& inst_pool,
                                                      ObjectPool<IR::Block>& block_pool,
                                                      const HostTranslateInfo& host_info,
                                                      IR::Program& source_program,
                                                      Shader::OutputTopology output_topology);

} // namespace Shader::Maxwell


code:
// SPDX-FileCopyrightText: Copyright 2021 yuzu Emulator Project
// SPDX-License-Identifier: GPL-2.0-or-later

#pragma once

namespace Shader {

// Try to keep entries here to a minimum
// They can accidentally change the cached information in a shader

/// Misc information about the host
struct HostTranslateInfo {
    bool support_float64{};      ///< True when the device supports 64-bit floats
    bool support_float16{};      ///< True when the device supports 16-bit floats
    bool support_int64{};        ///< True when the device supports 64-bit integers
    bool needs_demote_reorder{}; ///< True when the device needs DemoteToHelperInvocation reordered
    bool support_snorm_render_buffer{};  ///< True when the device supports SNORM render buffers
    bool support_viewport_index_layer{}; ///< True when the device supports gl_Layer in VS
    u32 min_ssbo_alignment{};            ///< Minimum alignment supported by the device for SSBOs
    bool support_geometry_shader_passthrough{}; ///< True when the device supports geometry
                                                ///< passthrough shaders
    bool support_conditional_barrier{}; ///< True when the device supports barriers in conditional
                                        ///< control flow
};

} // namespace Shader


code:
// SPDX-FileCopyrightText: Copyright 2021 yuzu Emulator Project
// SPDX-License-Identifier: GPL-2.0-or-later

#include "common/alignment.h"
#include "shader_recompiler/environment.h"
#include "shader_recompiler/frontend/ir/modifiers.h"
#include "shader_recompiler/frontend/ir/program.h"
#include "shader_recompiler/frontend/ir/value.h"
#include "shader_recompiler/ir_opt/passes.h"
#include "shader_recompiler/shader_info.h"

namespace Shader::Optimization {
namespace {
void AddConstantBufferDescriptor(Info& info, u32 index, u32 count) {
    if (count != 1) {
        throw NotImplementedException("Constant buffer descriptor indexing");
    }
    if ((info.constant_buffer_mask & (1U << index)) != 0) {
        return;
    }
    info.constant_buffer_mask |= 1U << index;

    auto& cbufs{info.constant_buffer_descriptors};
    cbufs.insert(std::ranges::lower_bound(cbufs, index, {}, &ConstantBufferDescriptor::index),
                 ConstantBufferDescriptor{
                     .index = index,
                     .count = 1,
                 });
}

void AddRegisterIndexedLdc(Info& info) {
    info.uses_cbuf_indirect = true;

    for (u32 i = 0; i < Info::MAX_INDIRECT_CBUFS; i++) {
        AddConstantBufferDescriptor(info, i, 1);

        // The shader can use any possible access size
        info.constant_buffer_used_sizes[i] = 0x10'000;
    }
}

u32 GetElementSize(IR::Type& used_type, Shader::IR::Opcode opcode) {
    switch (opcode) {
    case IR::Opcode::GetCbufU8:
    case IR::Opcode::GetCbufS8:
        used_type |= IR::Type::U8;
        return 1;
    case IR::Opcode::GetCbufU16:
    case IR::Opcode::GetCbufS16:
        used_type |= IR::Type::U16;
        return 2;
    case IR::Opcode::GetCbufU32:
        used_type |= IR::Type::U32;
        return 4;
    case IR::Opcode::GetCbufF32:
        used_type |= IR::Type::F32;
        return 4;
    case IR::Opcode::GetCbufU32x2:
        used_type |= IR::Type::U32x2;
        return 8;
    default:
        throw InvalidArgument("Invalid opcode {}", opcode);
    }
}

void GetPatch(Info& info, IR::Patch patch) {
    if (!IR::IsGeneric(patch)) {
        throw NotImplementedException("Reading non-generic patch {}", patch);
    }
    info.uses_patches.at(IR::GenericPatchIndex(patch)) = true;
}

void SetPatch(Info& info, IR::Patch patch) {
    if (IR::IsGeneric(patch)) {
        info.uses_patches.at(IR::GenericPatchIndex(patch)) = true;
        return;
    }
    switch (patch) {
    case IR::Patch::TessellationLodLeft:
    case IR::Patch::TessellationLodTop:
    case IR::Patch::TessellationLodRight:
    case IR::Patch::TessellationLodBottom:
        info.stores_tess_level_outer = true;
        break;
    case IR::Patch::TessellationLodInteriorU:
    case IR::Patch::TessellationLodInteriorV:
        info.stores_tess_level_inner = true;
        break;
    default:
        throw NotImplementedException("Set patch {}", patch);
    }
}

void CheckCBufNVN(Info& info, IR::Inst& inst) {
    const IR::Value cbuf_index{inst.Arg(0)};
    if (!cbuf_index.IsImmediate()) {
        info.nvn_buffer_used.set();
        return;
    }
    const u32 index{cbuf_index.U32()};
    if (index != 0) {
        return;
    }
    const IR::Value cbuf_offset{inst.Arg(1)};
    if (!cbuf_offset.IsImmediate()) {
        info.nvn_buffer_used.set();
        return;
    }
    const u32 offset{cbuf_offset.U32()};
    const u32 descriptor_size{0x10};
    const u32 upper_limit{info.nvn_buffer_base + descriptor_size * 16};
    if (offset >= info.nvn_buffer_base && offset < upper_limit) {
        const std::size_t nvn_index{(offset - info.nvn_buffer_base) / descriptor_size};
        info.nvn_buffer_used.set(nvn_index, true);
    }
}

void VisitUsages(Info& info, IR::Inst& inst) {
    switch (inst.GetOpcode()) {
    case IR::Opcode::CompositeConstructF16x2:
    case IR::Opcode::CompositeConstructF16x3:
    case IR::Opcode::CompositeConstructF16x4:
    case IR::Opcode::CompositeExtractF16x2:
    case IR::Opcode::CompositeExtractF16x3:
    case IR::Opcode::CompositeExtractF16x4:
    case IR::Opcode::CompositeInsertF16x2:
    case IR::Opcode::CompositeInsertF16x3:
    case IR::Opcode::CompositeInsertF16x4:
    case IR::Opcode::SelectF16:
    case IR::Opcode::BitCastU16F16:
    case IR::Opcode::BitCastF16U16:
    case IR::Opcode::PackFloat2x16:
    case IR::Opcode::UnpackFloat2x16:
    case IR::Opcode::ConvertS16F16:
    case IR::Opcode::ConvertS32F16:
    case IR::Opcode::ConvertS64F16:
    case IR::Opcode::ConvertU16F16:
    case IR::Opcode::ConvertU32F16:
    case IR::Opcode::ConvertU64F16:
    case IR::Opcode::ConvertF16S8:
    case IR::Opcode::ConvertF16S16:
    case IR::Opcode::ConvertF16S32:
    case IR::Opcode::ConvertF16S64:
    case IR::Opcode::ConvertF16U8:
    case IR::Opcode::ConvertF16U16:
    case IR::Opcode::ConvertF16U32:
    case IR::Opcode::ConvertF16U64:
    case IR::Opcode::ConvertF16F32:
    case IR::Opcode::ConvertF32F16:
    case IR::Opcode::FPAbs16:
    case IR::Opcode::FPAdd16:
    case IR::Opcode::FPCeil16:
    case IR::Opcode::FPFloor16:
    case IR::Opcode::FPFma16:
    case IR::Opcode::FPMul16:
    case IR::Opcode::FPNeg16:
    case IR::Opcode::FPRoundEven16:
    case IR::Opcode::FPSaturate16:
    case IR::Opcode::FPClamp16:
    case IR::Opcode::FPTrunc16:
    case IR::Opcode::FPOrdEqual16:
    case IR::Opcode::FPUnordEqual16:
    case IR::Opcode::FPOrdNotEqual16:
    case IR::Opcode::FPUnordNotEqual16:
    case IR::Opcode::FPOrdLessThan16:
    case IR::Opcode::FPUnordLessThan16:
    case IR::Opcode::FPOrdGreaterThan16:
    case IR::Opcode::FPUnordGreaterThan16:
    case IR::Opcode::FPOrdLessThanEqual16:
    case IR::Opcode::FPUnordLessThanEqual16:
    case IR::Opcode::FPOrdGreaterThanEqual16:
    case IR::Opcode::FPUnordGreaterThanEqual16:
    case IR::Opcode::FPIsNan16:
    case IR::Opcode::GlobalAtomicAddF16x2:
    case IR::Opcode::GlobalAtomicMinF16x2:
    case IR::Opcode::GlobalAtomicMaxF16x2:
    case IR::Opcode::StorageAtomicAddF16x2:
    case IR::Opcode::StorageAtomicMinF16x2:
    case IR::Opcode::StorageAtomicMaxF16x2:
        info.uses_fp16 = true;
        break;
    case IR::Opcode::CompositeConstructF64x2:
    case IR::Opcode::CompositeConstructF64x3:
    case IR::Opcode::CompositeConstructF64x4:
    case IR::Opcode::CompositeExtractF64x2:
    case IR::Opcode::CompositeExtractF64x3:
    case IR::Opcode::CompositeExtractF64x4:
    case IR::Opcode::CompositeInsertF64x2:
    case IR::Opcode::CompositeInsertF64x3:
    case IR::Opcode::CompositeInsertF64x4:
    case IR::Opcode::SelectF64:
    case IR::Opcode::BitCastU64F64:
    case IR::Opcode::BitCastF64U64:
    case IR::Opcode::PackDouble2x32:
    case IR::Opcode::UnpackDouble2x32:
    case IR::Opcode::FPAbs64:
    case IR::Opcode::FPAdd64:
    case IR::Opcode::FPCeil64:
    case IR::Opcode::FPFloor64:
    case IR::Opcode::FPFma64:
    case IR::Opcode::FPMax64:
    case IR::Opcode::FPMin64:
    case IR::Opcode::FPMul64:
    case IR::Opcode::FPNeg64:
    case IR::Opcode::FPRecip64:
    case IR::Opcode::FPRecipSqrt64:
    case IR::Opcode::FPRoundEven64:
    case IR::Opcode::FPSaturate64:
    case IR::Opcode::FPClamp64:
    case IR::Opcode::FPTrunc64:
    case IR::Opcode::FPOrdEqual64:
    case IR::Opcode::FPUnordEqual64:
    case IR::Opcode::FPOrdNotEqual64:
    case IR::Opcode::FPUnordNotEqual64:
    case IR::Opcode::FPOrdLessThan64:
    case IR::Opcode::FPUnordLessThan64:
    case IR::Opcode::FPOrdGreaterThan64:
    case IR::Opcode::FPUnordGreaterThan64:
    case IR::Opcode::FPOrdLessThanEqual64:
    case IR::Opcode::FPUnordLessThanEqual64:
    case IR::Opcode::FPOrdGreaterThanEqual64:
    case IR::Opcode::FPUnordGreaterThanEqual64:
    case IR::Opcode::FPIsNan64:
    case IR::Opcode::ConvertS16F64:
    case IR::Opcode::ConvertS32F64:
    case IR::Opcode::ConvertS64F64:
    case IR::Opcode::ConvertU16F64:
    case IR::Opcode::ConvertU32F64:
    case IR::Opcode::ConvertU64F64:
    case IR::Opcode::ConvertF32F64:
    case IR::Opcode::ConvertF64F32:
    case IR::Opcode::ConvertF64S8:
    case IR::Opcode::ConvertF64S16:
    case IR::Opcode::ConvertF64S32:
    case IR::Opcode::ConvertF64S64:
    case IR::Opcode::ConvertF64U8:
    case IR::Opcode::ConvertF64U16:
    case IR::Opcode::ConvertF64U32:
    case IR::Opcode::ConvertF64U64:
        info.uses_fp64 = true;
        break;
    default:
        break;
    }
    switch (inst.GetOpcode()) {
    case IR::Opcode::GetCbufU8:
    case IR::Opcode::GetCbufS8:
    case IR::Opcode::UndefU8:
    case IR::Opcode::LoadGlobalU8:
    case IR::Opcode::LoadGlobalS8:
    case IR::Opcode::WriteGlobalU8:
    case IR::Opcode::WriteGlobalS8:
    case IR::Opcode::LoadStorageU8:
    case IR::Opcode::LoadStorageS8:
    case IR::Opcode::WriteStorageU8:
    case IR::Opcode::WriteStorageS8:
    case IR::Opcode::LoadSharedU8:
    case IR::Opcode::LoadSharedS8:
    case IR::Opcode::WriteSharedU8:
    case IR::Opcode::SelectU8:
    case IR::Opcode::ConvertF16S8:
    case IR::Opcode::ConvertF16U8:
    case IR::Opcode::ConvertF32S8:
    case IR::Opcode::ConvertF32U8:
    case IR::Opcode::ConvertF64S8:
    case IR::Opcode::ConvertF64U8:
        info.uses_int8 = true;
        break;
    default:
        break;
    }
    switch (inst.GetOpcode()) {
    case IR::Opcode::GetCbufU16:
    case IR::Opcode::GetCbufS16:
    case IR::Opcode::UndefU16:
    case IR::Opcode::LoadGlobalU16:
    case IR::Opcode::LoadGlobalS16:
    case IR::Opcode::WriteGlobalU16:
    case IR::Opcode::WriteGlobalS16:
    case IR::Opcode::LoadStorageU16:
    case IR::Opcode::LoadStorageS16:
    case IR::Opcode::WriteStorageU16:
    case IR::Opcode::WriteStorageS16:
    case IR::Opcode::LoadSharedU16:
    case IR::Opcode::LoadSharedS16:
    case IR::Opcode::WriteSharedU16:
    case IR::Opcode::SelectU16:
    case IR::Opcode::BitCastU16F16:
    case IR::Opcode::BitCastF16U16:
    case IR::Opcode::ConvertS16F16:
    case IR::Opcode::ConvertS16F32:
    case IR::Opcode::ConvertS16F64:
    case IR::Opcode::ConvertU16F16:
    case IR::Opcode::ConvertU16F32:
    case IR::Opcode::ConvertU16F64:
    case IR::Opcode::ConvertF16S16:
    case IR::Opcode::ConvertF16U16:
    case IR::Opcode::ConvertF32S16:
    case IR::Opcode::ConvertF32U16:
    case IR::Opcode::ConvertF64S16:
    case IR::Opcode::ConvertF64U16:
        info.uses_int16 = true;
        break;
    default:
        break;
    }
    switch (inst.GetOpcode()) {
    case IR::Opcode::UndefU64:
    case IR::Opcode::LoadGlobalU8:
    case IR::Opcode::LoadGlobalS8:
    case IR::Opcode::LoadGlobalU16:
    case IR::Opcode::LoadGlobalS16:
    case IR::Opcode::LoadGlobal32:
    case IR::Opcode::LoadGlobal64:
    case IR::Opcode::LoadGlobal128:
    case IR::Opcode::WriteGlobalU8:
    case IR::Opcode::WriteGlobalS8:
    case IR::Opcode::WriteGlobalU16:
    case IR::Opcode::WriteGlobalS16:
    case IR::Opcode::WriteGlobal32:
    case IR::Opcode::WriteGlobal64:
    case IR::Opcode::WriteGlobal128:
    case IR::Opcode::SelectU64:
    case IR::Opcode::BitCastU64F64:
    case IR::Opcode::BitCastF64U64:
    case IR::Opcode::PackUint2x32:
    case IR::Opcode::UnpackUint2x32:
    case IR::Opcode::IAdd64:
    case IR::Opcode::ISub64:
    case IR::Opcode::INeg64:
    case IR::Opcode::ShiftLeftLogical64:
    case IR::Opcode::ShiftRightLogical64:
    case IR::Opcode::ShiftRightArithmetic64:
    case IR::Opcode::ConvertS64F16:
    case IR::Opcode::ConvertS64F32:
    case IR::Opcode::ConvertS64F64:
    case IR::Opcode::ConvertU64F16:
    case IR::Opcode::ConvertU64F32:
    case IR::Opcode::ConvertU64F64:
    case IR::Opcode::ConvertU64U32:
    case IR::Opcode::ConvertU32U64:
    case IR::Opcode::ConvertF16U64:
    case IR::Opcode::ConvertF32U64:
    case IR::Opcode::ConvertF64U64:
    case IR::Opcode::SharedAtomicExchange64:
    case IR::Opcode::GlobalAtomicIAdd64:
    case IR::Opcode::GlobalAtomicSMin64:
    case IR::Opcode::GlobalAtomicUMin64:
    case IR::Opcode::GlobalAtomicSMax64:
    case IR::Opcode::GlobalAtomicUMax64:
    case IR::Opcode::GlobalAtomicAnd64:
    case IR::Opcode::GlobalAtomicOr64:
    case IR::Opcode::GlobalAtomicXor64:
    case IR::Opcode::GlobalAtomicExchange64:
    case IR::Opcode::StorageAtomicIAdd64:
    case IR::Opcode::StorageAtomicSMin64:
    case IR::Opcode::StorageAtomicUMin64:
    case IR::Opcode::StorageAtomicSMax64:
    case IR::Opcode::StorageAtomicUMax64:
    case IR::Opcode::StorageAtomicAnd64:
    case IR::Opcode::StorageAtomicOr64:
    case IR::Opcode::StorageAtomicXor64:
    case IR::Opcode::StorageAtomicExchange64:
        info.uses_int64 = true;
        break;
    default:
        break;
    }
    switch (inst.GetOpcode()) {
    case IR::Opcode::WriteGlobalU8:
    case IR::Opcode::WriteGlobalS8:
    case IR::Opcode::WriteGlobalU16:
    case IR::Opcode::WriteGlobalS16:
    case IR::Opcode::WriteGlobal32:
    case IR::Opcode::WriteGlobal64:
    case IR::Opcode::WriteGlobal128:
    case IR::Opcode::GlobalAtomicIAdd32:
    case IR::Opcode::GlobalAtomicSMin32:
    case IR::Opcode::GlobalAtomicUMin32:
    case IR::Opcode::GlobalAtomicSMax32:
    case IR::Opcode::GlobalAtomicUMax32:
    case IR::Opcode::GlobalAtomicInc32:
    case IR::Opcode::GlobalAtomicDec32:
    case IR::Opcode::GlobalAtomicAnd32:
    case IR::Opcode::GlobalAtomicOr32:
    case IR::Opcode::GlobalAtomicXor32:
    case IR::Opcode::GlobalAtomicExchange32:
    case IR::Opcode::GlobalAtomicIAdd64:
    case IR::Opcode::GlobalAtomicSMin64:
    case IR::Opcode::GlobalAtomicUMin64:
    case IR::Opcode::GlobalAtomicSMax64:
    case IR::Opcode::GlobalAtomicUMax64:
    case IR::Opcode::GlobalAtomicAnd64:
    case IR::Opcode::GlobalAtomicOr64:
    case IR::Opcode::GlobalAtomicXor64:
    case IR::Opcode::GlobalAtomicExchange64:
    case IR::Opcode::GlobalAtomicIAdd32x2:
    case IR::Opcode::GlobalAtomicSMin32x2:
    case IR::Opcode::GlobalAtomicUMin32x2:
    case IR::Opcode::GlobalAtomicSMax32x2:
    case IR::Opcode::GlobalAtomicUMax32x2:
    case IR::Opcode::GlobalAtomicAnd32x2:
    case IR::Opcode::GlobalAtomicOr32x2:
    case IR::Opcode::GlobalAtomicXor32x2:
    case IR::Opcode::GlobalAtomicExchange32x2:
    case IR::Opcode::GlobalAtomicAddF32:
    case IR::Opcode::GlobalAtomicAddF16x2:
    case IR::Opcode::GlobalAtomicAddF32x2:
    case IR::Opcode::GlobalAtomicMinF16x2:
    case IR::Opcode::GlobalAtomicMinF32x2:
    case IR::Opcode::GlobalAtomicMaxF16x2:
    case IR::Opcode::GlobalAtomicMaxF32x2:
        info.stores_global_memory = true;
        [[fallthrough]];
    case IR::Opcode::LoadGlobalU8:
    case IR::Opcode::LoadGlobalS8:
    case IR::Opcode::LoadGlobalU16:
    case IR::Opcode::LoadGlobalS16:
    case IR::Opcode::LoadGlobal32:
    case IR::Opcode::LoadGlobal64:
    case IR::Opcode::LoadGlobal128:
        info.uses_int64 = true;
        info.uses_global_memory = true;
        info.used_constant_buffer_types |= IR::Type::U32 | IR::Type::U32x2;
        info.used_storage_buffer_types |= IR::Type::U32 | IR::Type::U32x2 | IR::Type::U32x4;
        break;
    case IR::Opcode::LoadLocal:
    case IR::Opcode::WriteLocal:
        info.uses_local_memory = true;
        break;
    default:
        break;
    }
    switch (inst.GetOpcode()) {
    case IR::Opcode::DemoteToHelperInvocation:
        info.uses_demote_to_helper_invocation = true;
        break;
    case IR::Opcode::GetAttribute:
    case IR::Opcode::GetAttributeU32:
        info.loads.mask[static_cast<size_t>(inst.Arg(0).Attribute())] = true;
        break;
    case IR::Opcode::SetAttribute:
        info.stores.mask[static_cast<size_t>(inst.Arg(0).Attribute())] = true;
        break;
    case IR::Opcode::GetPatch:
        GetPatch(info, inst.Arg(0).Patch());
        break;
    case IR::Opcode::SetPatch:
        SetPatch(info, inst.Arg(0).Patch());
        break;
    case IR::Opcode::GetAttributeIndexed:
        info.loads_indexed_attributes = true;
        break;
    case IR::Opcode::SetAttributeIndexed:
        info.stores_indexed_attributes = true;
        break;
    case IR::Opcode::SetFragColor:
        info.stores_frag_color[inst.Arg(0).U32()] = true;
        break;
    case IR::Opcode::SetSampleMask:
        info.stores_sample_mask = true;
        break;
    case IR::Opcode::SetFragDepth:
        info.stores_frag_depth = true;
        break;
    case IR::Opcode::WorkgroupId:
        info.uses_workgroup_id = true;
        break;
    case IR::Opcode::LocalInvocationId:
        info.uses_local_invocation_id = true;
        break;
    case IR::Opcode::InvocationId:
        info.uses_invocation_id = true;
        break;
    case IR::Opcode::InvocationInfo:
        info.uses_invocation_info = true;
        break;
    case IR::Opcode::SampleId:
        info.uses_sample_id = true;
        break;
    case IR::Opcode::IsHelperInvocation:
        info.uses_is_helper_invocation = true;
        break;
    case IR::Opcode::ResolutionDownFactor:
    case IR::Opcode::IsTextureScaled:
    case IR::Opcode::IsImageScaled:
        info.uses_rescaling_uniform = true;
        break;
    case IR::Opcode::LaneId:
        info.uses_subgroup_invocation_id = true;
        break;
    case IR::Opcode::ShuffleIndex:
    case IR::Opcode::ShuffleUp:
    case IR::Opcode::ShuffleDown:
    case IR::Opcode::ShuffleButterfly:
        info.uses_subgroup_shuffles = true;
        break;
    case IR::Opcode::GetCbufU8:
    case IR::Opcode::GetCbufS8:
    case IR::Opcode::GetCbufU16:
    case IR::Opcode::GetCbufS16:
    case IR::Opcode::GetCbufU32:
    case IR::Opcode::GetCbufF32:
    case IR::Opcode::GetCbufU32x2: {
        const IR::Value index{inst.Arg(0)};
        const IR::Value offset{inst.Arg(1)};
        if (index.IsImmediate()) {
            AddConstantBufferDescriptor(info, index.U32(), 1);
            u32 element_size = GetElementSize(info.used_constant_buffer_types, inst.GetOpcode());
            u32& size{info.constant_buffer_used_sizes[index.U32()]};
            if (offset.IsImmediate()) {
                size = Common::AlignUp(std::max(size, offset.U32() + element_size), 16u);
            } else {
                size = 0x10'000;
            }
        } else {
            AddRegisterIndexedLdc(info);
            GetElementSize(info.used_indirect_cbuf_types, inst.GetOpcode());
        }
        break;
    }
    case IR::Opcode::BindlessImageSampleImplicitLod:
    case IR::Opcode::BindlessImageSampleExplicitLod:
    case IR::Opcode::BindlessImageSampleDrefImplicitLod:
    case IR::Opcode::BindlessImageSampleDrefExplicitLod:
    case IR::Opcode::BindlessImageGather:
    case IR::Opcode::BindlessImageGatherDref:
    case IR::Opcode::BindlessImageFetch:
    case IR::Opcode::BindlessImageQueryDimensions:
    case IR::Opcode::BindlessImageQueryLod:
    case IR::Opcode::BindlessImageGradient:
    case IR::Opcode::BoundImageSampleImplicitLod:
    case IR::Opcode::BoundImageSampleExplicitLod:
    case IR::Opcode::BoundImageSampleDrefImplicitLod:
    case IR::Opcode::BoundImageSampleDrefExplicitLod:
    case IR::Opcode::BoundImageGather:
    case IR::Opcode::BoundImageGatherDref:
    case IR::Opcode::BoundImageFetch:
    case IR::Opcode::BoundImageQueryDimensions:
    case IR::Opcode::BoundImageQueryLod:
    case IR::Opcode::BoundImageGradient:
    case IR::Opcode::ImageGather:
    case IR::Opcode::ImageGatherDref:
    case IR::Opcode::ImageFetch:
    case IR::Opcode::ImageQueryDimensions:
    case IR::Opcode::ImageGradient: {
        const TextureType type{inst.Flags<IR::TextureInstInfo>().type};
        info.uses_sampled_1d |= type == TextureType::Color1D || type == TextureType::ColorArray1D;
        info.uses_sparse_residency |=
            inst.GetAssociatedPseudoOperation(IR::Opcode::GetSparseFromOp) != nullptr;
        break;
    }
    case IR::Opcode::ImageSampleImplicitLod:
    case IR::Opcode::ImageSampleExplicitLod:
    case IR::Opcode::ImageSampleDrefImplicitLod:
    case IR::Opcode::ImageSampleDrefExplicitLod:
    case IR::Opcode::ImageQueryLod: {
        const auto flags{inst.Flags<IR::TextureInstInfo>()};
        const TextureType type{flags.type};
        info.uses_sampled_1d |= type == TextureType::Color1D || type == TextureType::ColorArray1D;
        info.uses_shadow_lod |= flags.is_depth != 0;
        info.uses_sparse_residency |=
            inst.GetAssociatedPseudoOperation(IR::Opcode::GetSparseFromOp) != nullptr;
        break;
    }
    case IR::Opcode::ImageRead: {
        const auto flags{inst.Flags<IR::TextureInstInfo>()};
        info.uses_typeless_image_reads |= flags.image_format == ImageFormat::Typeless;
        info.uses_sparse_residency |=
            inst.GetAssociatedPseudoOperation(IR::Opcode::GetSparseFromOp) != nullptr;
        break;
    }
    case IR::Opcode::ImageWrite: {
        const auto flags{inst.Flags<IR::TextureInstInfo>()};
        info.uses_typeless_image_writes |= flags.image_format == ImageFormat::Typeless;
        info.uses_image_buffers |= flags.type == TextureType::Buffer;
        break;
    }
    case IR::Opcode::SubgroupEqMask:
    case IR::Opcode::SubgroupLtMask:
    case IR::Opcode::SubgroupLeMask:
    case IR::Opcode::SubgroupGtMask:
    case IR::Opcode::SubgroupGeMask:
        info.uses_subgroup_mask = true;
        break;
    case IR::Opcode::VoteAll:
    case IR::Opcode::VoteAny:
    case IR::Opcode::VoteEqual:
    case IR::Opcode::SubgroupBallot:
        info.uses_subgroup_vote = true;
        break;
    case IR::Opcode::FSwizzleAdd:
        info.uses_fswzadd = true;
        break;
    case IR::Opcode::DPdxFine:
    case IR::Opcode::DPdyFine:
    case IR::Opcode::DPdxCoarse:
    case IR::Opcode::DPdyCoarse:
        info.uses_derivatives = true;
        break;
    case IR::Opcode::LoadStorageU8:
    case IR::Opcode::LoadStorageS8:
    case IR::Opcode::WriteStorageU8:
    case IR::Opcode::WriteStorageS8:
        info.used_storage_buffer_types |= IR::Type::U8;
        break;
    case IR::Opcode::LoadStorageU16:
    case IR::Opcode::LoadStorageS16:
    case IR::Opcode::WriteStorageU16:
    case IR::Opcode::WriteStorageS16:
        info.used_storage_buffer_types |= IR::Type::U16;
        break;
    case IR::Opcode::LoadStorage32:
    case IR::Opcode::WriteStorage32:
    case IR::Opcode::StorageAtomicIAdd32:
    case IR::Opcode::StorageAtomicUMin32:
    case IR::Opcode::StorageAtomicUMax32:
    case IR::Opcode::StorageAtomicAnd32:
    case IR::Opcode::StorageAtomicOr32:
    case IR::Opcode::StorageAtomicXor32:
    case IR::Opcode::StorageAtomicExchange32:
        info.used_storage_buffer_types |= IR::Type::U32;
        break;
    case IR::Opcode::LoadStorage64:
    case IR::Opcode::WriteStorage64:
    case IR::Opcode::StorageAtomicIAdd32x2:
    case IR::Opcode::StorageAtomicSMin32x2:
    case IR::Opcode::StorageAtomicUMin32x2:
    case IR::Opcode::StorageAtomicSMax32x2:
    case IR::Opcode::StorageAtomicUMax32x2:
    case IR::Opcode::StorageAtomicAnd32x2:
    case IR::Opcode::StorageAtomicOr32x2:
    case IR::Opcode::StorageAtomicXor32x2:
    case IR::Opcode::StorageAtomicExchange32x2:
        info.used_storage_buffer_types |= IR::Type::U32x2;
        break;
    case IR::Opcode::LoadStorage128:
    case IR::Opcode::WriteStorage128:
        info.used_storage_buffer_types |= IR::Type::U32x4;
        break;
    case IR::Opcode::SharedAtomicSMin32:
        info.uses_atomic_s32_min = true;
        break;
    case IR::Opcode::SharedAtomicSMax32:
        info.uses_atomic_s32_max = true;
        break;
    case IR::Opcode::SharedAtomicInc32:
        info.uses_shared_increment = true;
        break;
    case IR::Opcode::SharedAtomicDec32:
        info.uses_shared_decrement = true;
        break;
    case IR::Opcode::SharedAtomicExchange64:
        info.uses_int64_bit_atomics = true;
        break;
    case IR::Opcode::GlobalAtomicInc32:
    case IR::Opcode::StorageAtomicInc32:
        info.used_storage_buffer_types |= IR::Type::U32;
        info.uses_global_increment = true;
        break;
    case IR::Opcode::GlobalAtomicDec32:
    case IR::Opcode::StorageAtomicDec32:
        info.used_storage_buffer_types |= IR::Type::U32;
        info.uses_global_decrement = true;
        break;
    case IR::Opcode::GlobalAtomicAddF32:
    case IR::Opcode::StorageAtomicAddF32:
        info.used_storage_buffer_types |= IR::Type::U32;
        info.uses_atomic_f32_add = true;
        break;
    case IR::Opcode::GlobalAtomicAddF16x2:
    case IR::Opcode::StorageAtomicAddF16x2:
        info.used_storage_buffer_types |= IR::Type::U32;
        info.uses_atomic_f16x2_add = true;
        break;
    case IR::Opcode::GlobalAtomicAddF32x2:
    case IR::Opcode::StorageAtomicAddF32x2:
        info.used_storage_buffer_types |= IR::Type::U32;
        info.uses_atomic_f32x2_add = true;
        break;
    case IR::Opcode::GlobalAtomicMinF16x2:
    case IR::Opcode::StorageAtomicMinF16x2:
        info.used_storage_buffer_types |= IR::Type::U32;
        info.uses_atomic_f16x2_min = true;
        break;
    case IR::Opcode::GlobalAtomicMinF32x2:
    case IR::Opcode::StorageAtomicMinF32x2:
        info.used_storage_buffer_types |= IR::Type::U32;
        info.uses_atomic_f32x2_min = true;
        break;
    case IR::Opcode::GlobalAtomicMaxF16x2:
    case IR::Opcode::StorageAtomicMaxF16x2:
        info.used_storage_buffer_types |= IR::Type::U32;
        info.uses_atomic_f16x2_max = true;
        break;
    case IR::Opcode::GlobalAtomicMaxF32x2:
    case IR::Opcode::StorageAtomicMaxF32x2:
        info.used_storage_buffer_types |= IR::Type::U32;
        info.uses_atomic_f32x2_max = true;
        break;
    case IR::Opcode::StorageAtomicSMin32:
        info.used_storage_buffer_types |= IR::Type::U32;
        info.uses_atomic_s32_min = true;
        break;
    case IR::Opcode::StorageAtomicSMax32:
        info.used_storage_buffer_types |= IR::Type::U32;
        info.uses_atomic_s32_max = true;
        break;
    case IR::Opcode::GlobalAtomicIAdd64:
    case IR::Opcode::GlobalAtomicSMin64:
    case IR::Opcode::GlobalAtomicUMin64:
    case IR::Opcode::GlobalAtomicSMax64:
    case IR::Opcode::GlobalAtomicUMax64:
    case IR::Opcode::GlobalAtomicAnd64:
    case IR::Opcode::GlobalAtomicOr64:
    case IR::Opcode::GlobalAtomicXor64:
    case IR::Opcode::GlobalAtomicExchange64:
    case IR::Opcode::StorageAtomicIAdd64:
    case IR::Opcode::StorageAtomicSMin64:
    case IR::Opcode::StorageAtomicUMin64:
    case IR::Opcode::StorageAtomicSMax64:
    case IR::Opcode::StorageAtomicUMax64:
    case IR::Opcode::StorageAtomicAnd64:
    case IR::Opcode::StorageAtomicOr64:
    case IR::Opcode::StorageAtomicXor64:
        info.used_storage_buffer_types |= IR::Type::U64 | IR::Type::U32x2;
        info.uses_int64_bit_atomics = true;
        break;
    case IR::Opcode::BindlessImageAtomicIAdd32:
    case IR::Opcode::BindlessImageAtomicSMin32:
    case IR::Opcode::BindlessImageAtomicUMin32:
    case IR::Opcode::BindlessImageAtomicSMax32:
    case IR::Opcode::BindlessImageAtomicUMax32:
    case IR::Opcode::BindlessImageAtomicInc32:
    case IR::Opcode::BindlessImageAtomicDec32:
    case IR::Opcode::BindlessImageAtomicAnd32:
    case IR::Opcode::BindlessImageAtomicOr32:
    case IR::Opcode::BindlessImageAtomicXor32:
    case IR::Opcode::BindlessImageAtomicExchange32:
    case IR::Opcode::BoundImageAtomicIAdd32:
    case IR::Opcode::BoundImageAtomicSMin32:
    case IR::Opcode::BoundImageAtomicUMin32:
    case IR::Opcode::BoundImageAtomicSMax32:
    case IR::Opcode::BoundImageAtomicUMax32:
    case IR::Opcode::BoundImageAtomicInc32:
    case IR::Opcode::BoundImageAtomicDec32:
    case IR::Opcode::BoundImageAtomicAnd32:
    case IR::Opcode::BoundImageAtomicOr32:
    case IR::Opcode::BoundImageAtomicXor32:
    case IR::Opcode::BoundImageAtomicExchange32:
    case IR::Opcode::ImageAtomicIAdd32:
    case IR::Opcode::ImageAtomicSMin32:
    case IR::Opcode::ImageAtomicUMin32:
    case IR::Opcode::ImageAtomicSMax32:
    case IR::Opcode::ImageAtomicUMax32:
    case IR::Opcode::ImageAtomicInc32:
    case IR::Opcode::ImageAtomicDec32:
    case IR::Opcode::ImageAtomicAnd32:
    case IR::Opcode::ImageAtomicOr32:
    case IR::Opcode::ImageAtomicXor32:
    case IR::Opcode::ImageAtomicExchange32:
        info.uses_atomic_image_u32 = true;
        break;
    default:
        break;
    }
}

void VisitFpModifiers(Info& info, IR::Inst& inst) {
    switch (inst.GetOpcode()) {
    case IR::Opcode::FPAdd16:
    case IR::Opcode::FPFma16:
    case IR::Opcode::FPMul16:
    case IR::Opcode::FPRoundEven16:
    case IR::Opcode::FPFloor16:
    case IR::Opcode::FPCeil16:
    case IR::Opcode::FPTrunc16: {
        const auto control{inst.Flags<IR::FpControl>()};
        switch (control.fmz_mode) {
        case IR::FmzMode::DontCare:
            break;
        case IR::FmzMode::FTZ:
        case IR::FmzMode::FMZ:
            info.uses_fp16_denorms_flush = true;
            break;
        case IR::FmzMode::None:
            info.uses_fp16_denorms_preserve = true;
            break;
        }
        break;
    }
    case IR::Opcode::FPAdd32:
    case IR::Opcode::FPFma32:
    case IR::Opcode::FPMul32:
    case IR::Opcode::FPRoundEven32:
    case IR::Opcode::FPFloor32:
    case IR::Opcode::FPCeil32:
    case IR::Opcode::FPTrunc32:
    case IR::Opcode::FPOrdEqual32:
    case IR::Opcode::FPUnordEqual32:
    case IR::Opcode::FPOrdNotEqual32:
    case IR::Opcode::FPUnordNotEqual32:
    case IR::Opcode::FPOrdLessThan32:
    case IR::Opcode::FPUnordLessThan32:
    case IR::Opcode::FPOrdGreaterThan32:
    case IR::Opcode::FPUnordGreaterThan32:
    case IR::Opcode::FPOrdLessThanEqual32:
    case IR::Opcode::FPUnordLessThanEqual32:
    case IR::Opcode::FPOrdGreaterThanEqual32:
    case IR::Opcode::FPUnordGreaterThanEqual32:
    case IR::Opcode::ConvertF16F32:
    case IR::Opcode::ConvertF64F32: {
        const auto control{inst.Flags<IR::FpControl>()};
        switch (control.fmz_mode) {
        case IR::FmzMode::DontCare:
            break;
        case IR::FmzMode::FTZ:
        case IR::FmzMode::FMZ:
            info.uses_fp32_denorms_flush = true;
            break;
        case IR::FmzMode::None:
            info.uses_fp32_denorms_preserve = true;
            break;
        }
        break;
    }
    default:
        break;
    }
}

void VisitCbufs(Info& info, IR::Inst& inst) {
    switch (inst.GetOpcode()) {
    case IR::Opcode::GetCbufU8:
    case IR::Opcode::GetCbufS8:
    case IR::Opcode::GetCbufU16:
    case IR::Opcode::GetCbufS16:
    case IR::Opcode::GetCbufU32:
    case IR::Opcode::GetCbufF32:
    case IR::Opcode::GetCbufU32x2: {
        CheckCBufNVN(info, inst);
        break;
    }
    default:
        break;
    }
}

void Visit(Info& info, IR::Inst& inst) {
    VisitUsages(info, inst);
    VisitFpModifiers(info, inst);
    VisitCbufs(info, inst);
}

void GatherInfoFromHeader(Environment& env, Info& info) {
    Stage stage{env.ShaderStage()};
    if (stage == Stage::Compute) {
        return;
    }
    const auto& header{env.SPH()};
    if (stage == Stage::Fragment) {
        if (!info.loads_indexed_attributes) {
            return;
        }
        for (size_t index = 0; index < IR::NUM_GENERICS; ++index) {
            const size_t offset{static_cast<size_t>(IR::Attribute::Generic0X) + index * 4};
            const auto vector{header.ps.imap_generic_vector[index]};
            info.loads.mask[offset + 0] = vector.x != PixelImap::Unused;
            info.loads.mask[offset + 1] = vector.y != PixelImap::Unused;
            info.loads.mask[offset + 2] = vector.z != PixelImap::Unused;
            info.loads.mask[offset + 3] = vector.w != PixelImap::Unused;
        }
        return;
    }
    if (info.loads_indexed_attributes) {
        for (size_t index = 0; index < IR::NUM_GENERICS; ++index) {
            const IR::Attribute attribute{IR::Attribute::Generic0X + index * 4};
            const auto mask = header.vtg.InputGeneric(index);
            for (size_t i = 0; i < 4; ++i) {
                info.loads.Set(attribute + i, mask[i]);
            }
        }
        for (size_t index = 0; index < 8; ++index) {
            const u16 mask{header.vtg.clip_distances};
            info.loads.Set(IR::Attribute::ClipDistance0 + index, ((mask >> index) & 1) != 0);
        }
        info.loads.Set(IR::Attribute::PrimitiveId, header.vtg.imap_systemb.primitive_array_id != 0);
        info.loads.Set(IR::Attribute::Layer, header.vtg.imap_systemb.rt_array_index != 0);
        info.loads.Set(IR::Attribute::ViewportIndex, header.vtg.imap_systemb.viewport_index != 0);
        info.loads.Set(IR::Attribute::PointSize, header.vtg.imap_systemb.point_size != 0);
        info.loads.Set(IR::Attribute::PositionX, header.vtg.imap_systemb.position_x != 0);
        info.loads.Set(IR::Attribute::PositionY, header.vtg.imap_systemb.position_y != 0);
        info.loads.Set(IR::Attribute::PositionZ, header.vtg.imap_systemb.position_z != 0);
        info.loads.Set(IR::Attribute::PositionW, header.vtg.imap_systemb.position_w != 0);
        info.loads.Set(IR::Attribute::PointSpriteS, header.vtg.point_sprite_s != 0);
        info.loads.Set(IR::Attribute::PointSpriteT, header.vtg.point_sprite_t != 0);
        info.loads.Set(IR::Attribute::FogCoordinate, header.vtg.fog_coordinate != 0);
        info.loads.Set(IR::Attribute::TessellationEvaluationPointU,
                       header.vtg.tessellation_eval_point_u != 0);
        info.loads.Set(IR::Attribute::TessellationEvaluationPointV,
                       header.vtg.tessellation_eval_point_v != 0);
        info.loads.Set(IR::Attribute::InstanceId, header.vtg.instance_id != 0);
        info.loads.Set(IR::Attribute::VertexId, header.vtg.vertex_id != 0);
        // TODO: Legacy varyings
    }
    if (info.stores_indexed_attributes) {
        for (size_t index = 0; index < IR::NUM_GENERICS; ++index) {
            const IR::Attribute attribute{IR::Attribute::Generic0X + index * 4};
            const auto mask{header.vtg.OutputGeneric(index)};
            for (size_t i = 0; i < 4; ++i) {
                info.stores.Set(attribute + i, mask[i]);
            }
        }
        for (size_t index = 0; index < 8; ++index) {
            const u16 mask{header.vtg.omap_systemc.clip_distances};
            const bool used{((mask >> index) & 1) != 0};
            info.stores.Set(IR::Attribute::ClipDistance0 + index, used);
            if (used) {
                info.used_clip_distances = static_cast<u32>(index) + 1;
            }
        }
        info.stores.Set(IR::Attribute::PrimitiveId,
                        header.vtg.omap_systemb.primitive_array_id != 0);
        info.stores.Set(IR::Attribute::Layer, header.vtg.omap_systemb.rt_array_index != 0);
        info.stores.Set(IR::Attribute::ViewportIndex, header.vtg.omap_systemb.viewport_index != 0);
        info.stores.Set(IR::Attribute::PointSize, header.vtg.omap_systemb.point_size != 0);
        info.stores.Set(IR::Attribute::PositionX, header.vtg.omap_systemb.position_x != 0);
        info.stores.Set(IR::Attribute::PositionY, header.vtg.omap_systemb.position_y != 0);
        info.stores.Set(IR::Attribute::PositionZ, header.vtg.omap_systemb.position_z != 0);
        info.stores.Set(IR::Attribute::PositionW, header.vtg.omap_systemb.position_w != 0);
        info.stores.Set(IR::Attribute::PointSpriteS, header.vtg.omap_systemc.point_sprite_s != 0);
        info.stores.Set(IR::Attribute::PointSpriteT, header.vtg.omap_systemc.point_sprite_t != 0);
        info.stores.Set(IR::Attribute::FogCoordinate, header.vtg.omap_systemc.fog_coordinate != 0);
        info.stores.Set(IR::Attribute::TessellationEvaluationPointU,
                        header.vtg.omap_systemc.tessellation_eval_point_u != 0);
        info.stores.Set(IR::Attribute::TessellationEvaluationPointV,
                        header.vtg.omap_systemc.tessellation_eval_point_v != 0);
        info.stores.Set(IR::Attribute::InstanceId, header.vtg.omap_systemc.instance_id != 0);
        info.stores.Set(IR::Attribute::VertexId, header.vtg.omap_systemc.vertex_id != 0);
        // TODO: Legacy varyings
    }
}
} // Anonymous namespace

void CollectShaderInfoPass(Environment& env, IR::Program& program) {
    Info& info{program.info};
    const u32 base{[&] {
        switch (program.stage) {
        case Stage::VertexA:
        case Stage::VertexB:
            return 0x110u;
        case Stage::TessellationControl:
            return 0x210u;
        case Stage::TessellationEval:
            return 0x310u;
        case Stage::Geometry:
            return 0x410u;
        case Stage::Fragment:
            return 0x510u;
        case Stage::Compute:
            return 0x310u;
        }
        throw InvalidArgument("Invalid stage {}", program.stage);
    }()};
    info.nvn_buffer_base = base;

    for (IR::Block* const block : program.post_order_blocks) {
        for (IR::Inst& inst : block->Instructions()) {
            Visit(info, inst);
        }
    }
    GatherInfoFromHeader(env, info);
}

} // namespace Shader::Optimization